`DeviceSegmentedRadixSortKernel` uses `volatile` temporary storage

[bernhardmgruber]

DeviceSegmentedRadixSortKernel uses volatile temporary storage:

cccl/cub/cub/device/dispatch/kernels/kernel_segmented_radix_sort.cuh

Lines 142 to 151 in a0a6bfd

	__shared__ union
	{
	typename BlockUpsweepT::TempStorage upsweep;
	typename BlockDownsweepT::TempStorage downsweep;
	struct
	{
	volatile SegmentSizeT reverse_counts_in[RADIX_DIGITS];
	volatile SegmentSizeT reverse_counts_out[RADIX_DIGITS];
	typename DigitScanT::TempStorage scan;
	};

This is at best a code smell, but probably either buggy and inefficient (disabled optimizations). The use of volatile should be replaced by proper loads and stores, atomic or regular, as needed and in accordance with the memory model.