[CUDA] Code example states that it should fail, but it does not

[carlosgalvezp]

I'm testing the code example provided here, but it does not appear to fail anymore, contrary to the documentation. Should it be updated?

https://godbolt.org/z/5bs8h1vWK

Note that I had to modify it slightly, the #endif probably was meant to be placed at the end of the struct.

[Artem-B]

The example is... a bit out of date. A lot has changed in the way clang handles CUDA code since it was written.
Specifically, originally, the documentation implied that the failure would be reported during the host compilation when it would see the S.foo() call in the device function, but no suitable function overload would be found.

#ifdef __CUDA_ARCH__
struct S {
  __device__ void foo() { /* use device_only */ }
  int device_only;
};
#else
struct S {
  __host__ void foo() { /* use host_only */ }
  double host_only;
};
#endif

__device__ void test() {
  S s;
  // clang generates an error here, because during host compilation, we
  // have ifdef'ed away the __device__ overload of S::foo().  The __device__
  // overload must be present *even during host compilation*.
  s.foo();
}

That error still happens. However, for some of the the "wrong side" errors (e.g errors in the GPU-only code during host compilation), clang defers the diagnostics and only emits them if we ever attempt to generate code for such a function. In this case, clang would find a host variant of S::foo, and would issue a deferred diagnostics, reporting that one can't call a host function from the GPU code. However, because clang never generates any code for the GPU functions during host compilation, you never get to see that error. If you change the kind of error to be something else (i.e. add an argument to the host version of the function), the error becomes unambiguous and will be reported during the host compilation.

That said, the documentation is still correct that such code is a bad practice, as it breaks compiler assumptions that it sees identical source code during the hsot and the GPU compilation. While you may be able to get by in some cases, with sufficiently complex code it will likely buy one more trouble than it's worth. Think of it as an ODR violation, only for host/GPU boundary, instead of the 'classic' case which affects host-side object files.

[carlosgalvezp]

Understood, thank you for the detailed explanation!

[carlosgalvezp]

By the way, partially related, is the following code OK on Clang? The example above talks about classes/objects with different members only, so I don't know if it applies to plain functions.

__host__ void foo_cpu() { printf("CPU\n"); }
__device__ void foo_gpu() { printf("GPU\n"); }

__host__ __device__ void foo()
{
#ifdef __CUDA_ARCH__
  foo_gpu();
#else
  foo_cpu();
#endif
}

The goal is essentially to have one common function foo that works on H, D and HD contexts, and also is compatible with NVCC. We currently have the complex conditional includes presented in the example but wonder if it could be simplified in this way.

[Artem-B]

It should work, but without specific details it's hard to tell if that will work well enough. The host/gpu code divergence will certainly be observable, but I can't tell whether it would matter in your case.

[carlosgalvezp]

Yes I meant more from an ODR/UB perspective. One concrete example is implementing a custom assertFail function that calls __trap(); in device code, and std::abort(); in host code. Or a log2 wrapper that calls __log2f in device mode but std::log2 in host mode.