             // Smem needed for block exchange
             typename BlockExchangeT::TempStorage exchange;
 
-        };
-
-        union 
-        {
             // Smem needed for block-wide reduction
             typename BlockReduceT::TempStorage reduce;
 

            scan_item.key = row_indices[0];
            scan_item.value = 0.0;
        }

        // Scan downsweep and scatter
        ValueT* s_partials = &temp_storage.merge_items[0].nonzero;

            row_indices[ITEM] = thread_current_coord.x;
        }

        __syncthreads();

        // Block-wide reduce-value-by-segment
        KeyValuePairT       tile_carry;
        ReduceBySegmentOpT  scan_op;
