discrepancy between clang and nvcc regarding std::is_invocable_v

[yxsamliu]

https://godbolt.org/z/3cP5v4cbd

Basically, nvcc does not check availability by host/device attributes in std::is_invocable_v but clang does. This makes std::is_invocable_v returns false for a host function in a kernel for clang but not for nvcc.

Open this issue to discuss whether this should be treated as a clang bug.

@Artem-B

[Artem-B]

It looks like the same issue as what we've recently seen with concepts and template deduction guidelines.

In order to determine callability, we need to plumb through the caller context, and we currently don't (or assume global/host context, as it appears to be the case in the reproducer above.)

AFAICT, NVCC gets it wrong in the opposite direction and claims that everything is callable from everywhere:
https://godbolt.org/z/b4Wja18Eo

It may end up being the least bad choice here, as the user may instantiate a type based on std::is_invocable_v and the same type may be applicable to GPU and the host and we can't give two different answers within the same compilation for the same type we're creating with std::is_invocable_v and the given function object.

I guess it boils down to what we want std::is_invocable_v to mean? That the function object will really callable at runtime in the current sub-compilation? Or that it may potentially be callable in some subcompilation?

Ideally, within your reproducer we'd want std::is_invocable_v to be false for the host-only function objects within GPU-only functions and vice versa.

[ldionne]

FWIW we have experience with using strict availability attributes to influence how things SFINAE-away (which is what's happening here under the hood), and it has been the cause of a lot of issues. I am not well versed in CUDA, but from the outside I would strongly recommend that __host__ and __device__ does not influence what std::is_invocable returns.

[Artem-B]

That suggests that we may need to stick with the "true if it may be callable somewhere" interpretation, which is what NVCC appears to do.

On the positive side, it should be consistent for both host and GPU sub-compilations.
The downside is that it will be rather counter-intuitive to see "true" for a function object that is obviously not callable (e.g. an object with a GPU-only operator() used in an assert() within a host-only function).

We may need to grow some sort of target-specific extension std::is_callable_on_host and std::is_callable_on_gpu which could be used to check callability more specifically, where it matters.

[ldionne]

I think something like this is quite desirable as we work towards running e.g. PSTL algorithms on GPUs (see for example #66968). However, it would probably operate at a different level than something like std::is_invocable, which clearly operates at the C++ "front-end" level in the sense that it checks whether some expression is syntactically valid.

[yxsamliu]

The reduced test case is here https://godbolt.org/z/xfYs16M35

clang tries to substitute __is_invocable<DeviceAdd,int, int>

for that, it tries to substitute _S_test<DeviceAdd,int, int>

for that it tries to substitute std::declval<DeviceAdd>()(std::declval<int, int>())

which ends up as a call DeviceAdd(0,0) with the current function _S_test. IsAllowedCUDACall thinks that it is a call of device function in a host function. This call is discarded, which causes substitution failure.

adding __device__ __host__ to _S_test makes the test pass.

In a sense, clang's behavior is reasonable. When we do substitutions for decltype(std::declval<_F>()(std::declval<_Args>()...)) _S_test(int);, we may want to limit _F to functions that are available to _S_test, because we could have defined functions of the same name but different signatures for device and host. It is not so useful to do substitution with functions not available, since it will just result in inconsistencies after the type is really instantiated.

However, we are now facing the issue that <type_traits> are written in C++ and all function templates are assumed to be host by default. constexpr can make them implicitly host device but in this case we have a function that is not constexpr.

One workaround for this issue is to add a wrapper header for <type_traits> where we wrap the original header with #pragma clang force_cuda_host_device begin and #pragma clang force_cuda_host_device end.

Another solution is to disable availability check during template substitution, but this could cause consistency issues and regressions.