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

        // 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;
        }
    }


            if (accumulate)
            {
                // Move down (accumulate)
                running_total += tile_nonzeros[thread_current_coord.y];
                ++thread_current_coord.y;
            }
            else
            {
                // Move right (reset)
                row_partials[thread_current_coord.x] = running_total;
                running_total = 0.0;
                ++thread_current_coord.x;
            }
        }

        __syncthreads(); // Perf sync

        // Block-wide reduce-value-by-key
        int                 thread_num_rows = thread_end_coord.x - thread_start_coord.x;
        ItemOffsetPairT     scan_item(thread_num_rows, running_total);
        ItemOffsetPairT     tile_aggregate;
        ItemOffsetPairT     scan_identity(0.0, 0);
        ReduceBySegmentOpT  scan_op;
        BlockScanT(temp_storage.scan).ExclusiveScan(scan_item, scan_item, scan_identity, scan_op, tile_aggregate);

        // Update block carry
        block_carry = tile_aggregate.value;

        // Inter-thread fix-up
        if (thread_num_rows > 0)
            row_partials[thread_start_coord.x] += scan_item.value;      

        __syncthreads();

        // Scatter
        for (int item = threadIdx.x; item < tile_num_rows; item += BLOCK_THREADS)
        {
            d_vector_y[tile_coord.x + item] = row_partials[item];

        }

        // Update the block coordinate
        tile_coord = tile_end_coord;
    }

