[llvm-bugs] [Bug 47277] New: Do not throw an error for calling a __host__ function from a __device__ function during host compilation

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=47277

            Bug ID: 47277
           Summary: Do not throw an error for calling a __host__ function
                    from a __device__ function during host compilation
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: CUDA
          Assignee: unassignedclangbugs at nondot.org
          Reporter: drohr at jwdt.org
                CC: llvm-bugs at lists.llvm.org

Created attachment 23876
  --> https://bugs.llvm.org/attachment.cgi?id=23876&action=edit
testcase

We have a problem in our application using
#ifdef __CUDA_ARCH__
to use different logging backends for device and host.
Attached is a simplified example.

If __CUDA_ARCH__ is set, we just use printf(...) in our LOG() macro.
On the host when __CUDA_ARCH__ is not set, we use a logging library, which I
just called LogHost() in the testcase.
This works well during device compilation, but it fails during host
compilation.
What seems to happen is:
- The preprocessor parses the file with __CUDA_ARCH__ not set, so the LOG()
macro uses the LogHost() function.
- That makes the __device__ void test() function call LogHost(), but LogHost is
only a host function, not a __host__ __device__ function. (Which I want to be
like this, in the end LogHost shall come from an external library.)
- Now, even during host compilation, clang seems to check whether device
functions call host functions and bails out in that case.
My testcase fails with:
test.cu:12:3: error: no matching function for call to 'LogHost'
  LOG("test\n");
  ^~~~~~~~~~~~~
test.cu:4:15: note: candidate function not viable: call to __host__ function
from __device__ function
__host__ void LogHost(const char*) {}

I am wondering why this should fail. Basically, there is no need to check this
during host compilation at all. And funny enough, if I change
__device__ void test()
into
__host__ __device__ void test()
it actually works. This seems quite strange to me, since also in this case
test() is a __device__ function (in addition to being a host function) calling
LogHost. From this I would conclude that clang already has a mechanism to
ignore that the __device__ function test calls the __host__ function LogHost,
if test is also a __host__ function. Could this be extended to the case where
test() is only a __device__ function?

In addition, I want to mention that the testcase compiles without error with
nvcc, and I think it would make much sense if clang would remain compatible to
nvcc here.

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20200821/c4dc4a48/attachment.html>