[Issue]: [HIP] Report error about device/global function while device-side code has already been generated

[aywala]

Problem Description

#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 hipcc. It report an error when compiling for host.
https://godbolt.org/z/dGG5zEhEc

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 hipcc 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, hipcc can not handle correctly. I know __CUDACC__ will be ok, but __CUDA_ARCH__\ __HIP_DEVICE_COMPILE__ should be ok too. This is a problem caused by the design of compiler not caused by the language standard.

Operating System

Ubuntu22.04

CPU

R5 7500F

GPU

gfx906

ROCm Version

ROCm 6.1.0

ROCm Component

No response

Steps to Reproduce

No response

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

[shiltian]

There is a reason __host__ exists, and that is exactly it. I don't think this is an issue. The code in the description doesn't conform with the programming model IMO.