[CUDA] When we emit an error that might have been deferred, also print a callstack.

Summary:
Previously, when you did something not allowed in a host+device function
and then caused it to be codegen'ed, we would print out an error telling
you that you did something bad, but we wouldn't tell you how we decided
that the function needed to be codegen'ed.

This change causes us to print out a callstack when emitting deferred
errors.  This is immensely helpful when debugging highly-templated code,
where it's often unclear how a function became known-emitted.

We only print the callstack once per function, after we print the all
deferred errors.

This patch also switches all of our hashtables to using canonical
FunctionDecls instead of regular FunctionDecls.  This prevents a number
of bugs, some of which are caught by tests added here, in which we
assume that two FDs for the same function have the same pointer value.

Reviewers: rnk

Subscribers: cfe-commits, tra

Differential Revision: https://reviews.llvm.org/D25704

llvm-svn: 284647

7 files changed, 51 insertions, 4 deletions

diff --git a/clang/test/SemaCUDA/bad-calls-on-same-line.cu b/clang/test/SemaCUDA/bad-calls-on-same-line.cu
index e91baff5d28..53d5e128234 100644
--- a/clang/test/SemaCUDA/bad-calls-on-same-line.cu
+++ b/clang/test/SemaCUDA/bad-calls-on-same-line.cu
@@ -35,5 +35,7 @@ inline __host__ __device__ void hd() {
 void host_fn() {
   hd<int>();
   hd<double>();  // expected-note {{function template specialization 'hd<double>'}}
+  // expected-note@-1 {{called by 'host_fn'}}
   hd<float>();  // expected-note {{function template specialization 'hd<float>'}}
+  // expected-note@-1 {{called by 'host_fn'}}
 }
diff --git a/clang/test/SemaCUDA/call-device-fn-from-host.cu b/clang/test/SemaCUDA/call-device-fn-from-host.cu
index ab88338b80d..dc227226101 100644
--- a/clang/test/SemaCUDA/call-device-fn-from-host.cu
+++ b/clang/test/SemaCUDA/call-device-fn-from-host.cu
@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - -verify
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN:   -verify -verify-ignore-unexpected=note
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
index bb6ea230fa2..d484af14172 100644
--- a/clang/test/SemaCUDA/call-host-fn-from-device.cu
+++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -1,5 +1,5 @@
 // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
-// RUN:   -emit-llvm -o /dev/null -verify
+// RUN:   -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
 // are emitted during codegen.
diff --git a/clang/test/SemaCUDA/call-stack-for-deferred-err.cu b/clang/test/SemaCUDA/call-stack-for-deferred-err.cu
new file mode 100644
index 00000000000..ddcaabf4ef5
--- /dev/null
+++ b/clang/test/SemaCUDA/call-stack-for-deferred-err.cu
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// We should emit an error for hd_fn's use of a VLA.  This would have been
+// legal if hd_fn were never codegen'ed on the device, so we should also print
+// out a callstack showing how we determine that hd_fn is known-emitted.
+//
+// Compare to no-call-stack-for-deferred-err.cu.
+
+inline __host__ __device__ void hd_fn(int n);
+inline __device__ void device_fn2() { hd_fn(42); } // expected-note {{called by 'device_fn2'}}
+
+__global__ void kernel() { device_fn2(); } // expected-note {{called by 'kernel'}}
+
+inline __host__ __device__ void hd_fn(int n) {
+  int vla[n]; // expected-error {{variable-length array}}
+}
diff --git a/clang/test/SemaCUDA/exceptions.cu b/clang/test/SemaCUDA/exceptions.cu
index 73d2b9d084e..49568ecac7e 100644
--- a/clang/test/SemaCUDA/exceptions.cu
+++ b/clang/test/SemaCUDA/exceptions.cu
@@ -50,3 +50,6 @@ inline __host__ __device__ void hd3() {
 }
 
 __device__ void call_hd3() { hd3(); }
+#ifdef __CUDA_ARCH__
+// expected-note@-2 {{called by 'call_hd3'}}
+#endif
diff --git a/clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu b/clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu
new file mode 100644
index 00000000000..6dc98695c1e
--- /dev/null
+++ b/clang/test/SemaCUDA/no-call-stack-for-immediate-errs.cu
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// Here we should dump an error about the VLA in device_fn, but we should not
+// print a callstack indicating how device_fn becomes known-emitted, because
+// it's an error to use a VLA in any __device__ function, even one that doesn't
+// get emitted.
+
+inline __device__ void device_fn(int n);
+inline __device__ void device_fn2() { device_fn(42); }
+
+__global__ void kernel() { device_fn2(); }
+
+inline __device__ void device_fn(int n) {
+  int vla[n]; // expected-error {{variable-length array}}
+}
diff --git a/clang/test/SemaCUDA/trace-through-global.cu b/clang/test/SemaCUDA/trace-through-global.cu
index 7a9b8dc72b5..065342fdd11 100644
--- a/clang/test/SemaCUDA/trace-through-global.cu
+++ b/clang/test/SemaCUDA/trace-through-global.cu
@@ -35,10 +35,16 @@ __global__ void kernel(int) { hd2(); }
 template <typename T>
 void launch_kernel() {
   kernel<<<0, 0>>>(T());
-  hd1();
-  hd3(T());
+
+  // Notice that these two diagnostics are different: Because the call to hd1
+  // is not dependent on T, the call to hd1 comes from 'launch_kernel', while
+  // the call to hd3, being dependent, comes from 'launch_kernel<int>'.
+  hd1(); // expected-note {{called by 'launch_kernel'}}
+  hd3(T()); // expected-note {{called by 'launch_kernel<int>'}}
 }
 
 void host_fn() {
   launch_kernel<int>();
+  // expected-note@-1 {{called by 'host_fn'}}
+  // expected-note@-2 {{called by 'host_fn'}}
 }