         };
 
         // The first partial reduction scattered by this thread block
-        SizeT first_partial;
+        PartialReduction first_partial;
     };
 
     //---------------------------------------------------------------------

    // Parameterized BlockScan type for reduce-value-by-segment scan
    typedef BlockScan<PartialReduction, BLOCK_THREADS, SCAN_ALGORITHM> BlockScan;

    // A cache item that can either be a value or a segment-ID
    union CacheItem
    {
        SizeT segment_idx;
        Value value;            // TODO: fix for non-default-constructible value types
    };

    // Shared memory type for this threadblock
    struct TempStorage
    {

                    partial_reductions[ITEM].partial = d_values[thread_value_idx];
                    thread_value_idx++;

                    if ((valid_value = (thread_value_idx < next_thread_value_idx)))
                        d_value_offsets[thread_value_idx];
                }
            }

            // Update tile starting indices to the last thread's end indices
            block_segment_idx = temp_storage.segment_idx[BLOCK_THREADS];
            block_value_idx = temp_storage.value_idx[BLOCK_THREADS];

            // 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.
