Repository: pytorch
Issue: Torch fails to build with cub from CCCL v2.3.2 or newer, due to -D__CUDA_NO_HALF_OPERATORS__
Labels: module: build,module: cuda,triaged

### 🐛 Describe the bug

From https://github.com/NVIDIA/cccl/issues/1766:

We've encountered another interesting conundrum, trying to build torch.

Torch builds with a bunch of __CUDA_NO_{HALF/BFLOAT16}... defines https://github.com/pytorch/pytorch/blob/faf0015052ee37db718bc5efa6673e0c25be1e8d/cmake/Dependencies.cmake#L1609

It works well enough, at least up until the CUB version included with cuda-12.4 (v2.3.0?).
However, attempting to build it with cub v2.3.2 runs into an issue. Apparently the newer version of CUB needs some of the fp16/bf16 overloads disabled by the torch build flags.

```
In file included from third_party/py/torch/aten/src/ATen/cuda/cub.cu:2:
In file included from third_party/py/torch/aten/src/ATen/cuda/cub.cuh:15:
In file included from third_party/gpus/cccl/cub/cub/cub.cuh:64:
In file included from third_party/gpus/cccl/cub/cub/device/device_histogram.cuh:52:
third_party/gpus/cccl/cub/cub/device/dispatch/dispatch_histogram.cuh:779:35: error: invalid operands to binary expression ('__half' and '__half')
  779 |                    (return sample >= min_level && sample < max_level;),
      |                            ~~~~~~ ^  ~~~~~~~~~
...
blaze-out/k8-opt-cuda12/bin/third_party/gpus/cuda/_virtual_includes/_stage/third_party/gpus/cuda/include/cuda_bf16.hpp:310:52: note: candidate function not viable: no known conversion from '__half' to 'const __nv_bfloat16' for 1st argument
  310 | __CUDA_HOSTDEVICE__ __CUDA_BF16_FORCEINLINE__ bool operator>=(const __nv_bfloat16 &lh, const __nv_bfloat16 &rh) { return __hge(lh, rh); }
      |                                                    ^          ~~~~~~~~~~~~~~~~~~~~~~~
blaze-out/k8-opt-cuda12/bin/third_party/gpus/cuda/_virtual_includes/_stage/third_party/gpus/cuda/include/cuda_bf16.hpp:367:52: note: candidate function not viable: no known conversion from '__half' to 'const __nv_bfloat162' for 1st argument
  367 | __CUDA_HOSTDEVICE__ __CUDA_BF16_FORCEINLINE__ bool operator>=(const __nv_bfloat162 &lh, const __nv_bfloat162 &rh) { return __hbge2(lh, rh); }
      |                                                    ^          ~~~~~~~~~~~~~~~~~~~~~~~~
```
Note that the ambiguous overload here is a red herring. The real problem is that the correct overloads were not found because they were disabled.

If CUDA-provided overloads are enabled, then torch build fails:
https://github.com/pytorch/pytorch/actions/runs/8991128640/job/24733299645?pr=125707

```
[ 89%] Building CUDA object caffe2/CMakeFiles/torch_cuda.dir/__/aten/src/ATen/native/cuda/ActivationHardshrinkKernel.cu.o
/var/lib/jenkins/workspace/aten/src/ATen/native/cuda/ActivationHardshrinkKernel.cu(23): error: more than one operator ">=" matches these operands:
            built-in operator "arithmetic >= arithmetic"
            function "operator>=(const __half &, const __half &)" (declared at line 311 of /usr/local/cuda/include/cuda_fp16.hpp)
            operand types are: scalar_t >= c10::Half
```

So, we somehow need to have the cake (have CUDA-provided overloads), and eat it (have those overloads disabled), to keep both cub and torch happy.

Ideally it would be great if CUDA headers would stash the overloads in some known namespace, which would allow cub to pull them in into its own namespace, but it's not an option, those overloads are just preprocessed away.

Edit: Not sure if it's possible, though, as the __CUDA_NO_* macros also disable member functions, not just free-standing ones.

Fixing torch to work w/o disabling CUDA-provided overloads for `__half` would be nice. 

@voznesenskym @malfet

### Versions

The issue was encountered with v2.3.0a0 36c1cc962aaef854d2388a5ecfde230d40bcc1d6