[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