[clang] f38a9e5 - [CUDA] Allow local static variables with target attributes.

Artem Belevich via cfe-commits cfe-commits at lists.llvm.org
Mon Nov 2 14:38:21 PST 2020


Author: Artem Belevich
Date: 2020-11-02T14:37:13-08:00
New Revision: f38a9e51178add132d2c8ae160787fb2175a48a4

URL: https://github.com/llvm/llvm-project/commit/f38a9e51178add132d2c8ae160787fb2175a48a4
DIFF: https://github.com/llvm/llvm-project/commit/f38a9e51178add132d2c8ae160787fb2175a48a4.diff

LOG: [CUDA] Allow local static variables with target attributes.

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

Differential Revision: https://reviews.llvm.org/D88345

Added: 
    

Modified: 
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/lib/Sema/SemaDecl.cpp
    clang/lib/Sema/SemaDeclAttr.cpp
    clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
    clang/test/SemaCUDA/bad-attributes.cu
    clang/test/SemaCUDA/device-var-init.cu

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 7555a5ebbd34..83d968d3ae95 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8179,10 +8179,6 @@ def err_dynamic_var_init : Error<
     "__device__, __constant__, and __shared__ variables.">;
 def err_shared_var_init : Error<
     "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'">;
 def err_cuda_vla : Error<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
@@ -8190,7 +8186,7 @@ def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
 def err_cuda_host_shared : Error<
     "__shared__ local variables not allowed in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
-def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
+def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">;
 def err_cuda_ovl_target : Error<
   "%select{__device__|__global__|__host__|__host__ __device__}0 function %1 "
   "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;

diff  --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 1dcf7a16a897..6af2c188ba50 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -13172,32 +13172,9 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) {
     }
   }
 
-  if (VD->isStaticLocal()) {
+  if (VD->isStaticLocal())
     CheckStaticLocalForDllExport(VD);
 
-    if (dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) {
-      // 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();
-      }();
-    }
-  }
-
   // Perform check for initializers of device-side global variables.
   // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
   // 7.5). We must also apply the same checks to all __shared__

diff  --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ce816c1ab4b7..71fccc4c68e7 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -4394,8 +4394,8 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL))
     return;
   const auto *VD = cast<VarDecl>(D);
-  if (!VD->hasGlobalStorage()) {
-    S.Diag(AL.getLoc(), diag::err_cuda_nonglobal_constant);
+  if (VD->hasLocalStorage()) {
+    S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
     return;
   }
   D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL));
@@ -4456,6 +4456,20 @@ static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
     D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
 }
 
+static void handleDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (checkAttrMutualExclusion<CUDAGlobalAttr>(S, D, AL)) {
+    return;
+  }
+
+  if (const auto *VD = dyn_cast<VarDecl>(D)) {
+    if (VD->hasLocalStorage()) {
+      S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
+      return;
+    }
+  }
+  D->addAttr(::new (S.Context) CUDADeviceAttr(S.Context, AL));
+}
+
 static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *Fn = cast<FunctionDecl>(D);
   if (!Fn->isInlineSpecified()) {
@@ -7526,8 +7540,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
     handleGlobalAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDADevice:
-    handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
-                                                                        AL);
+    handleDeviceAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDAHost:
     handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);

diff  --git a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
index 9cb1c6804a48..d259d3dd78cf 100644
--- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu
+++ b/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 @@ inline __device__ void devfun(const int ** b) {
 
 __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 @@ __global__ void kernel(int *a, const int **b) {
   a[5] = x5;
   b[0] = &w;
   b[1] = &z2;
+  b[2] = &local_static_constant;
+  b[3] = &local_static_device;
   devfun(b);
 }
 

diff  --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu
index d72f74471c23..a990598e7b7e 100644
--- a/clang/test/SemaCUDA/bad-attributes.cu
+++ b/clang/test/SemaCUDA/bad-attributes.cu
@@ -64,11 +64,11 @@ __global__ static inline void foobar() {};
 
 __constant__ int global_constant;
 void host_fn() {
-  __constant__ int c; // expected-error {{__constant__ variables must be global}}
+  __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
   __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}}
 }
 __device__ void device_fn() {
-  __constant__ int c; // expected-error {{__constant__ variables must be global}}
+  __constant__ int c; // expected-error {{__constant__ and __device__ are not allowed on non-static local variables}}
 }
 
 typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}

diff  --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu
index dd5d19a6a2e8..88350f56651c 100644
--- a/clang/test/SemaCUDA/device-var-init.cu
+++ b/clang/test/SemaCUDA/device-var-init.cu
@@ -24,6 +24,12 @@ __constant__ int c_v_f = f();
 
 __shared__ T s_t_i = {2};
 // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__device__ T d_t_i = {2};
+__constant__ T c_t_i = {2};
+
+__device__ ECD d_ecd_i{};
+__shared__ ECD s_ecd_i{};
+__constant__ ECD c_ecd_i{};
 
 __device__ EC d_ec_i(3);
 // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
@@ -196,34 +202,218 @@ __shared__ T_FA_NED s_t_fa_ned;
 __constant__ T_FA_NED c_t_fa_ned;
 // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
-// Verify that only __shared__ local variables may be static on device
-// side and that they are not allowed to be initialized.
+// Verify that local variables may be static on device
+// side and that they conform to the initialization constraints.
+// __shared__ can't be initialized at all and others don't support dynamic initialization.
 __device__ void df_sema() {
-  static __shared__ NCFS s_ncfs;
-  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
-  static __shared__ UC s_uc;
-  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
-  static __shared__ NED s_ned;
-  // 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'}}
   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'}}
   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'}}
   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'}}
+
+
+  // __shared__ does not need to be explicitly static.
+  __shared__ int lsi;
+  // __constant__ and __device__ can not be non-static local
+  __constant__ int lci;
+  // expected-error at -1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+  __device__ int ldi;
+  // expected-error at -1 {{__constant__ and __device__ are not allowed on non-static local variables}}
+
+  // Same test cases as for the globals above.
+
+  static __device__ int d_v_f = f();
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ int s_v_f = f();
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ int c_v_f = f();
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __shared__ T s_t_i = {2};
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __device__ T d_t_i = {2};
+  static __constant__ T c_t_i = {2};
+
+  static __device__ ECD d_ecd_i;
+  static __shared__ ECD s_ecd_i;
+  static __constant__ ECD c_ecd_i;
+
+  static __device__ EC d_ec_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ EC s_ec_i(3);
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ EC c_ec_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ EC d_ec_i2 = {3};
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ EC s_ec_i2 = {3};
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ EC c_ec_i2 = {3};
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ ETC d_etc_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ ETC s_etc_i(3);
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ ETC c_etc_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ ETC d_etc_i2 = {3};
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ ETC s_etc_i2 = {3};
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ ETC c_etc_i2 = {3};
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ UC d_uc;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UC s_uc;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UC c_uc;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ UD d_ud;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UD s_ud;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UD c_ud;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ ECI d_eci;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ ECI s_eci;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ ECI c_eci;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NEC d_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NEC s_nec;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NEC c_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NED d_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NED s_ned;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NED c_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NCV d_ncv;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NCV s_ncv;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NCV c_ncv;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ VD d_vd;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ VD s_vd;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ VD c_vd;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NCF d_ncf;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NCF s_ncf;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NCF c_ncf;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __shared__ NCFS s_ncfs;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+
+  static __device__ UTC d_utc;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UTC s_utc;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UTC c_utc;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ UTC d_utc_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ UTC s_utc_i(3);
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ UTC c_utc_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NETC d_netc;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NETC s_netc;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NETC c_netc;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ NETC d_netc_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ NETC s_netc_i(3);
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ NETC c_netc_i(3);
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ EC_I_EC1 d_ec_i_ec1;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ EC_I_EC1 s_ec_i_ec1;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ EC_I_EC1 c_ec_i_ec1;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_V_T d_t_v_t;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_V_T s_t_v_t;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_V_T c_t_v_t;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_B_NEC d_t_b_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_B_NEC s_t_b_nec;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_B_NEC c_t_b_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_F_NEC d_t_f_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_F_NEC s_t_f_nec;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_F_NEC c_t_f_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_FA_NEC d_t_fa_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_FA_NEC s_t_fa_nec;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_FA_NEC c_t_fa_nec;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_B_NED d_t_b_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_B_NED s_t_b_ned;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_B_NED c_t_b_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_F_NED d_t_f_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_F_NED s_t_f_ned;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_F_NED c_t_f_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+  static __device__ T_FA_NED d_t_fa_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+  static __shared__ T_FA_NED s_t_fa_ned;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  static __constant__ T_FA_NED c_t_fa_ned;
+  // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 }
 
 __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'}}
-#endif
 }
 
 inline __host__ __device__ void hd_emitted_host_only() {


        


More information about the cfe-commits mailing list