Author: Yaxun (Sam) Liu Date: 2020-08-10T19:02:49-04:00 New Revision: fb04d7b4a69831f6b999b1776da738557b108e0d
URL: https://github.com/llvm/llvm-project/commit/fb04d7b4a69831f6b999b1776da738557b108e0d DIFF: https://github.com/llvm/llvm-project/commit/fb04d7b4a69831f6b999b1776da738557b108e0d.diff LOG: [CUDA][HIP] Do not externalize implicit constant static variable Differential Revision: https://reviews.llvm.org/D85686 Added: Modified: clang/include/clang/AST/ASTContext.h clang/lib/AST/ASTContext.cpp clang/lib/Sema/SemaExpr.cpp clang/test/CodeGenCUDA/static-device-var-no-rdc.cu Removed: ################################################################################ diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 78207a4aad31..b5fb1a4da480 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -3034,6 +3034,9 @@ OPT_LIST(V) /// Return a new OMPTraitInfo object owned by this context. OMPTraitInfo &getNewOMPTraitInfo(); + /// Whether a C++ static variable may be externalized. + bool mayExternalizeStaticVar(const Decl *D) const; + /// Whether a C++ static variable should be externalized. bool shouldExternalizeStaticVar(const Decl *D) const; diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 61e39dc176a9..7947de3a31fc 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -11191,10 +11191,17 @@ clang::operator<<(const DiagnosticBuilder &DB, return DB << "a prior #pragma section"; } -bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const { +bool ASTContext::mayExternalizeStaticVar(const Decl *D) const { return !getLangOpts().GPURelocatableDeviceCode && - (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) && + ((D->hasAttr<CUDADeviceAttr>() && + !D->getAttr<CUDADeviceAttr>()->isImplicit()) || + (D->hasAttr<CUDAConstantAttr>() && + !D->getAttr<CUDAConstantAttr>()->isImplicit())) && isa<VarDecl>(D) && cast<VarDecl>(D)->isFileVarDecl() && - cast<VarDecl>(D)->getStorageClass() == SC_Static && + cast<VarDecl>(D)->getStorageClass() == SC_Static; +} + +bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const { + return mayExternalizeStaticVar(D) && CUDAStaticDeviceVarReferencedByHost.count(cast<VarDecl>(D)); } diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 0d7f81e0d01a..79340a9c7d7b 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -17910,8 +17910,7 @@ static void DoMarkVarDeclReferenced(Sema &SemaRef, SourceLocation Loc, // This also requires the reference of the static device/constant variable by // host code to be visible in the device compilation for the compiler to be // able to externalize the static device/constant variable. - if ((Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>()) && - Var->isFileVarDecl() && Var->getStorageClass() == SC_Static) { + if (SemaRef.getASTContext().mayExternalizeStaticVar(Var)) { auto *CurContext = SemaRef.CurContext; if (!CurContext || !isa<FunctionDecl>(CurContext) || cast<FunctionDecl>(CurContext)->hasAttr<CUDAHostAttr>() || diff --git a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu index 1aea467c2d49..c7beb4c7e1ac 100644 --- a/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu +++ b/clang/test/CodeGenCUDA/static-device-var-no-rdc.cu @@ -1,11 +1,11 @@ // REQUIRES: x86-registered-target // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \ +// 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 %s -// RUN: %clang_cc1 -triple x86_64-gnu-linux \ +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ // RUN: -emit-llvm -o - -x hip %s | FileCheck \ // RUN: -check-prefixes=HOST %s @@ -53,6 +53,12 @@ static __constant__ int y; // DEV-NOT: @_ZL1z static int z; +// Test implicit static constant variable, which should not be externalized. +// HOST-DAG: @_ZL2z2 = internal constant i32 456 +// DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 456 + +static constexpr int z2 = 456; + // Test static device variable in inline function, which should not be // externalized nor registered. // DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat @@ -72,6 +78,7 @@ __global__ void kernel(int *a, const int **b) { a[4] = x4; a[5] = x5; b[0] = &w; + b[1] = &z2; devfun(b); } @@ -81,11 +88,12 @@ __host__ __device__ void hdf(int *a) { int* getDeviceSymbol(int *x); -void foo(int *a) { +void foo(const int **a) { getDeviceSymbol(&x); getDeviceSymbol(&x5); getDeviceSymbol(&y); z = 123; + a[0] = &z2; } // HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]] _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
