 
         chain_coverage_kernel<<<n_blocks, block_dim>>>(
             n_chains,
+            nvbio::plain_view( mem->chain_reads ),
             nvbio::plain_view( mem->chain_offsets ),
             nvbio::plain_view( mem->chain_lengths ),
             nvbio::plain_view( mem->mems ),

            pipeline->chunk,
            n_chains,
            nvbio::plain_view( mem->chain_reads ),
            nvbio::plain_view( mem->chain_offsets ),
            nvbio::plain_view( mem->chain_lengths ),
            nvbio::plain_view( chain_index ),
            nvbio::plain_view( chain_ranges ),
            nvbio::plain_view( chain_weights ),

__global__
void chain_coverage_kernel(
    const uint32                                    n_chains,           // the number of chains
    const uint32*                                   chain_reads,        // the chain reads
    const uint32*                                   chain_offsets,      // the chain offsets
    const uint32*                                   chain_lengths,      // the chain lengths
    const mem_state::mem_type*                      mems,               // the MEMs for this chunk of reads
    const uint32*                                   mems_index,         // a sorting index into the MEMs specifying their processing order
          uint2*                                    chain_ranges,       // the output chain ranges
          uint64*                                   chain_weights)      // the output chain weights
{
    const uint32 chain_id = threadIdx.x + blockIdx.x * blockDim.x;
    if (chain_id >= n_chains)
        return;

    const uint32 read  = chain_reads[ chain_id ];
    const uint32 begin = chain_offsets[ chain_id ];
    const uint32 end   = chain_lengths[ chain_id ] + begin;

