diff options
author | Justin Lebar <jlebar@google.com> | 2016-10-13 18:45:17 +0000 |
---|---|---|
committer | Justin Lebar <jlebar@google.com> | 2016-10-13 18:45:17 +0000 |
commit | 75a3d319e9bdad831c47f449d7406a5ca61781c5 (patch) | |
tree | bb1c0a489146a723cce6e82735a2b3265a36f790 /test/SemaCUDA | |
parent | a1f58c85829f2c491e9cf32445f8c8287f11d80f (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.cu | 6 | ||||
-rw-r--r-- | test/SemaCUDA/static-vars-hd.cu | 20 |
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(); } |