[FEA]: Support FP8 in Block Radix Sort

[gevtushenko]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

CUB

Is your feature request related to a problem? Please describe.

CUB radix sort doesn't support FP8. The code below doesn't compile:

__global__ void kernel() {
  using block_radix_sort_t = cub::BlockRadixSort<__nv_fp8_e4m3, 32, 1>;
  using temp_storage_t = block_radix_sort_t::TempStorage;

  __shared__ temp_storage_t temp_storage;

  __nv_fp8_e4m3 fp8[1]{__nv_fp8_e4m3(static_cast<float>(threadIdx.x))};

  block_radix_sort_t(temp_storage).SortDescending(fp8);
}

Describe the solution you'd like

FP8 support is documented and tested.

Describe alternatives you've considered

No response

Additional context

No response