     const read_chunk                                chunk,              // the current sub-batch
     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 uint32*                                   chain_index,        // the chain order
     const uint2*                                    chain_ranges,       // the chain ranges
-    const uint32*                                   chain_weights,      // the chain weights
+    const uint64*                                   chain_weights,      // the chain weights
     const float                                     mask_level,         // input option
     const float                                     chain_drop_ratio,   // input option
     const uint32                                    min_seed_len,       // input option

__global__
void chain_coverage_kernel(
    const uint32                                    n_chains,           // the number of chains
    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
          uint32*                                   chain_weights)      // the output chain weights
{
    const uint32 chain_id = threadIdx.x + blockIdx.x * blockDim.x;
    if (chain_id >= n_chains)
        return;

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



    // write out the outputs
    chain_ranges[ chain_id ]  = range;
    chain_weights[ chain_id ] = uint64( weight ) | (uint64( read ) << 32);
}

// filter the chains belonging to each read
