summaryrefslogtreecommitdiff
path: root/test/CodeGenCUDA
diff options
context:
space:
mode:
authorArtem Belevich <tra@google.com>2016-05-09 19:36:08 +0000
committerArtem Belevich <tra@google.com>2016-05-09 19:36:08 +0000
commita93fdcf11017c5e79d5e130d7f981763b218d8bf (patch)
treee96f03079fb321a2269dd22505d033180f15e643 /test/CodeGenCUDA
parente3595c5ee135331dc89f3a8dce00956122fc028e (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.cu8
-rw-r--r--test/CodeGenCUDA/device-var-init.cu8
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)