For __CUDA_ARCH__ >= 350 report that merge_sort is not supported.
Port CUDA's radix sort to __device__ code. Still needs some __CUDA_ARCH__ guarding.