[CUDA] Fixed false error reporting in case of calling H->G->HD->D.

[tra]

Summary

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.

clang/lib/Sema/SemaCUDA.cpp

bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   // the caller is a global initializer?
   FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
   if (!Caller)
     return true;
 
   // 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)
+  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 {
+  } 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
     // omitting at the call to the kernel from the callgraph.  This ensures
     // that, when compiling for host, only HD functions actually called from the
     // host get marked as known-emitted.
     if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
       CUDACallGraph[Caller].insert({Callee, Loc});

clang/test/SemaCUDA/call-device-fn-from-host.cu

__host__ __device__ void fn_ptr() {
   auto* ptr = &device_fn;
   // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
 }
 
 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>>>(); }