[CUDA][HIP] Report error about global/device function while device code has already been generated

[aywala]

#include <hip/hip_runtime.h>
#include <stdio.h>

#ifdef __HIP_DEVICE_COMPILE__
    #define HD __host__ __device__
#else
    #define HD
#endif

HD void foo(){
    printf("execute foo\n");
}

__global__ void kernel(){
    foo();
}

int main() {
    kernel<<<1,1>>>();
    hipDeviceSynchronize();
}
}

This code sample is error with clang when compiling for host target.

hip_macro.cpp:19:5: error: no matching function for call to 'foo'
    foo();
    ^~~
hip_macro.cpp:10:9: note: candidate function not viable: call to __host__ function from __global__ function
HD void foo(){

#include <stdio.h>

#ifdef __CUDA_ARCH__
    #define HD __host__ __device__
#else
    #define HD
#endif

HD void foo() {
    printf("execute foo\n");
}

__global__  void kernel(){
    foo(); 
}

int main() {
    kernel<<<1,1>>>();
    cudaDeviceSynchronize();
}

This code sample is ok with nvcc.

Why does clang parse device function once again when compiling for host target? And send an error for device function but the device code has already been generated. That does not make sense.

Real world problem:
Some header files have something like

#ifdef __CUDA_ARCH__
#define HD __host__ __device__
#else
    #define HD
#endif

When these header files are included in cuda source file, clang can not handle correctly. I know __CUDACC__ will be ok, but __CUDA_ARCH__ should be ok too. This is a problem caused by the design of compiler not caused by the language standard.

[aywala]

This is not a problem of LLVM. It is a problem of HIP.

[Artem-B]

It is not a problem with HIP/CUDA/clang or LLVM.

The error you get is intentional.

I can refer you to CUDA programming manual which explicitly states:
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cuda-arch

The type signature of the following entities shall not depend on whether __CUDA_ARCH__ is defined or not, or on a particular value of __CUDA_ARCH__

...
The compiler does not guarantee that a diagnostic will be generated for the unsupported uses of __CUDA_ARCH__ described above.

For clang, __host__ and __device__ attributes are part of the function signature and are expected to be the same for both the host and device compilations, and it is correct to diagnose an invalid call in your example. For NVCC it's an undefined behavior which it does not guarantee to diagnose, so your example is an UB which happens to work.

The source code in your example needs some changes to be compileable with clang.

[aywala]

I found this problem when compiling https://github.com/NVIDIA/caffe. I have a question, why doesn't skip device functions when compiling for host target? NVCC only passes host functions to host compiler. This difference may cause many portability problems.

[Artem-B]

With C++ pretty much any difference in source code will be observable during compilation. It's impossible to compile sufficiently complicated C++ code if host and device see different sources. I.e. whether a function is callable may determine which template variant gets instantiated. NVCC has to work around that.