Author: tra Date: Fri Mar 23 12:49:03 2018 New Revision: 328362 URL: http://llvm.org/viewvc/llvm-project?rev=328362&view=rev Log: [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 Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=328362&r1=328361&r2=328362&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Fri Mar 23 12:49:03 2018 @@ -790,9 +790,12 @@ bool Sema::CheckCUDACall(SourceLocation // 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 Modified: cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu?rev=328362&r1=328361&r2=328362&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu (original) +++ cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu Fri Mar 23 12:49:03 2018 @@ -83,3 +83,10 @@ template <typename T> __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>>>(); } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits