From 0b2af1a2889423bb797856841ac81cf10d01c696 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 26 Jan 2021 17:11:01 -0500 Subject: [NFC][CUDA] Refactor registering device variable Extract registering device variable to CUDA runtime codegen function since it will be called in multiple places. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D95558 --- clang/lib/CodeGen/CodeGenModule.cpp | 55 ++----------------------------------- 1 file changed, 2 insertions(+), 53 deletions(-) (limited to 'clang/lib/CodeGen/CodeGenModule.cpp') diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index c9cf4076579b..b133d2a84a59 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4297,59 +4297,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, (D->hasAttr() || D->hasAttr())) GV->setExternallyInitialized(true); } else { - // Host-side shadows of external declarations of device-side - // global variables become internal definitions. These have to - // be internal in order to prevent name conflicts with global - // host variables with the same name in a different TUs. - if (D->hasAttr() || D->hasAttr()) { - Linkage = llvm::GlobalValue::InternalLinkage; - // Shadow variables and their properties must be registered with CUDA - // runtime. Skip Extern global variables, which will be registered in - // the TU where they are defined. - // - // Don't register a C++17 inline variable. The local symbol can be - // discarded and referencing a discarded local symbol from outside the - // comdat (__cuda_register_globals) is disallowed by the ELF spec. - // TODO: Reject __device__ constexpr and __device__ inline in Sema. - if (!D->hasExternalStorage() && !D->isInline()) - getCUDARuntime().registerDeviceVar(D, *GV, !D->hasDefinition(), - D->hasAttr()); - } else if (D->hasAttr()) { - // __shared__ variables are odd. Shadows do get created, but - // they are not registered with the CUDA runtime, so they - // can't really be used to access their device-side - // counterparts. It's not clear yet whether it's nvcc's bug or - // a feature, but we've got to do the same for compatibility. - Linkage = llvm::GlobalValue::InternalLinkage; - } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() || - D->getType()->isCUDADeviceBuiltinTextureType()) { - // Builtin surfaces and textures and their template arguments are - // also registered with CUDA runtime. - Linkage = llvm::GlobalValue::InternalLinkage; - const ClassTemplateSpecializationDecl *TD = - cast( - D->getType()->getAs()->getDecl()); - const TemplateArgumentList &Args = TD->getTemplateArgs(); - if (TD->hasAttr()) { - assert(Args.size() == 2 && - "Unexpected number of template arguments of CUDA device " - "builtin surface type."); - auto SurfType = Args[1].getAsIntegral(); - if (!D->hasExternalStorage()) - getCUDARuntime().registerDeviceSurf(D, *GV, !D->hasDefinition(), - SurfType.getSExtValue()); - } else { - assert(Args.size() == 3 && - "Unexpected number of template arguments of CUDA device " - "builtin texture type."); - auto TexType = Args[1].getAsIntegral(); - auto Normalized = Args[2].getAsIntegral(); - if (!D->hasExternalStorage()) - getCUDARuntime().registerDeviceTex(D, *GV, !D->hasDefinition(), - TexType.getSExtValue(), - Normalized.getZExtValue()); - } - } + getCUDARuntime().internalizeDeviceSideVar(D, Linkage); + getCUDARuntime().handleVarRegistration(D, *GV); } } -- cgit v1.2.3