2 files changed, 59 insertions, 0 deletions
diff --git a/clang/test/SemaCUDA/exceptions-host-device.cu b/clang/test/SemaCUDA/exceptions-host-device.cu
new file mode 100644
index 00000000000..3f323469ac1
--- /dev/null
+++ b/clang/test/SemaCUDA/exceptions-host-device.cu
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -verify %s -S -o /dev/null
+// RUN: %clang_cc1 -fcxx-exceptions -verify -DHOST %s -S -o /dev/null
+
+#include "Inputs/cuda.h"
+
+// Check that it's an error to use 'try' and 'throw' from a __host__ __device__
+// function if and only if it's codegen'ed for device.
+
+#ifdef HOST
+// expected-no-diagnostics
+#endif
+
+__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'}}
+#endif
+}
+
+// No error, never instantiated on device.
+inline __host__ __device__ void hd2() {
+  throw NULL;
+  try {} catch(void*) {}
+}
+void call_hd2() { hd2(); }
+
+// Error, instantiated on device.
+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'}}
+#endif
+}
+__device__ void call_hd3() { hd3(); }
diff --git a/clang/test/SemaCUDA/exceptions.cu b/clang/test/SemaCUDA/exceptions.cu
new file mode 100644
index 00000000000..75586bc4a97
--- /dev/null
+++ b/clang/test/SemaCUDA/exceptions.cu
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+void host() {
+  throw NULL;
+  try {} catch(void*) {}
+}
+__device__ void device() {
+  throw NULL;
+  // expected-error@-1 {{cannot use 'throw' in __device__ function 'device'}}
+  try {} catch(void*) {}
+  // expected-error@-1 {{cannot use 'try' in __device__ function 'device'}}
+}
+__global__ void kernel() {
+  throw NULL;
+  // expected-error@-1 {{cannot use 'throw' in __global__ function 'kernel'}}
+  try {} catch(void*) {}
+  // expected-error@-1 {{cannot use 'try' in __global__ function 'kernel'}}
+}