diff options
| author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2021-01-26 17:11:01 -0500 |
|---|---|---|
| committer | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2021-02-03 14:29:51 -0500 |
| commit | 0b2af1a2889423bb797856841ac81cf10d01c696 (patch) | |
| tree | 3fd03338b91695155d7bf9cc7ecf2ff9a6870e17 /clang/lib/CodeGen/CodeGenModule.cpp | |
| parent | 4a35941dbde63964240d30309e2b72d8be8fabb6 (diff) | |
[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
Diffstat (limited to 'clang/lib/CodeGen/CodeGenModule.cpp')
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 55 |
1 files changed, 2 insertions, 53 deletions
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<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>())) 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<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) { - 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<CUDAConstantAttr>()); - } else if (D->hasAttr<CUDASharedAttr>()) { - // __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<ClassTemplateSpecializationDecl>( - D->getType()->getAs<RecordType>()->getDecl()); - const TemplateArgumentList &Args = TD->getTemplateArgs(); - if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) { - 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); } } |
