[PATCH] D88345: [CUDA] Allow `static const {__constant__, __device__}` variables.
Artem Belevich via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 25 16:32:00 PDT 2020
tra created this revision.
tra added reviewers: jlebar, yaxunl.
Herald added subscribers: sanjoy.google, bixia.
Herald added a project: clang.
tra requested review of this revision.
While CUDA documentation claims that such variables are not allowed[1], NVCC has
been accepting them since CUDA-10.0 [2] and some headers in CUDA-11 rely on this
working.
1. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#static-variables-function
2. https://godbolt.org/z/zsodzc
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D88345
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/Sema/SemaDecl.cpp
clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
clang/test/SemaCUDA/device-var-init.cu
Index: clang/test/SemaCUDA/device-var-init.cu
===================================================================
--- clang/test/SemaCUDA/device-var-init.cu
+++ clang/test/SemaCUDA/device-var-init.cu
@@ -207,22 +207,20 @@
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
static __device__ int ds;
- // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
static __constant__ int dc;
- // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
static int v;
- // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables may be marked 'static'}}
static const int cv = 1;
static const __device__ int cds = 1;
- // expected-error at -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 at -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 at -2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
+ // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables or const variables may be marked 'static'}}
#endif
}
Index: clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
===================================================================
--- clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
+++ clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
@@ -13,6 +13,8 @@
// Test function scope static device variable, which should not be externalized.
// DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
+// DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42
+// DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43
// Check a static device variable referenced by host function is externalized.
// DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
@@ -78,6 +80,8 @@
__global__ void kernel(int *a, const int **b) {
const static int w = 1;
+ const static __constant__ int local_static_constant = 42;
+ const static __device__ int local_static_device = 43;
a[0] = x;
a[1] = y;
a[2] = x2;
@@ -86,6 +90,8 @@
a[5] = x5;
b[0] = &w;
b[1] = &z2;
+ b[2] = &local_static_constant;
+ b[3] = &local_static_device;
devfun(b);
}
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -13162,10 +13162,7 @@
[&]() {
if (!getLangOpts().CUDA)
return;
- if (VD->hasAttr<CUDASharedAttr>())
- return;
- if (VD->getType().isConstQualified() &&
- !(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+ if (VD->hasAttr<CUDASharedAttr>() || VD->getType().isConstQualified())
return;
if (CUDADiagIfDeviceCode(VD->getLocation(),
diag::err_device_static_local_var)
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8155,8 +8155,8 @@
"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 or const variables without device "
- "memory qualifier may be marked 'static'">;
+ "function, only __shared__ variables or const variables "
+ "may be marked 'static'">;
def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D88345.294451.patch
Type: text/x-patch
Size: 4707 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200925/9f06b8c9/attachment.bin>
More information about the cfe-commits
mailing list