Make _CCCL_GLOBAL_CONSTANT inline when RDC is enabled#7682
Make _CCCL_GLOBAL_CONSTANT inline when RDC is enabled#7682davebayer wants to merge 1 commit intoNVIDIA:mainfrom
_CCCL_GLOBAL_CONSTANT inline when RDC is enabled#7682Conversation
| #if _CCCL_DEVICE_COMPILATION() && !_CCCL_CUDA_COMPILER(NVHPC) | ||
| # define _CCCL_GLOBAL_CONSTANT _CCCL_DEVICE constexpr | ||
| # if _CCCL_HAS_RDC() | ||
| # define _CCCL_GLOBAL_CONSTANT _CCCL_DEVICE inline constexpr | ||
| # else // ^^^ _CCCL_HAS_RDC() ^^^ / vvv !_CCCL_HAS_RDC() vvv | ||
| # define _CCCL_GLOBAL_CONSTANT _CCCL_DEVICE constexpr | ||
| # endif // ^^^ !_CCCL_HAS_RDC() ^^^ | ||
| #else // ^^^ _CCCL_DEVICE_COMPILATION() && !_CCCL_CUDA_COMPILER(NVHPC) ^^^ / | ||
| // vvv !_CCCL_DEVICE_COMPILATION() || _CCCL_CUDA_COMPILER(NVHPC) vvv | ||
| # define _CCCL_GLOBAL_CONSTANT inline constexpr |
There was a problem hiding this comment.
@miscco Why do we treat nvhpc differently? The objects don't live in device memory, if we don't make them __device__ with nvhpc
There was a problem hiding this comment.
The reason is that NVHPC has a single compilation pass, so it will see that the global object is used on device and instantiate it there
There was a problem hiding this comment.
@dkolsen-pgi could you tell us what's exactly happening under the hood?
What I observe is that if I set the execution space of a CPO to:
- nothing it seems that the object is not part of the generated PTX and if I do
&cpoin device code, it returnnullptr __device__, the object is visible in the PTX.
See https://ce.nvidia.com/z/7d77Kj. I'm not sure whether this is a problem, but it's not consistent with what other compilers produce. Because with all other compilers, we basically create the object twice - 1x on __host__ and 1x on __device__
There was a problem hiding this comment.
NVHPC's handling of constexpr variables has bugs. I can't fully explain what you are seeing.
The basic cases work well when the constexpr variable doesn't have any __host__ or __device__ annotations:
- A constexpr scalar variable whose address is never needed will be constant folded in both host and device code. An actual object won't be generated in either host or device because it isn't needed.
- A CPO or other stateless function object can be called just fine in both host and device code. The object isn't created correctly in device code, but it doesn't matter because the object is never actually accessed.
The thing that doesn't work is when the constexpr variable has a meaningful value and that value is accessed via a pointer to the object. Using it in host code is fine, but using it in device code usually won't compile, with nvc++ complaining about accessing a global variable from device code. If the constexpr variable is explicitly marked __device__, it will work correctly in device code, but wrong results might happen in host code.
We haven't fixed these problems because it's not easy to do and the existing behavior works well enough for the code that nvc++ needs to compile.
I think constexpr inline, without any __device__ is generally the best thing to do for NVHPC.
🥳 CI Workflow Results🟩 Finished in 3h 54m: Pass: 100%/95 | Total: 4d 12h | Max: 3h 54m | Hits: 16%/249655See results here. |
|
question: Have we convinced ourselves this isn't going to cause insidious ODR issues? |
4 participants
When compiling libcu++, it puts significant amount of objects (mostly CPOs) with static linkage to the global memory. This PR makes the symbols
inline(.weak) when RDC is enabled.Demo: https://godbolt.org/z/8ofPq8Wo1