[PATCH] D25150: [CUDA] Allow static variables in __host__ __device__ functions, so long as they're never codegen'ed for device.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 13 11:54:41 PDT 2016
This revision was automatically updated to reflect the committed changes.
Closed by commit rL284145: [CUDA] Allow static variables in __host__ __device__ functions, so long as… (authored by jlebar).
Changed prior to commit:
https://reviews.llvm.org/D25150?vs=73182&id=74566#toc
Repository:
rL LLVM
https://reviews.llvm.org/D25150
Files:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/SemaCUDA/device-var-init.cu
cfe/trunk/test/SemaCUDA/static-vars-hd.cu
Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6741,8 +6741,8 @@
def err_shared_var_init : Error<
"initialization is not supported for __shared__ variables.">;
def err_device_static_local_var : Error<
- "Within a __device__/__global__ function, "
- "only __shared__ variables may be marked \"static\"">;
+ "within a %select{__device__|__global__|__host__|__host__ __device__}0 "
+ "function, only __shared__ variables may be marked 'static'">;
def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
Index: cfe/trunk/test/SemaCUDA/static-vars-hd.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/static-vars-hd.cu
+++ cfe/trunk/test/SemaCUDA/static-vars-hd.cu
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -S -o /dev/null -verify %s
+// RUN: %clang_cc1 -fcxx-exceptions -S -o /dev/null -D HOST -verify %s
+
+#include "Inputs/cuda.h"
+
+#ifdef HOST
+// expected-no-diagnostics
+#endif
+
+__host__ __device__ void f() {
+ static int x = 42;
+#ifndef HOST
+ // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+#endif
+}
+
+inline __host__ __device__ void g() {
+ static int x = 42; // no error on device because this is never codegen'ed there.
+}
+void call_g() { g(); }
Index: cfe/trunk/test/SemaCUDA/device-var-init.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu
@@ -207,9 +207,9 @@
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
static __device__ int ds;
- // expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
static __constant__ int dc;
- // expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
static int v;
- // expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
}
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -10677,12 +10677,11 @@
// 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 && getLangOpts().CUDAIsDevice &&
- (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>()) &&
- !VD->hasAttr<CUDASharedAttr>()) {
- Diag(VD->getLocation(), diag::err_device_static_local_var);
+ if (getLangOpts().CUDA && !VD->hasAttr<CUDASharedAttr>() &&
+ CUDADiagIfDeviceCode(VD->getLocation(),
+ diag::err_device_static_local_var)
+ << CurrentCUDATarget())
VD->setInvalidDecl();
- }
}
}
@@ -10696,7 +10695,7 @@
if (Init && VD->hasGlobalStorage()) {
if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
VD->hasAttr<CUDASharedAttr>()) {
- assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
+ assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
bool AllowedInit = false;
if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
AllowedInit =
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D25150.74566.patch
Type: text/x-patch
Size: 4076 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20161013/b4d9bdb5/attachment.bin>
More information about the cfe-commits
mailing list