[CUDA] Fix false-positive in known-emitted handling.

Previously: When compiling for host, our constructed call graph went
*through* kernel calls.  This meant that if we had

  host calls kernel calls HD

we would incorrectly mark the HD function as known-emitted on the host
side, and thus perform host-side checks on it.

Fixing this exposed another issue, wherein when marking a function as
known-emitted, we also need to traverse the callgraph of its template,
because non-dependent calls are attached to a function's template, not
its instantiation.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284355 91177308-0d34-0410-b5e6-96231b3b80d8

1 files changed, 44 insertions, 0 deletions

diff --git a/test/SemaCUDA/trace-through-global.cu b/test/SemaCUDA/trace-through-global.cu
new file mode 100644
index 0000000000..7a9b8dc72b
--- /dev/null
+++ b/test/SemaCUDA/trace-through-global.cu
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+// Check that it's OK for kernels to call HD functions that call device-only
+// functions.
+
+#include "Inputs/cuda.h"
+
+__device__ void device_fn(int) {}
+// expected-note@-1 {{declared here}}
+// expected-note@-2 {{declared here}}
+
+inline __host__ __device__ int hd1() {
+  device_fn(0);  // expected-error {{reference to __device__ function}}
+  return 0;
+}
+
+inline __host__ __device__ int hd2() {
+  // No error here because hd2 is only referenced from a kernel.
+  device_fn(0);
+  return 0;
+}
+
+inline __host__ __device__ void hd3(int) {
+  device_fn(0);  // expected-error {{reference to __device__ function 'device_fn'}}
+}
+inline __host__ __device__ void hd3(double) {}
+
+inline __host__ __device__ void hd4(int) {}
+inline __host__ __device__ void hd4(double) {
+  device_fn(0);  // No error; this function is never called.
+}
+
+__global__ void kernel(int) { hd2(); }
+
+template <typename T>
+void launch_kernel() {
+  kernel<<<0, 0>>>(T());
+  hd1();
+  hd3(T());
+}
+
+void host_fn() {
+  launch_kernel<int>();
+}