             tile_row_end_offsets[tile_num_rows + threadIdx.x] = num_nonzeros;
         }
 
-        // Load a tile of nonzeros
         #pragma unroll
         for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
         {
             int item_idx                = threadIdx.x + (ITEM * BLOCK_THREADS);
             item_idx                    = CUB_MIN(item_idx, tile_num_nonzeros - 1);
-
+        
             OffsetT column_index        = d_matrix_column_indices[tile_coord.y + item_idx];
             ValueT  matrix_value        = d_matrix_values[tile_coord.y + item_idx];
             ValueT  vector_value        = d_vector_x[column_index];
-
             tile_nonzeros[item_idx]     = matrix_value * vector_value;
         }
 
         __syncthreads();
 
-        // Compute the thread's merge path segment
-        ItemOffsetPairT thread_segment[ITEMS_PER_THREAD];
+        // Traverse the thread's merge path segment
+        ValueT*         row_partials = reinterpret_cast<ValueT*>(temp_storage.merge_items);
+        ValueT          running_total = (threadIdx.x == 0) ? block_carry : 0.0;
         CoordinateT     thread_current_coord = thread_start_coord;
-        OffsetT         flags[ITEMS_PER_THREAD];
 
         #pragma unroll
         for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)

        // Output the tile's merge-path carry
        if (threadIdx.x == 0)
        {
            d_tile_carry_rows[blockIdx.x] = block_carry.offset;
            d_tile_carry_values[blockIdx.x] = block_carry.value;
        }
    }


     */
    __device__ __forceinline__ void ConsumeTile(
        CoordinateT&        tile_coord,
        ValueT&             block_carry)
    {
        __syncthreads();

