[CUDA] Add Sema::CUDADiagBuilder and Sema::CUDADiagIf{Device,Host}Code().

Summary:
Together these let you easily create diagnostics that

 - are never emitted for host code
 - are always emitted for __device__ and __global__ functions, and
 - are emitted for __host__ __device__ functions iff these functions are
   codegen'ed.

At the moment there are only three diagnostics that need this treatment,
but I have more to add, and it's not sustainable to write code for emitting
every such diagnostic twice, and from a special wrapper in SemaCUDA.cpp.

While we're at it, don't emit the function name in
err_cuda_device_exceptions: It's not necessary to print it, and making
this work in the new framework in the face of a null value for
dyn_cast<FunctionDecl>(CurContext) isn't worth the effort.

Reviewers: rnk

Subscribers: cfe-commits, tra

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

llvm-svn: 284143

2 files changed, 8 insertions, 8 deletions

diff --git a/clang/test/SemaCUDA/exceptions-host-device.cu b/clang/test/SemaCUDA/exceptions-host-device.cu
index 3f323469ac1..9e18634e905 100644
--- a/clang/test/SemaCUDA/exceptions-host-device.cu
+++ b/clang/test/SemaCUDA/exceptions-host-device.cu
@@ -14,8 +14,8 @@ __host__ __device__ void hd1() {
   throw NULL;
   try {} catch(void*) {}
 #ifndef HOST
-  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function 'hd1'}}
-  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function 'hd1'}}
+  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
+  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
 #endif
 }
 
@@ -31,8 +31,8 @@ inline __host__ __device__ void hd3() {
   throw NULL;
   try {} catch(void*) {}
 #ifndef HOST
-  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function 'hd3'}}
-  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function 'hd3'}}
+  // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
+  // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
 #endif
 }
 __device__ void call_hd3() { hd3(); }
diff --git a/clang/test/SemaCUDA/exceptions.cu b/clang/test/SemaCUDA/exceptions.cu
index 75586bc4a97..9ed9b697720 100644
--- a/clang/test/SemaCUDA/exceptions.cu
+++ b/clang/test/SemaCUDA/exceptions.cu
@@ -9,13 +9,13 @@ void host() {
 }
 __device__ void device() {
   throw NULL;
-  // expected-error@-1 {{cannot use 'throw' in __device__ function 'device'}}
+  // expected-error@-1 {{cannot use 'throw' in __device__ function}}
   try {} catch(void*) {}
-  // expected-error@-1 {{cannot use 'try' in __device__ function 'device'}}
+  // expected-error@-1 {{cannot use 'try' in __device__ function}}
 }
 __global__ void kernel() {
   throw NULL;
-  // expected-error@-1 {{cannot use 'throw' in __global__ function 'kernel'}}
+  // expected-error@-1 {{cannot use 'throw' in __global__ function}}
   try {} catch(void*) {}
-  // expected-error@-1 {{cannot use 'try' in __global__ function 'kernel'}}
+  // expected-error@-1 {{cannot use 'try' in __global__ function}}
 }