<html>
<head>
<base href="https://bugs.llvm.org/">
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW - Do not throw an error for calling a __host__ function from a __device__ function during host compilation"
href="https://bugs.llvm.org/show_bug.cgi?id=47277">47277</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>Do not throw an error for calling a __host__ function from a __device__ function during host compilation
</td>
</tr>
<tr>
<th>Product</th>
<td>clang
</td>
</tr>
<tr>
<th>Version</th>
<td>trunk
</td>
</tr>
<tr>
<th>Hardware</th>
<td>PC
</td>
</tr>
<tr>
<th>OS</th>
<td>Linux
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>enhancement
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>CUDA
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedclangbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>drohr@jwdt.org
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr></table>
<p>
<div>
<pre>Created <span class=""><a href="attachment.cgi?id=23876" name="attach_23876" title="testcase">attachment 23876</a> <a href="attachment.cgi?id=23876&action=edit" title="testcase">[details]</a></span>
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.</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are on the CC list for the bug.</li>
</ul>
</body>
</html>