From b2926d8649d8def7f527f6afd945985a445c629a Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Fri, 23 Mar 2018 19:49:03 +0000 Subject: [PATCH] [CUDA] Fixed false error reporting in case of calling H->G->HD->D. Launching a kernel from the host code does not generate code for the kernel itself. This fixes an issue with clang erroneously reporting an error for a HD->D call from within the kernel. Differential Revision: https://reviews.llvm.org/D44837 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@328362 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Sema/SemaCUDA.cpp | 9 ++++++--- test/SemaCUDA/call-device-fn-from-host.cu | 7 +++++++ 2 files changed, 13 insertions(+), 3 deletions(-) diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp index cac5f68227..ccd93fa48b 100644 --- a/lib/Sema/SemaCUDA.cpp +++ b/lib/Sema/SemaCUDA.cpp @@ -790,9 +790,12 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) { // If the caller is known-emitted, mark the callee as known-emitted. // Otherwise, mark the call in our call graph so we can traverse it later. bool CallerKnownEmitted = IsKnownEmitted(*this, Caller); - if (CallerKnownEmitted) - MarkKnownEmitted(*this, Caller, Callee, Loc); - else { + if (CallerKnownEmitted) { + // Host-side references to a __global__ function refer to the stub, so the + // function itself is never emitted and therefore should not be marked. + if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global) + MarkKnownEmitted(*this, Caller, Callee, Loc); + } else { // If we have // host fn calls kernel fn calls host+device, // the HD function does not get instantiated on the host. We model this by diff --git a/test/SemaCUDA/call-device-fn-from-host.cu b/test/SemaCUDA/call-device-fn-from-host.cu index 3ea013f3a7..26215d581d 100644 --- a/test/SemaCUDA/call-device-fn-from-host.cu +++ b/test/SemaCUDA/call-device-fn-from-host.cu @@ -83,3 +83,10 @@ template __host__ __device__ void fn_ptr_template() { auto* ptr = &device_fn; // Not an error because the template isn't instantiated. } + +// Launching a kernel from a host function does not result in code generation +// for it, so calling HD function which calls a D function should not trigger +// errors. +static __host__ __device__ void hd_func() { device_fn(); } +__global__ void kernel() { hd_func(); } +void host_func(void) { kernel<<<1, 1>>>(); } -- 2.40.0