[Clang] [CUDA] unsupported declare CUDA host device function and define host/device separate

[Acthink-Yang]

https://godbolt.org/z/K44e7esn1

__host__ __device__ void func();  // error

// __device__ void func(); // ok
// __host__ void func(); // ok


__host__ void func() {}

__device__ void func(void) {}

error: host function 'func' cannot overload host device function 'func'

I found "HD functions cannot be overloaded by H or D functions with the same signature" in this doc:

https://llvm.org/docs/CompileCudaWithLLVM.html#overloading-based-on-host-and-device-attributes

but I want to known declear HD, than separate define H/D is OK?

[Acthink-Yang]

I found the check in clang/lib/Sema/SemaCUDA.cpp with this comment:

>     // Don't allow HD and global functions to overload other functions with the
>     // same signature.  We allow overloading based on CUDA attributes so that
>     // functions can have different implementations on the host and device, but
>     // HD/global functions "exist" in some sense on both the host and device, so
>     // should have the same implementation on both sides.

I've found that hosts and devices often have different implementations:

#define __CUDA_HOSTDEVICE_FP16_DECL__ static __host__ __device__ __inline__
__CUDA_HOSTDEVICE_FP16_DECL__ __half __float2half(const float a)
{
    __half val;
#if defined(__CUDA_ARCH__)
    asm("{  cvt.rn.f16.f32 %0, %1;}\n" : "=h"(__HALF_TO_US(val)) : "f"(a));
#else
    __half_raw r;
    unsigned int sign = 0U;
    unsigned int remainder = 0U;
    r.x = __internal_float2half(a, sign, remainder);
    if ((remainder > 0x80000000U) || ((remainder == 0x80000000U) && ((r.x & 0x1U) != 0U))) {
        r.x++;
    }
    val = r;
#endif
    return val;
}

[llvmbot]

@llvm/issue-subscribers-clang-frontend

Author: QiYue (QiYueFeiXue)

https://godbolt.org/z/K44e7esn1

__host__ __device__ void func();  // error

// __device__ void func(); // ok
// __host__ void func(); // ok


__host__ void func() {}

__device__ void func(void) {}

error: host function 'func' cannot overload host device function 'func'

I found "HD functions cannot be overloaded by H or D functions with the same signature" in this doc:

https://llvm.org/docs/CompileCudaWithLLVM.html#overloading-based-on-host-and-device-attributes

but I want to known declear HD, than separate define H/D is OK?

[Artem-B]

This behavior is by design and is not going to change.

I've found that hosts and devices often have different implementations:

That's an NVCC-ism. NVCC does not have a notion of target attribute overloads, so preprocessing inside of HD functions is the only way they can handle different H/D implementations. This was the major reason we went with the overloading approach. We still support NVCC's way of doing things, but it does not mesh with overloads. It's an either-or choice:

• either a single HD function with ifdefs inside. (the only way to do it with NVCC)
• or separate H and D overloads.

To a degree, you can mix HD with H/D functions is by relying on existing C++ overloading mechanisms. E.g. by adding a dummy function argument to H and D functions.