r278767 - [CUDA] Fix "declared here" note on deferred wrong-side errors.

[Justin Lebar via cfe-commits]

Author: jlebar
Date: Mon Aug 15 19:48:21 2016
New Revision: 278767

URL: http://llvm.org/viewvc/llvm-project?rev=278767&view=rev
Log:
[CUDA] Fix "declared here" note on deferred wrong-side errors.

Previously we weren't deferring these "declared here" notes, which is
obviously wrong.

Modified:
    cfe/trunk/lib/Sema/SemaCUDA.cpp
    cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu
    cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu

Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=278767&r1=278766&r2=278767&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Mon Aug 15 19:48:21 2016
@@ -499,11 +499,16 @@ bool Sema::CheckCUDACall(SourceLocation
   if (Pref == Sema::CFP_WrongSide) {
     // We have to do this odd dance to create our PartialDiagnostic because we
     // want its storage to be allocated with operator new, not in an arena.
-    PartialDiagnostic PD{PartialDiagnostic::NullDiagnostic()};
-    PD.Reset(diag::err_ref_bad_target);
-    PD << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
-    Caller->addDeferredDiag({Loc, std::move(PD)});
-    Diag(Callee->getLocation(), diag::note_previous_decl) << Callee;
+    PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()};
+    ErrPD.Reset(diag::err_ref_bad_target);
+    ErrPD << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
+    Caller->addDeferredDiag({Loc, std::move(ErrPD)});
+
+    PartialDiagnostic NotePD{PartialDiagnostic::NullDiagnostic()};
+    NotePD.Reset(diag::note_previous_decl);
+    NotePD << Callee;
+    Caller->addDeferredDiag({Callee->getLocation(), std::move(NotePD)});
+
     // This is not immediately an error, so return true.  The deferred errors
     // will be emitted if and when Caller is codegen'ed.
     return true;

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=278767&r1=278766&r2=278767&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu (original)
+++ cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu Mon Aug 15 19:48:21 2016
@@ -6,10 +6,18 @@
 #include "Inputs/cuda.h"
 
 __device__ void device_fn() {}
+// expected-note at -1 {{'device_fn' declared here}}
+// expected-note at -2 {{'device_fn' declared here}}
+// expected-note at -3 {{'device_fn' declared here}}
+// expected-note at -4 {{'device_fn' declared here}}
+// expected-note at -5 {{'device_fn' declared here}}
 
 struct S {
   __device__ S() {}
+  // expected-note at -1 {{'S' declared here}}
+  // expected-note at -2 {{'S' declared here}}
   __device__ ~S() { device_fn(); }
+  // expected-note at -1 {{'~S' declared here}}
   int x;
 };
 
@@ -24,6 +32,7 @@ struct T {
   __host__ __device__ void hd3();
 
   __device__ void d() {}
+  // expected-note at -1 {{'d' declared here}}
 };
 
 __host__ __device__ void T::hd3() {

Modified: cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu?rev=278767&r1=278766&r2=278767&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu (original)
+++ cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu Mon Aug 15 19:48:21 2016
@@ -6,10 +6,19 @@
 #include "Inputs/cuda.h"
 
 extern "C" void host_fn() {}
+// expected-note at -1 {{'host_fn' declared here}}
+// expected-note at -2 {{'host_fn' declared here}}
+// expected-note at -3 {{'host_fn' declared here}}
+// expected-note at -4 {{'host_fn' declared here}}
+// expected-note at -5 {{'host_fn' declared here}}
+// expected-note at -6 {{'host_fn' declared here}}
 
 struct S {
   S() {}
+  // expected-note at -1 {{'S' declared here}}
+  // expected-note at -2 {{'S' declared here}}
   ~S() { host_fn(); }
+  // expected-note at -1 {{'~S' declared here}}
   int x;
 };
 
@@ -24,6 +33,7 @@ struct T {
   __host__ __device__ void hd3();
 
   void h() {}
+  // expected-note at -1 {{'h' declared here}}
 };
 
 __host__ __device__ void T::hd3() {