Author: jlebar Date: Thu Oct 13 13:45:17 2016 New Revision: 284145 URL: http://llvm.org/viewvc/llvm-project?rev=284145&view=rev Log: [CUDA] Allow static variables in __host__ __device__ functions, so long as they're never codegen'ed for device.
Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25150 Added: cfe/trunk/test/SemaCUDA/static-vars-hd.cu Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/lib/Sema/SemaDecl.cpp cfe/trunk/test/SemaCUDA/device-var-init.cu Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=284145&r1=284144&r2=284145&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Oct 13 13:45:17 2016 @@ -6741,8 +6741,8 @@ def err_dynamic_var_init : Error< 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">; Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=284145&r1=284144&r2=284145&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu Oct 13 13:45:17 2016 @@ -10677,12 +10677,11 @@ Sema::FinalizeDeclaration(Decl *ThisDecl // 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 @@ Sema::FinalizeDeclaration(Decl *ThisDecl 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 = Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=284145&r1=284144&r2=284145&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/device-var-init.cu (original) +++ cfe/trunk/test/SemaCUDA/device-var-init.cu Thu Oct 13 13:45:17 2016 @@ -207,9 +207,9 @@ __device__ void df_sema() { // expected-error@-1 {{initialization is not supported for __shared__ variables.}} static __device__ int ds; - // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} + // expected-error@-1 {{within a __device__ 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"}} + // expected-error@-1 {{within a __device__ 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"}} + // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}} } Added: cfe/trunk/test/SemaCUDA/static-vars-hd.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/static-vars-hd.cu?rev=284145&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/static-vars-hd.cu (added) +++ cfe/trunk/test/SemaCUDA/static-vars-hd.cu Thu Oct 13 13:45:17 2016 @@ -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@-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(); } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
