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

        scan_item.value = running_total;
        scan_item.key = thread_current_coord.x;

        BlockScanT(temp_storage.scan).ExclusiveScan(scan_item, scan_item, scan_op, tile_carry);
        if (threadIdx.x == 0)
        {
            scan_item.key = row_indices[0];

        CoordinateT tile_end_coord       = d_tile_coordinates[tile_idx + 1];

        // Consume multi-segment tile
        KeyValuePairT tile_carry = ConsumeTile2(
            tile_idx,
            tile_start_coord,
            tile_end_coord,
