[CUDA] Emit errors for wrong-side calls made on the same line as non-wrong-side calls.

Summary:
This fixes two related bugs:

1) Previously, if you had a non-wrong side call at some source code
location L, we wouldn't emit errors for wrong-side calls that appeared
at L.

2) We'd only emit one wrong-side error per source code location, when we
actually want to emit it twice if we hit this line more than once due to
e.g. template instantiation.

Reviewers: tra

Subscribers: rnk, cfe-commits

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

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

1 files changed, 39 insertions, 0 deletions

diff --git a/test/SemaCUDA/bad-calls-on-same-line.cu b/test/SemaCUDA/bad-calls-on-same-line.cu
new file mode 100644
index 0000000000..e91baff5d2
--- /dev/null
+++ b/test/SemaCUDA/bad-calls-on-same-line.cu
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+// The hd function template is instantiated three times.
+//
+// Two of those instantiations call a device function, which is an error when
+// compiling for host.  Clang should report both errors.
+
+#include "Inputs/cuda.h"
+
+template <typename T>
+struct Selector {};
+
+template <>
+struct Selector<int> {
+  __host__ void f() {}
+};
+
+template <>
+struct Selector<float> {
+  __device__ void f() {} // expected-note {{declared here}}
+};
+
+template <>
+struct Selector<double> {
+  __device__ void f() {} // expected-note {{declared here}}
+};
+
+template <typename T>
+inline __host__ __device__ void hd() {
+  Selector<T>().f();
+  // expected-error@-1 {{reference to __device__ function}}
+  // expected-error@-2 {{reference to __device__ function}}
+}
+
+void host_fn() {
+  hd<int>();
+  hd<double>();  // expected-note {{function template specialization 'hd<double>'}}
+  hd<float>();  // expected-note {{function template specialization 'hd<float>'}}
+}