[BUG]: isfinite breaks NVRTC since 2.8.0

[leofang]

Is this a duplicate?

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

Type of Bug

Compile-time Error

Component

libcu++

Describe the bug

Based on internal discussion, __builtin_isfinite does not exist in NVRTC 12.0-12.2.

How to Reproduce

1. Create a conda env

conda create -n my_env -y python=3.12 cuda-core cuda-version=12.0
conda activate my_env

2. Run the follow Python script

from cuda.core.experimental import Program, Device

dev = Device(0)
dev.set_current()

code2 = r"""
#include <cuda/std/cmath>


__global__ void my_kernel(double* s) {
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  double y = s[1]; 
  if (tid == 0) {
     auto k = isfinite(y);
     if (k) {
        s[0] += 1;
     }
  }
}
"""

prog = Program(code2, code_type="c++")
mod = prog.compile(target_type="cubin", options=("-std=c++17", "-I/path/to/local/cccl/libcudacxx/include/"))

3. Repeat steps 1 & 2 for 12.1-12.8 via

conda install -y cuda-version=12.X  # X=1,2,..., no 7

Expected behavior

Everything should work on Linux/Windows.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[leofang]

(Fixed the issue title; I meant to say 2.7.0 works fine but not 2.8.0 or main)

[leofang]

Notes:

1. Try and make nvrtc on windows pass #3623 was an attempt to fix this, but it did not touch the root cause and did not work
2. I do not understand why the NVRTC CI tests did not catch this. Seems like a gap to me.
3. this seems to work for NVRTC on both Linux and Windows, and for all CUDA 12.x:

diff --git a/libcudacxx/include/cuda/std/__cccl/builtin.h b/libcudacxx/include/cuda/std/__cccl/builtin.h
index 3a5fda2f0..e3683fc58 100644
--- a/libcudacxx/include/cuda/std/__cccl/builtin.h
+++ b/libcudacxx/include/cuda/std/__cccl/builtin.h
@@ -244,9 +244,9 @@
 #endif // _CCCL_CHECK_BUILTIN(isfinite)
 
 // Below 11.7 nvcc treats the builtin as a host only function
-#if _CCCL_CUDACC_BELOW(11, 7)
+#if _CCCL_CUDACC_BELOW(12, 3)
 #  undef _CCCL_BUILTIN_ISFINITE
-#endif // _CCCL_CUDACC_BELOW(11, 7)
+#endif // _CCCL_CUDACC_BELOW(12, 3)
 
 #if _CCCL_CHECK_BUILTIN(builtin_isinf) || _CCCL_COMPILER(GCC)
 #  define _CCCL_BUILTIN_ISINF(...) __builtin_isinf(__VA_ARGS__)