@miscco Why do we treat nvhpc differently? The objects don't live in device memory, if we don't make them__device__ with nvhpc

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

@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&cpo in device code, it returnnullptr
• __device__, the object is visible in the PTX.

Seehttps://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__

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 thinkconstexpr inline, without any__device__ is generally the best thing to do for NVHPC.