summaryrefslogtreecommitdiff
path: root/test/SemaCUDA
diff options
context:
space:
mode:
authorJustin Lebar <jlebar@google.com>2016-10-13 18:45:17 +0000
committerJustin Lebar <jlebar@google.com>2016-10-13 18:45:17 +0000
commit75a3d319e9bdad831c47f449d7406a5ca61781c5 (patch)
treebb1c0a489146a723cce6e82735a2b3265a36f790 /test/SemaCUDA
parenta1f58c85829f2c491e9cf32445f8c8287f11d80f (diff)
[CUDA] Allow static variables in __host__ __device__ functions, so long as they're never codegen'ed for device.
Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25150 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284145 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/SemaCUDA')
-rw-r--r--test/SemaCUDA/device-var-init.cu6
-rw-r--r--test/SemaCUDA/static-vars-hd.cu20
2 files changed, 23 insertions, 3 deletions
diff --git a/test/SemaCUDA/device-var-init.cu b/test/SemaCUDA/device-var-init.cu
index d807a51b65..122dfca423 100644
--- a/test/SemaCUDA/device-var-init.cu
+++ b/test/SemaCUDA/device-var-init.cu
@@ -207,9 +207,9 @@ __device__ void df_sema() {
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
static __device__ int ds;
- // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
static __constant__ int dc;
- // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
static int v;
- // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
}
diff --git a/test/SemaCUDA/static-vars-hd.cu b/test/SemaCUDA/static-vars-hd.cu
new file mode 100644
index 0000000000..70cc0417f0
--- /dev/null
+++ b/test/SemaCUDA/static-vars-hd.cu
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -S -o /dev/null -verify %s
+// RUN: %clang_cc1 -fcxx-exceptions -S -o /dev/null -D HOST -verify %s
+
+#include "Inputs/cuda.h"
+
+#ifdef HOST
+// expected-no-diagnostics
+#endif
+
+__host__ __device__ void f() {
+ static int x = 42;
+#ifndef HOST
+ // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+#endif
+}
+
+inline __host__ __device__ void g() {
+ static int x = 42; // no error on device because this is never codegen'ed there.
+}
+void call_g() { g(); }