         {
             OffsetT row_offset              = wd_row_end_offsets[tile_coord.x + row - 1];
             OffsetT row_end_offset          = wd_row_end_offsets[tile_coord.x + row];
-            int     local_row_end_offset    = row_end_offset - tile_coord.y - 1;
+            int     local_row_end_offset    = row_end_offset - (tile_coord.y - 1);
 
             ValueT value = temp_storage.values[local_row_end_offset];
             temp_storage.values[local_row_end_offset] = NAN_TOKEN;

            scan_op,
            tile_aggregate);

        __syncthreads();

        KeyValuePairT scan_items_out[ITEMS_PER_THREAD];

        // Exclusive scan in registers with prefix
        ThreadScanExclusive(scan_items, scan_items_out, scan_op, thread_partial);

        // Compact segment totals
        for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
        {
            if (scan_items[ITEM].key)
            {
                s_saved[scan_items_out[ITEM].key] += scan_items_out[ITEM].value;
            }
        }

        __syncthreads();

        // Store row totals
        #pragma unroll 1
        for (int row = threadIdx.x; row < tile_num_rows; row += BLOCK_THREADS)
        {
//                CubLog("Storing row %d: %f\n", tile_coord.x + row, value);
            spmv_params.d_vector_y[tile_coord.x + row] = s_saved[row];
        }
    }


};


        {
            scan_items[ITEM].key        = 0;
            scan_items[ITEM].value      = 0.0;

            int     local_nonzero_idx   = (threadIdx.x * ITEMS_PER_THREAD) + ITEM;
            ValueT  value               = temp_storage.values[local_nonzero_idx];

            if (local_nonzero_idx < tile_num_nonzeros)
            {
                if (value != value)
                    scan_items[ITEM].key = 1;
                else
                    scan_items[ITEM].value = value;
            }
        }

