Index: clang/lib/AST/ASTContext.cpp =================================================================== --- clang/lib/AST/ASTContext.cpp +++ clang/lib/AST/ASTContext.cpp @@ -11438,9 +11438,9 @@ ((D->hasAttr() && !D->getAttr()->isImplicit()) || (D->hasAttr() && - !D->getAttr()->isImplicit())) && - isa(D) && cast(D)->isFileVarDecl() && - cast(D)->getStorageClass() == SC_Static; + !D->getAttr()->isImplicit()) || + D->hasAttr()) && + isa(D) && cast(D)->getStorageClass() == SC_Static; } bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const { Index: clang/lib/CodeGen/CGDecl.cpp =================================================================== --- clang/lib/CodeGen/CGDecl.cpp +++ clang/lib/CodeGen/CGDecl.cpp @@ -11,6 +11,7 @@ //===----------------------------------------------------------------------===// #include "CGBlocks.h" +#include "CGCUDARuntime.h" #include "CGCXXABI.h" #include "CGCleanup.h" #include "CGDebugInfo.h" @@ -414,15 +415,41 @@ llvm::GlobalVariable *var = cast(addr->stripPointerCasts()); + // CUDA/HIP: need to register static device variable declared in host + // or host device functions. + if (getLangOpts().CUDA && !getLangOpts().CUDAIsDevice && CurFuncDecl) { + if (auto *FD = dyn_cast(CurFuncDecl)) { + if (!FD->hasAttr() && + (!FD->hasAttr() || FD->hasAttr())) + CGM.getCUDARuntime().handleVarRegistration(&D, *var); + } + } + // CUDA's local and local static __shared__ variables should not // have any non-empty initializers. This is ensured by Sema. // Whatever initializer such variable may have when it gets here is // a no-op and should not be emitted. bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice && D.hasAttr(); - // If this value has an initializer, emit it. - if (D.getInit() && !isCudaSharedVar) + // HIP static managed variables need to be emitted as declarations in device + // compilation in host or host device functions. + bool isUndefManagedVar = false; + if (getLangOpts().CUDAIsDevice && D.hasAttr() && + CurFuncDecl) { + if (auto *FD = dyn_cast(CurFuncDecl)) { + if (!FD->hasAttr() && + (!FD->hasAttr() || FD->hasAttr())) { + isUndefManagedVar = true; + } + } + } + if (isUndefManagedVar) { + var->setInitializer(nullptr); + var->setLinkage(llvm::GlobalValue::ExternalLinkage); + } else if (D.getInit() && !isCudaSharedVar) { + // If this value has an initializer, emit it. var = AddInitializerToStaticVarDecl(D, var); + } var->setAlignment(alignment.getAsAlign()); Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -94,6 +94,42 @@ llvm_unreachable("invalid C++ ABI kind"); } +// Helper class for emitting device-side static variables created in host-side +// functions for CUDA/HIP. While we do not emit host-side functions on device, +// we still need to emit the static variables the host code will expect to see +// on the device. +class CUDAStaticDeviceVarEmitter + : public StmtVisitor { +public: + CodeGenFunction CGF; + CUDAStaticDeviceVarEmitter(CodeGenModule &CGM) : CGF(CGM) {} + void Visit(Stmt *S) { + if (!S) + return; + if (auto *DS = dyn_cast(S)) { + for (auto &&D : DS->decls()) { + if (auto *VD = dyn_cast(D)) { + if (VD->hasAttr() || + VD->hasAttr()) { + llvm::GlobalValue::LinkageTypes Linkage = + CGF.CGM.getLLVMLinkageVarDefinition(VD, /*IsConstant=*/false); + return CGF.EmitStaticVarDecl(*VD, Linkage); + } + } + } + } + for (auto &&SS : S->children()) + Visit(SS); + } + void runOn(const FunctionDecl *FD) { + assert(CGF.getLangOpts().CUDAIsDevice); + assert(!FD->hasAttr() && !FD->hasAttr()); + assert(FD->hasBody()); + CGF.CurFuncDecl = FD; + Visit(FD->getBody()); + } +}; + CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO, const PreprocessorOptions &PPO, const CodeGenOptions &CGO, llvm::Module &M, @@ -2748,8 +2784,16 @@ !Global->hasAttr() && !Global->hasAttr() && !Global->getType()->isCUDADeviceBuiltinSurfaceType() && - !Global->getType()->isCUDADeviceBuiltinTextureType()) + !Global->getType()->isCUDADeviceBuiltinTextureType()) { + if (auto *FD = dyn_cast(Global)) { + if (FD->hasBody()) { + // Emit static device or constant variables for host functions. + CUDAStaticDeviceVarEmitter E(*this); + E.runOn(FD); + } + } return; + } } else { // We need to emit host-side 'shadows' for all global // device-side variables because the CUDA runtime needs their Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -7244,6 +7244,25 @@ // Handle attributes prior to checking for duplicates in MergeVarDecl ProcessDeclAttributes(S, NewVD, D); + // CUDA/HIP: Function-scope static variables in device or global functions + // have implicit device or constant attribute. Function-scope static variables + // in host device functions have implicit device or constant attribute in + // device compilation only. + if (getLangOpts().CUDA && SC == SC_Static) { + FunctionDecl *CurFD = getCurFunctionDecl(); + if (CurFD && + (CurFD->hasAttr() || + CurFD->hasAttr()) && + (getLangOpts().CUDAIsDevice || !CurFD->hasAttr()) && + !NewVD->hasAttr() && + !NewVD->hasAttr()) { + if (NewVD->isConstexpr() || NewVD->getType().getQualifiers().hasConst()) + NewVD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext())); + else if (!NewVD->hasAttr()) + NewVD->addAttr(CUDADeviceAttr::CreateImplicit(getASTContext())); + } + } + if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice || getLangOpts().SYCLIsDevice) { if (EmitTLSUnsupportedError && Index: clang/test/AST/ast-dump-func-scope-static-var.cu =================================================================== --- /dev/null +++ clang/test/AST/ast-dump-func-scope-static-var.cu @@ -0,0 +1,53 @@ +// RUN: %clang_cc1 -std=c++11 -ast-dump -x hip %s | FileCheck %s +// RUN: %clang_cc1 -std=c++11 -ast-dump -fcuda-is-device -x hip %s | FileCheck %s + +#include "Inputs/cuda.h" + +// CHECK-LABEL: FunctionDecl {{.*}} fun1 +// CHECK: VarDecl {{.*}} a 'int' static +// CHECK-NOT: CUDADeviceAttr +// CHECK: VarDecl {{.*}} b 'int' static +// CHECK-NEXT: CUDADeviceAttr {{.*}}cuda.h +// CHECK: VarDecl {{.*}} c 'const int' static cinit +// CHECK-NOT: CUDADeviceAttr +// CHECK: VarDecl {{.*}} d 'const int' static constexpr cinit +// CHECK-NOT: CUDADeviceAttr +// CHECK: VarDecl {{.*}} e 'int' static cinit +// CHECK: CUDAConstantAttr {{.*}}cuda.h +// CHECK: VarDecl {{.*}} f 'int' static cinit +// CHECK: HIPManagedAttr {{.*}}cuda.h +// CHECK: CUDADeviceAttr {{.*}}Implicit +// CHECK-NOT: CUDADeviceAttr +void fun1() { + static int a; + static __device__ int b; + static const int c = 1; + static constexpr int d = 1; + static __constant__ int e = 1; + static __managed__ int f = 1; +} + +// CHECK-LABEL: FunctionDecl {{.*}} fun2 +// CHECK: VarDecl {{.*}} a 'int' static +// CHECK-NEXT: CUDADeviceAttr {{.*}}Implicit +// CHECK: VarDecl {{.*}} b 'int' static +// CHECK-NEXT: CUDADeviceAttr {{.*}}cuda.h +// CHECK: VarDecl {{.*}} c 'const int' static cinit +// CHECK: CUDAConstantAttr {{.*}}Implicit +// CHECK-NOT: CUDADeviceAttr +// CHECK: VarDecl {{.*}} d 'const int' static constexpr cinit +// CHECK: CUDAConstantAttr {{.*}}Implicit +// CHECK-NOT: CUDADeviceAttr +// CHECK: VarDecl {{.*}} e 'int' static cinit +// CHECK: CUDAConstantAttr {{.*}}cuda.h +// CHECK: VarDecl {{.*}} f 'int' static cinit +// CHECK: HIPManagedAttr {{.*}}cuda.h +// CHECK: CUDADeviceAttr {{.*}}Implicit +__device__ void fun2() { + static int a; + static __device__ int b; + static const int c = 1; + static constexpr int d = 1; + static __constant__ int e = 1; + static __managed__ int f = 1; +} Index: clang/test/CodeGenCUDA/func-scope-static-var.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/func-scope-static-var.cu @@ -0,0 +1,168 @@ +// REQUIRES: x86-registered-target, amdgpu-registered-target + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=DEV,NORDC %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \ +// RUN: -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=DEV,RDC %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=HOST %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-llvm -fgpu-rdc -o - -x hip %s | FileCheck \ +// RUN: -check-prefixes=HOST %s + +#include "Inputs/cuda.h" + +// In device functions, static device variables are not externalized nor shadowed. +// Static managed variable behaves like a normal static device variable. + +// DEV: @_ZZ4fun1vE1a = internal addrspace(1) global i32 1 +// HOST-NOT: @_ZZ4fun1vE1a +// DEV: @_ZZ4fun1vE1b = internal addrspace(1) global i32 2 +// HOST-NOT: @_ZZ4fun1vE1b +// DEV: @_ZZ4fun1vE1c = internal addrspace(4) constant i32 3 +// HOST-NOT: @_ZZ4fun1vE1c +// DEV: @_ZZ4fun1vE1d = internal addrspace(4) constant i32 4 +// HOST-NOT: @_ZZ4fun1vE1d +// DEV: @_ZZ4fun1vE1e = internal addrspace(4) global i32 5 +// HOST-NOT: @_ZZ4fun1vE1e +// DEV: @_ZZ4fun1vE1f = internal addrspace(1) global i32 6 +// HOST-NOT: @_ZZ4fun1vE1f +__device__ int fun1() { + static int a = 1; + static __device__ int b = 2; + static const int c = 3; + static constexpr int d = 4; + static __constant__ int e = 5; + static __managed__ int f = 6; + return a + b + c + d + e + f; +} + +// Assuming this function accepts a device pointer and does some work. +__host__ __device__ int work(int *x); + +// In host function, static device variables are externalized if used and shadowed. + +// DEV-NOT: @_ZZ4fun2vE1a +// HOST: @_ZZ4fun2vE1a = internal global i32 1 +// NORDC: @_ZZ4fun2vE1b = dso_local addrspace(1) global i32 2 +// RDC: @_ZZ4fun2vE1b = internal addrspace(1) global i32 2 +// HOST: @_ZZ4fun2vE1b = internal global i32 2 +// DEV-NOT: @_ZZ4fun2vE1c +// HOST: @_ZZ4fun2vE1c = internal constant i32 3 +// DEV-NOT: @_ZZ4fun2vE1d +// HOST: @_ZZ4fun2vE1d = internal constant i32 4 +// NORDC: @_ZZ4fun2vE1e = dso_local addrspace(4) global i32 5 +// RDC: @_ZZ4fun2vE1e = internal addrspace(4) global i32 5 +// HOST: @_ZZ4fun2vE1e = internal global i32 5 +// DEV: @_ZZ4fun2vE1f = internal addrspace(1) global i32* addrspacecast (i32 addrspace(1)* @_ZZ4fun2vE1b to i32*) +// HOST: @_ZZ4fun2vE1f = internal global i32* @_ZZ4fun2vE1b +// NORDC: @_ZZ4fun2vE1b_0 = dso_local addrspace(1) global i32 6 +// RDC: @_ZZ4fun2vE1b_0 = internal addrspace(1) global i32 6 +// HOST: @_ZZ4fun2vE1b_0 = internal global i32 6 +// NORDC: @_ZZ4fun2vE1g = dso_local addrspace(1) externally_initialized global i32 undef +// RDC: @_ZZ4fun2vE1g = external dso_local addrspace(1) global i32 +// HOST: @_ZZ4fun2vE1g = internal global i32 7 +int fun2() { + static int a = 1; + static __device__ int b = 2; + static const int c = 3; + static constexpr int d = 4; + static __constant__ int e = 5; + static __device__ int *f = &b; + for (int i = 0; i < 10; i++) { + static __device__ int b = 6; + work(&b); + } + static __managed__ int g = 7; + return a + c + d + work(&e) + g; +} + +// In host device function, explicit static device variables are externalized +// if used and registered. Static variables w/o attributes are implicit device +// variables in device compilation and host variables in host compilation. +// The variable emitted in host compilation is not the shadow variable of the +// variable emitted in device compilation. + +// DEV: @_ZZ4fun3vE1a = internal addrspace(1) global i32 1 +// HOST: @_ZZ4fun3vE1a = internal global i32 1 +// NORDC: @_ZZ4fun3vE1b = dso_local addrspace(1) global i32 2 +// RDC: @_ZZ4fun3vE1b = internal addrspace(1) global i32 2 +// HOST: @_ZZ4fun3vE1b = internal global i32 2 +// DEV: @_ZZ4fun3vE1c = internal addrspace(4) constant i32 3 +// HOST: @_ZZ4fun3vE1c = internal constant i32 3 +// DEV: @_ZZ4fun3vE1d = internal addrspace(4) constant i32 4 +// HOST: @_ZZ4fun3vE1d = internal constant i32 4 +// NORDC: @_ZZ4fun3vE1e = dso_local addrspace(4) global i32 5 +// RDC: @_ZZ4fun3vE1e = internal addrspace(4) global i32 5 +// HOST: @_ZZ4fun3vE1e = internal global i32 5 +// DEV: @_ZZ4fun3vE1f = internal addrspace(1) global i32* addrspacecast (i32 addrspace(1)* @_ZZ4fun3vE1b to i32*) +// HOST: @_ZZ4fun3vE1f = internal global i32* @_ZZ4fun3vE1b +// NORDC: @_ZZ4fun3vE1b_0 = dso_local addrspace(1) global i32 6 +// RDC: @_ZZ4fun3vE1b_0 = internal addrspace(1) global i32 6 +// HOST: @_ZZ4fun3vE1b_0 = internal global i32 6 +// NORDC: @_ZZ4fun3vE1g = dso_local addrspace(1) externally_initialized global i32 undef +// RDC: @_ZZ4fun3vE1g = external dso_local addrspace(1) global i32 +// HOST: @_ZZ4fun3vE1g = internal global i32 7 +__host__ __device__ int fun3() { + static int a = 1; + static __device__ int b = 2; + static const int c = 3; + static constexpr int d = 4; + static __constant__ int e = 5; + static __device__ int *f = &b; + for (int i = 0; i < 10; i++) { + static __device__ int b = 6; + work(&b); + } + static __managed__ int g = 7; + return a + c + d + work(&e) + g; +} + +// In kernels, static device variables are not externalized nor shadowed +// since they cannot be accessed by host code. Static managed variable behaves +// like a normal static device variable. + +// DEV: @_ZZ4fun4vE1a = internal addrspace(1) global i32 1 +// HOST-NOT: @_ZZ4fun4vE1a +// DEV: @_ZZ4fun4vE1b = internal addrspace(1) global i32 2 +// HOST-NOT: @_ZZ4fun4vE1b +// DEV: @_ZZ4fun4vE1c = internal addrspace(4) constant i32 3 +// HOST-NOT: @_ZZ4fun4vE1c +// DEV: @_ZZ4fun4vE1d = internal addrspace(4) constant i32 4 +// HOST-NOT: @_ZZ4fun4vE1d +// DEV: @_ZZ4fun4vE1e = internal addrspace(4) global i32 5 +// HOST-NOT: @_ZZ4fun4vE1e +// DEV: @_ZZ4fun4vE1f = internal addrspace(1) global i32 6 +// HOST-NOT: @_ZZ4fun4vE1f +__global__ void fun4() { + static int a = 1; + static __device__ int b = 2; + static const int c = 3; + static constexpr int d = 4; + static __constant__ int e = 5; + static __managed__ int f = 6; +} + +// HOST-NOT: call void @__hipRegisterVar({{.*}}@_ZZ4fun1vE1f +// HOST-NOT: call void @__hipRegisterManagedVar({{.*}}@_ZZ4fun1vE1f +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun2vE1b +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun2vE1e +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun2vE1f +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun2vE1b_0 +// HOST: call void @__hipRegisterManagedVar({{.*}}@_ZZ4fun2vE1g +// HOST-NOT: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1a +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1b +// HOST-NOT: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1c +// HOST-NOT: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1d +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1e +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1f +// HOST: call void @__hipRegisterVar({{.*}}@_ZZ4fun3vE1b_0 +// HOST: call void @__hipRegisterManagedVar({{.*}}@_ZZ4fun3vE1g +// HOST-NOT: call void @__hipRegisterVar({{.*}}@_ZZ4fun4vE1f +// HOST-NOT: call void @__hipRegisterManagedVar({{.*}}@_ZZ4fun4vE1f 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 @@ -14,7 +14,7 @@ // 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 +// DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(4) constant i32 43 // Check a static device variable referenced by host function is externalized. // DEV-DAG: @_ZL1x ={{.*}} addrspace(1) externally_initialized global i32 0 Index: clang/test/SemaCUDA/func-scope-static-var.cu =================================================================== --- /dev/null +++ clang/test/SemaCUDA/func-scope-static-var.cu @@ -0,0 +1,115 @@ +// RUN: %clang_cc1 -std=c++14 -fsyntax-only -verify=host,com -x hip %s +// RUN: %clang_cc1 -std=c++14 -fsyntax-only -fcuda-is-device -verify=dev,com -x hip %s +// RUN: %clang_cc1 -std=c++14 -fsyntax-only -fgpu-rdc -verify=host,com -x hip %s +// RUN: %clang_cc1 -std=c++14 -fsyntax-only -fgpu-rdc -fcuda-is-device -verify=dev,com -x hip %s + +#include "Inputs/cuda.h" + +struct A { + static int a; + static __device__ int fun(); +}; + +int A::a; +__device__ int A::fun() { + return a; + // dev-error@-1 {{reference to __host__ variable 'a' in __device__ function}} +} + +// Assuming this function accepts a pointer to a device variable and calculate some result. +__device__ __host__ int work(const int *x); + +int fun1(int x) { + static __device__ int a = sizeof(a); + static __device__ int b = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const __device__ int c = sizeof(a); + static constexpr __device__ int d = sizeof(a); + static __constant__ __device__ int e = sizeof(a); + static __managed__ __device__ int f = sizeof(a); + static int a2 = sizeof(a); + static int b2 = x; + static const int c2 = sizeof(a); + static constexpr int d2 = sizeof(a); + static __constant__ int e2 = sizeof(a); + static __managed__ int f2 = sizeof(a); + return work(&a) + work(&b) + work(&c) + work(&d) + work(&e) + f + a2 + b2 + c2 + d2 + work(&e2) + f2; +} + +__device__ int fun2(int x) { + static __device__ int a = sizeof(a); + static __device__ int b = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const __device__ int c = sizeof(a); + static constexpr __device__ int d = sizeof(a); + static __constant__ __device__ int e = sizeof(a); + static __managed__ __device__ int f = sizeof(a); + static int a2 = sizeof(a); + static int b2 = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const int c2 = sizeof(a); + static constexpr int d2 = sizeof(a); + static __constant__ int e2 = sizeof(a); + static __managed__ int f2 = sizeof(a); + return a + b + c + d + e + f + a2 + b2 + c2 + d2 + e2 + f2; +} + +__device__ __host__ int fun3(int x) { + static __device__ int a = sizeof(a); + static __device__ int b = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const __device__ int c = sizeof(a); + static constexpr __device__ int d = sizeof(a); + static __constant__ __device__ int e = sizeof(a); + static __managed__ __device__ int f = sizeof(a); + static int a2 = sizeof(a); + static int b2 = x; + // dev-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const int c2 = sizeof(a); + static constexpr int d2 = sizeof(a); + static __constant__ int e2 = sizeof(a); + static __managed__ int f2 = sizeof(a); + return work(&a) + work(&b) + work(&c) + work(&d) + work(&e) + f + a2 + b2 + c2 + d2 + work(&e2) + f2; +} + +template +__device__ __host__ int fun4(T x) { + static __device__ int a = sizeof(x); + static __device__ int b = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const __device__ int c = sizeof(x); + static constexpr __device__ int d = sizeof(x); + static __constant__ __device__ int e = sizeof(a); + static __managed__ __device__ int f = sizeof(a); + static int a2 = sizeof(x); + static int b2 = x; + // dev-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const int c2 = sizeof(x); + static constexpr int d2 = sizeof(x); + static __constant__ int e2 = sizeof(a); + static __managed__ int f2 = sizeof(a); + return work(&a) + work(&b) + work(&c) + work(&d) + work(&e) + f + a2 + b2 + c2 + d2 + work(&e2) + f2; +} + +__device__ __host__ int fun4_caller() { + return fun4(1); + // com-note@-1 {{in instantiation of function template specialization 'fun4' requested here}} +} + +__global__ void fun5(int x, int *y) { + static __device__ int a = sizeof(a); + static __device__ int b = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const __device__ int c = sizeof(a); + static constexpr __device__ int d = sizeof(a); + static __constant__ __device__ int e = sizeof(a); + static __managed__ __device__ int f = sizeof(a); + static int a2 = sizeof(a); + static int b2 = x; + // com-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static const int c2 = sizeof(a); + static constexpr int d2 = sizeof(a); + static __constant__ int e2 = sizeof(a); + static __managed__ int f2 = sizeof(a); + *y = a + b + c + d + e + f + a2 + b2 + c2 + d2 + e2 + f2; +}