| Issue |
117831
|
| Summary |
[CUDA][HIP] Report error about device function while device code has already been generated
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
aywala
|
```cpp
#ifdef __CUDA_ARCH__
#define HD __attribute__((host)) __attribute__((device))
#else
#define HD
#endif
HD void foo(){
}
__attribute__((device)) void kernel(){
foo();
}
```
This code sample is ok with nvcc, but error with clang.
https://godbolt.org/z/nnf6jrGoj
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.
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
