 
     if (threadIdx.x == 0)
     {
-        if (gridDim.x == 1)
-        {
-            // Update the segment reduction for the region's last segment
-            d_output[last_tuple.key] = last_tuple.value;
-        }
-        else
-        {
+        if (gridDim.x > 1)
+        {
+            // Special case where the first segment written and the carry-out are for the same segment
+            if (first_tuple.key == last_tuple.key)
+            {
+                first_tuple.value = identity;
+            }
+
             // Write the first and last partial products from this thread block so
             // that they can be subsequently "fixed up" in the next kernel.
             d_tuple_partials[blockIdx.x * 2]          = first_tuple;

                        if ((valid_segment = (thread_segment_idx < next_thread_segment_idx)))
                            segment_end_offset = d_segment_end_offsets[thread_segment_idx];
                    }
                    else if (valid_value && !prefer_segment)
                    {
                        // Consume this value index



        segment_offsets.push_back(consumed);
    }
}


