 
     typedef nvbio::cuda::ldg_pointer<uint32>                                        base_iterator;
 
-    typedef nvbio::PackedStringLoader<base_iterator,4,false,cache_type>             pattern_loader_type;
-    typedef typename pattern_loader_type::iterator                                  pattern_iterator;
-    typedef nvbio::vector_wrapper<pattern_iterator>                                 pattern_string;
-
-    typedef nvbio::PackedStringLoader<base_iterator,2,false,uncached_tag_type>      text_loader_type;
-    typedef typename text_loader_type::iterator                                     text_iterator;
-    typedef nvbio::vector_wrapper<text_iterator>                                    text_string;
+    typedef nvbio::PackedStringLoader<base_iterator,4,io::ReadData::HI_BITS,cache_type> pattern_loader_type;
+    typedef typename pattern_loader_type::iterator                                      pattern_iterator;
+    typedef nvbio::vector_wrapper<pattern_iterator>                                     pattern_string;
+
+    typedef nvbio::PackedStringLoader<base_iterator,2,false,uncached_tag_type>          text_loader_type;
+    typedef typename text_loader_type::iterator                                         text_iterator;
+    typedef nvbio::vector_wrapper<text_iterator>                                        text_string;
 
     // an alignment context
     struct context_type


// A simple kernel to test the speed of alignment without the possible overheads of the BatchAlignmentScore interface
//
template <uint32 BLOCKDIM, uint32 MAX_REF_LEN, typename aligner_type, typename score_type>
__global__ void alignment_test_kernel(const aligner_type aligner, const uint32 N_probs, const uint32 M, const uint32 N, const uint32* strptr, const uint32* refptr, score_type* score)
{
    const uint32 tid = blockIdx.x * BLOCKDIM + threadIdx.x;

    typedef lmem_cache_tag_type                                                 lmem_cache_type;
    typedef nvbio::cuda::ldg_pointer<uint32>                                    base_iterator;

    typedef nvbio::PackedStringLoader<base_iterator,4,false,lmem_cache_type>    pattern_loader_type;
    typedef typename pattern_loader_type::iterator                              pattern_iterator;
    typedef nvbio::vector_wrapper<pattern_iterator>                             pattern_string;

    typedef nvbio::PackedStringLoader<base_iterator,2,false,lmem_cache_type>    text_loader_type;
    typedef typename text_loader_type::iterator                                 text_iterator;
    typedef nvbio::vector_wrapper<text_iterator>                                text_string;

    pattern_loader_type pattern_loader;
    pattern_string pattern = pattern_string( M, pattern_loader.load( strptr, tid * M, tid < N_probs ? M : 0u ) );

    text_loader_type text_loader;
    text_string text = text_string( N, text_loader.load( strptr, tid * N, tid < N_probs ? N : 0u ) );

    aln::BestSink<int32> sink;

    aln::alignment_score<MAX_REF_LEN>(
        aligner,
        pattern,
        aln::trivial_quality_string(),

            score_dvec );

        // test the ThreadParallelScheduler
        batch_score_profile<aln::ThreadParallelScheduler>( stream );

        // test the StagedThreadParallelScheduler
        //batch_score_profile<aln::StagedThreadParallelScheduler>( stream );
    }
    {
        const uint32 BLOCKDIM = 128;
        const uint32 N_BLOCKS = (n_tasks + BLOCKDIM-1) / BLOCKDIM;

        Timer timer;
        timer.start();

        // enact the batch
        alignment_test_kernel<BLOCKDIM,400> <<<N_BLOCKS,BLOCKDIM>>>(
            aligner,
            n_tasks,
            offsets_dvec,
            pattern_dvec,
            text_len,
            text_dvec,
            score_dvec );

        cudaDeviceSynchronize();

        timer.stop();

        const float time = timer.seconds();

        fprintf(stderr,"  %5.1f", 1.0e-9f * float(n_tasks*uint64(max_pattern_len)*uint64(text_len))/time );
    }
    fprintf(stderr, " GCUPS\n");
}
