diff options
author | Artem Belevich <tra@google.com> | 2016-05-09 19:36:08 +0000 |
---|---|---|
committer | Artem Belevich <tra@google.com> | 2016-05-09 19:36:08 +0000 |
commit | a93fdcf11017c5e79d5e130d7f981763b218d8bf (patch) | |
tree | e96f03079fb321a2269dd22505d033180f15e643 /test/CodeGenCUDA | |
parent | e3595c5ee135331dc89f3a8dce00956122fc028e (diff) |
[CUDA] Only __shared__ variables can be static local on device side.
According to CUDA programming guide (v7.5):
> E.2.9.4: Within the body of a device or global function, only
> shared variables may be declared with static storage class.
Differential Revision: http://reviews.llvm.org/D20034
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@268962 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r-- | test/CodeGenCUDA/address-spaces.cu | 8 | ||||
-rw-r--r-- | test/CodeGenCUDA/device-var-init.cu | 8 |
2 files changed, 8 insertions, 8 deletions
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index 31cba958e1..4670f46db9 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -38,14 +38,6 @@ __device__ void foo() { // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*) k++; - static int li; - // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*) - li++; - - __constant__ int lj; - // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*) - lj++; - __shared__ int lk; // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*) lk++; diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu index 864cc6daee..23c9fe1376 100644 --- a/test/CodeGenCUDA/device-var-init.cu +++ b/test/CodeGenCUDA/device-var-init.cu @@ -368,6 +368,14 @@ __device__ void df() { T_F_NEC t_f_nec; T_FA_NEC t_fa_nec; static __shared__ UC s_uc; +#if ERROR_CASE + static __device__ int ds; + // expected-error@-1 {{Within a __device__/__global__ 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"}} + static int v; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} +#endif } // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec) |