[CUDA] Disallow exceptions in device code.

Reviewers: tra

Subscribers: cfe-commits, jhen

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

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

2 files changed, 59 insertions, 0 deletions

diff --git a/test/SemaCUDA/exceptions-host-device.cu b/test/SemaCUDA/exceptions-host-device.cu
new file mode 100644
index 0000000000..3f323469ac
--- /dev/null
+++ b/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/test/SemaCUDA/exceptions.cu b/test/SemaCUDA/exceptions.cu
new file mode 100644
index 0000000000..75586bc4a9
--- /dev/null
+++ b/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'}}
+}