diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-07-28 03:05:25 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-07-28 03:05:25 +0000 |
| commit | a4005e13f72ece696cce5d190fb73abec116d18a (patch) | |
| tree | 99111c1e5acba99421c03668201f96acaba2e38e /clang/test/SemaCUDA/device-var-init.cu | |
| parent | 39e5137f43917417a2ce6bb663de39af005bd452 (diff) | |
| download | bcm5719-llvm-a4005e13f72ece696cce5d190fb73abec116d18a.tar.gz bcm5719-llvm-a4005e13f72ece696cce5d190fb73abec116d18a.zip | |
[CUDA][HIP] Allow function-scope static const variable
CUDA 8.0 E.3.9.4 says: 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.
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, which
can be emitted as a global variable in constant address space.
Currently clang only allows function-scope static variable with
__shared__ qualifier.
This patch also allows function-scope static const variable without
device memory qualifier and emits it as a global variable in constant
address space.
Differential Revision: https://reviews.llvm.org/D49931
llvm-svn: 338188
Diffstat (limited to 'clang/test/SemaCUDA/device-var-init.cu')
| -rw-r--r-- | clang/test/SemaCUDA/device-var-init.cu | 13 |
1 files changed, 9 insertions, 4 deletions
diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu index 46cb90da2ec..dd5d19a6a2e 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -207,17 +207,22 @@ __device__ void df_sema() { // 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 } |
