             block_value_idx :                               // First thread starts at the block's start
             temp_storage.thread_value_idx[threadIdx.x];     // Other threads start at their predecessor's end
 
-        // Barrier for smem reuse
-        __syncthreads();
-
         // Check if first segment end-offset is in range
         bool valid_segment = (thread_segment_idx < next_thread_segment_idx);
 

            next_tile_segment_idx   = temp_storage.thread_segment_idx[BLOCK_THREADS];
            next_tile_value_idx     = temp_storage.thread_value_idx[BLOCK_THREADS];

            if (block_segment_idx == next_tile_segment_idx)
            {
                // There are no segment end-offsets in this tile.  Perform a
                // simple block-wide reduction and accumulate the result into
                // the running total.
                SingleSegmentTile(next_tile_value_idx, block_value_idx, block_segment_idx);
            }
//          else if (block_value_idx == next_tile_value_idx)
//            {
//                // There are no values in this tile (only empty segments).
//                EmptySegmentsTile(next_tile_segment_idx, block_segment_idx);
//            }
            else
            {
                // Merge the tile's segment and value indices
                MultiSegmentTile(block_segment_idx, block_value_idx, next_thread_segment_idx, next_thread_value_idx);

            else if (valid_value)
            {
                // Consume this value index
                value_offsets[ITEM] = thread_value_idx;
                thread_value_idx++;

                // Read next value offset (if valid)
                valid_value = (thread_value_idx < next_thread_value_idx);
            }
        }

        // Load values
        Value values[ITEMS_PER_THREAD];

        #pragma unroll
        for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
        {
            values[ITEM] = (value_offsets[ITEM] == -1) ?
                identity :
                d_values[value_offsets[ITEM]];
        }

        // Reduce within thread segments
        #pragma unroll
        for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
        {
            if (segment_ids[ITEM] != -1)
            {
                // Consume this segment index
                d_output[segment_ids[ITEM]] = pairs[1].value;

                if (pairs[0].key == segment_ids[ITEM])
                    pairs[0].value = pairs[1].value;

                pairs[1].key    = segment_ids[ITEM];
                pairs[1].value  = identity;
            }

            pairs[1].value = reduction_op(pairs[1].value, values[ITEM]);
        }

        // Barrier for smem reuse
        __syncthreads();

        // Use prefix scan to reduce values by segment-id.  The segment-reductions end up in items flagged as segment-tails.
        KeyValuePair block_aggregate;
        BlockScan(temp_storage.scan).InclusiveScan(
            pairs,                          // Scan input
            pairs,                          // Scan output
            scan_op,                        // Scan operator
            block_aggregate,                // Block-wide total (unused)
            prefix_op);                     // Prefix operator for seeding the block-wide scan with the running total


/*
        // Check if first segment end-offset is in range
        bool valid_segment = (thread_segment_idx < next_thread_segment_idx);

        // Check if first value offset is in range
        bool valid_value = (thread_value_idx < next_thread_value_idx);

        // Load first segment end-offset
        Offset segment_end_offset = (valid_segment) ?
            d_segment_end_offsets[thread_segment_idx] :
            num_values;                                                     // Out of range (the last segment end-offset is one-past the last value offset)

        // Load first value offset
        Offset value_offset = (valid_value) ?
            d_value_offsets[thread_value_idx] :
            num_values;                                                     // Out of range (one-past the last value offset)

        // Assemble segment-demarcating tail flags and partial reduction tuples
        TailFlag        tail_flags[ITEMS_PER_THREAD];
        KeyValuePair    partial_reductions[ITEMS_PER_THREAD];

        #pragma unroll
        for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
        {
            // Default tuple and flag values
            partial_reductions[ITEM].key    = thread_segment_idx;
            partial_reductions[ITEM].value  = identity;
            tail_flags[ITEM]                = 0;

            // Whether or not we slide (a) right along the segment path or (b) down the value path
            if (valid_segment && (!valid_value || (segment_end_offset <= value_offset)))
            {
                // Consume this segment index

                // Set tail flag noting the end of the segment
                tail_flags[ITEM] = 1;

                // Increment segment index
                thread_segment_idx++;

                // Read next segment end-offset (if valid)
                if ((valid_segment = (thread_segment_idx < next_thread_segment_idx)))
                    segment_end_offset = d_segment_end_offsets[thread_segment_idx];
            }
            else if (valid_value)
            {
                // Consume this value index

                // Update the tuple's value with the value at this index.
                partial_reductions[ITEM].value = d_values[value_offset];

                // Increment value index
                thread_value_idx++;
