[BUG]: NVRTC / NVHPC

[gevtushenko]

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

When using nvhpc 23.9 with nvrtc, there's an issue with numeric limits:

/opt/src/cccl/libcudacxx/include/cuda/std/detail/libcxx/include/limits(361): error: floating constant is out of range
      _LIBCUDACXX_INLINE_VISIBILITY static constexpr type denorm_min() noexcept {return __FLT_DENORM_MIN__;}
                                                                                        ^

/opt/src/cccl/libcudacxx/include/cuda/std/detail/libcxx/include/limits(413): error: floating constant is out of range
      _LIBCUDACXX_INLINE_VISIBILITY static constexpr type denorm_min() noexcept {return __DBL_DENORM_MIN__;}

How to Reproduce

#include <memory>
#include <iostream>
#include <nvrtc.h>
#include <cuda.h>

int main()
{
  nvrtcProgram prog{};

  const char* src =
    "#include <cuda/std/limits>                                                                  \n"
    "                                                                                            \n"
    "extern \"C\" __global__ void kernel(int *ptr, int *errors)                                  \n"
    "{                                                                                           \n"
    "}                                                                                           \n";

  const char* name = "test";

  if (NVRTC_SUCCESS != nvrtcCreateProgram(&prog, src, name, 0, nullptr, nullptr)) {
    std::cout << "error 1" << std::endl;
  }

  constexpr int num_includes = 3;
  const char* includes[num_includes] = {
    "-I/opt/src/cccl/libcudacxx/include",
    "-I/usr/local/cuda/include",
    "-arch=sm_52"
  };

  std::size_t log_size{};
  nvrtcResult compile_result = nvrtcCompileProgram(prog, num_includes, includes);

  if (NVRTC_SUCCESS != nvrtcGetProgramLogSize(prog, &log_size)) {
    std::cout << "error 2" << std::endl;
  }

  std::unique_ptr<char[]> log{ new char[log_size] };
  if (NVRTC_SUCCESS != nvrtcGetProgramLog(prog, log.get())) {
    std::cout << "error 3" << std::endl;
  }

  std::cout << "nvrtc log = " << log.get() << std::endl;
}

When compiled with nvhpc:

/opt/nvidia/hpc_sdk/Linux_x86_64/23.9/compilers/bin/nvc++ catch2_test_nvrtc.cu -I/opt/src/cccl/libcudacxx/include/ -O3 -lcuda -lnvrtc -L/usr/local/cuda/lib64

produces:

LD_LIBRARY_PATH=/usr/local/cuda/lib64:${LD_LIBRARY_PATH} ./a.out
nvrtc log = /opt/src/cccl/libcudacxx/include/cuda/std/detail/libcxx/include/limits(361): error: floating constant is out of range
      _LIBCUDACXX_INLINE_VISIBILITY static constexpr type denorm_min() noexcept {return __FLT_DENORM_MIN__;}
                                                                                        ^

/opt/src/cccl/libcudacxx/include/cuda/std/detail/libcxx/include/limits(413): error: floating constant is out of range
      _LIBCUDACXX_INLINE_VISIBILITY static constexpr type denorm_min() noexcept {return __DBL_DENORM_MIN__;}
                                                                                        ^

2 errors detected in the compilation of "test".

I tried using nvrtc that's shipped with nvhpc, seems to be the same issue there.

Expected behavior

NVRTC is able to compile the code.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

/usr/local/cuda/bin/nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Fri_Nov__3_17:16:49_PDT_2023
Cuda compilation tools, release 12.3, V12.3.103
Build cuda_12.3.r12.3/compiler.33492891_0

[gevtushenko]

Should be addressed by nvbug 4256559