[FEA] Consider disabling --expt-relaxed-constexpr

[jrhemstad]

Is your feature request related to a problem? Please describe.

--expt-relaxed-constexpr is a convenient way to reuse existing constexpr host code, e.g., things like std::max.

However, it can lead to some pretty surprising behavior. Consider:

constexpr int bar(int j){
    if(j<0){
        throw;
    }
    return 42;
}
__global__ void kernel(int * i){
    *i = bar(-1);
}

https://godbolt.org/z/frb8c6cd7

One might expect this to fail to compile as throw is not valid in device code. However, not only does it happily compile, but it just stores the value 42.

This example looks pretty harmless:

int foo(int i){
    return i * 2;
}
constexpr int bar(int j){
    if(j<0){
        return foo(j);
    }
    return 42;
}
__global__ void kernel( int * i){
    *i = bar(-1);
}

But this too results in an ill-formed program without a diagnostic.

https://godbolt.org/z/aTzGaMrGd

Describe the solution you'd like

We should think pretty hard about if we want to risk such egregious undefined behavior in libcudf.

As such, we may want to consider moving towards disabling --expt-relaxed-constexpr. At the very least, we should be preferring CUDA_HOST_DEVICE_CALLABLE whenever possible (for functions that need be called from both host and device).

Additional Context

The only place it is 100% safe to use a constexpr function in device code with --expt-relaxed-constexpr is when used in a context that requires constant evaluation. Then it will fail to compile if the constexpr function contains things that would result in an ill-formed program: https://godbolt.org/z/47qfnPnc9

[harrism]

@karthikeyann please be aware of this with respect to #7713

[harrism]

At the very least, we should be preferring CUDA_HOST_DEVICE_CALLABLE whenever possible.

No, we should prefer CUDA_DEVICE_CALLABLE or no annotation whenever possible. We should use CUDA_HOST_DEVICE_CALLABLE only for functions that are required to be called from both host and device.

[jrhemstad]

only for functions that are required to be called from both host and device.

This was implied. The whole point of --expt-relaxed-constexpr is to call existing host constexpr functions from device, which implies it can be/is called from both host and device.

[harrism]

I want to be very explicit because it is not obvious to everyone (a fact that is evident in our code).

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[harrism]

Still valid

[PointKernel]

This seems to be unblocked with the recent CCCL update.

The changes in NVIDIA cuCollections PR #595 show that cuco can now build and run properly without relaxed constexprs.

[vyasr]

NVIDIA/cuCollections#595 and rapidsai/cuspatial#1494 inspired me to look into this again. It's slow going (but low effort) to investigate, so I'll report back when I get compilation far enough to make a determination on whether removing this setting will be possible in cudf. I've run into one very odd case that seemed like a compiler bug, but other than that it seems largely like a lot of tedious but not difficult work.