"initialization is not supported for __shared__ variables.">;
def err_device_static_local_var : Error<
"within a %select{__device__|__global__|__host__|__host__ __device__}0 "
- "function, only __shared__ variables may be marked 'static'">;
+ "function, only __shared__ variables or const variables without device "
+ "memory qualifier may be marked 'static'">;
def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
return LangAS::cuda_constant;
else if (D && D->hasAttr<CUDASharedAttr>())
return LangAS::cuda_shared;
+ else if (D && D->hasAttr<CUDADeviceAttr>())
+ return LangAS::cuda_device;
+ else if (D && D->getType().isConstQualified())
+ return LangAS::cuda_constant;
else
return LangAS::cuda_device;
}
NewAttr->setInherited(true);
VD->addAttr(NewAttr);
}
- // CUDA E.2.9.4: Within the body of a __device__ or __global__
- // function, only __shared__ variables may be declared with
- // static storage class.
- if (getLangOpts().CUDA && !VD->hasAttr<CUDASharedAttr>() &&
- CUDADiagIfDeviceCode(VD->getLocation(),
- diag::err_device_static_local_var)
- << CurrentCUDATarget())
- VD->setInvalidDecl();
+ // CUDA 8.0 E.3.9.4: Within the body of a __device__ or __global__
+ // function, only __shared__ variables or variables without any device
+ // memory qualifiers may be declared with static storage class.
+ // Note: It is unclear how a function-scope non-const static variable
+ // without device memory qualifier is implemented, therefore only static
+ // const variable without device memory qualifier is allowed.
+ [&]() {
+ if (!getLangOpts().CUDA)
+ return;
+ if (VD->hasAttr<CUDASharedAttr>())
+ return;
+ if (VD->getType().isConstQualified() &&
+ !(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+ return;
+ if (CUDADiagIfDeviceCode(VD->getLocation(),
+ diag::err_device_static_local_var)
+ << CurrentCUDATarget())
+ VD->setInvalidDecl();
+ }();
}
}
// CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
// CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
+// CHECK: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5]
+// CHECK: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123
+
// We should not emit global initializers for device-side variables.
// CHECK-NOT: @__cxx_global_var_init
static __shared__ ETC s_etc;
// CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
+ static const int const_array[] = {1, 2, 3, 4, 5};
+ static const int const_int = 123;
+
// anchor point separating constructors and destructors
df(); // CHECK: call void @_Z2dfv()
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
static __device__ int ds;
- // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
static __constant__ int dc;
- // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
static int v;
- // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ static const int cv = 1;
+ static const __device__ int cds = 1;
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ static const __constant__ int cdc = 1;
+ // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
}
__host__ __device__ void hd_sema() {
static int x = 42;
#ifdef __CUDA_ARCH__
- // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
#endif
}