     for (uint32 i = begin + 1; i < end; ++i)
     {
         const uint2  i_span = chain_ranges[ chain_index[i] ];
-        const uint32 i_w    = chain_weights[ chain_index[i] ];
+        const uint32 i_w    = chain_weights[ i ] & 0xFFFFFFFFu;               // already sorted as chain_index
 
         uint32 j;
         for (j = begin; j < begin + n; ++j)
         {
             const uint2  j_span = chain_ranges[ chain_index[j] ];
-            const uint32 j_w    = chain_weights[ chain_index[j] ];
+            const uint32 j_w    = chain_weights[ j ] & 0xFFFFFFFFu;           // already sorted as chain_index
 
             const uint32 max_begin = nvbio::max( i_span.x, j_span.x );
             const uint32 min_end   = nvbio::min( i_span.y, j_span.y );

__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;



        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 ),
