Author: Joseph Huber Date: 2026-02-24T09:11:31-06:00 New Revision: 9f2351c27e2e7417880dd99fc8516fff713a5ffe
URL: https://github.com/llvm/llvm-project/commit/9f2351c27e2e7417880dd99fc8516fff713a5ffe DIFF: https://github.com/llvm/llvm-project/commit/9f2351c27e2e7417880dd99fc8516fff713a5ffe.diff LOG: [HIP] Do not apply 'externally_initialized' to constant device variables (#182157) Summary: >From the Language reference: > By default, global initializers are optimized by assuming that global variables defined within the module are not modified from their initial values before the start of the global initializer. This is true even for variables potentially accessible from outside the module, including those with external linkage or appearing in @llvm.used or dllexported variables. This assumption may be suppressed by marking the variable with externally_initialized. This is intended because device programs can be modified beyond the normal lifetime expected by the optimization pipeline. However, for constant variables we should be able to safely assume that these are truly constant within the module. In the vast majority of cases these will not get externally visible symbols, but even `extern const` uses we should assert that the user should not be writing them if they are marked const. Added: Modified: clang/lib/CodeGen/CodeGenModule.cpp clang/test/CodeGenCUDA/const-var.cu clang/test/CodeGenCUDA/constexpr-variables.cu clang/test/CodeGenCUDA/host-used-device-var.cu Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index bff2a5ba5d245..35ee108cdc4fc 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6261,7 +6261,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())." if (LangOpts.CUDA) { if (LangOpts.CUDAIsDevice) { - if (Linkage != llvm::GlobalValue::InternalLinkage && + if (Linkage != llvm::GlobalValue::InternalLinkage && !D->isConstexpr() && + !D->getType().isConstQualified() && (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() || D->getType()->isCUDADeviceBuiltinSurfaceType() || D->getType()->isCUDADeviceBuiltinTextureType())) diff --git a/clang/test/CodeGenCUDA/const-var.cu b/clang/test/CodeGenCUDA/const-var.cu index 70d4df18dfeef..7bb18590f2804 100644 --- a/clang/test/CodeGenCUDA/const-var.cu +++ b/clang/test/CodeGenCUDA/const-var.cu @@ -15,8 +15,8 @@ // Both are promoted to device side. // DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1 -// DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) externally_initialized constant ptr addrspacecast (ptr addrspace(4) @_ZN5Test1L1aE to ptr) -// DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) externally_initialized constant ptr addrspacecast (ptr addrspace(4) @_ZN5Test1L1aE to ptr) +// DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) constant ptr addrspacecast (ptr addrspace(4) @_ZN5Test1L1aE to ptr) +// DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) constant ptr addrspacecast (ptr addrspace(4) @_ZN5Test1L1aE to ptr) // DEV-DAG: @_ZN5Test12b2E = addrspace(1) externally_initialized global i32 1 // HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1 // HOST-DAG: @_ZN5Test11B2p1E = constant ptr @_ZN5Test1L1aE diff --git a/clang/test/CodeGenCUDA/constexpr-variables.cu b/clang/test/CodeGenCUDA/constexpr-variables.cu index 7ae56341cdf57..c528e6f2a5295 100644 --- a/clang/test/CodeGenCUDA/constexpr-variables.cu +++ b/clang/test/CodeGenCUDA/constexpr-variables.cu @@ -16,10 +16,10 @@ namespace B { __constant__ const int &use_B_b = B::b; struct Q { - // CXX14: @_ZN1Q2k2E = {{.*}}externally_initialized constant i32 6 + // CXX14: @_ZN1Q2k2E = {{.*}}constant i32 6 // CXX17: @_ZN1Q2k2E = internal {{.*}}constant i32 6 // CXX14: @_ZN1Q2k1E = available_externally {{.*}}constant i32 5 - // CXX17: @_ZN1Q2k1E = {{.*}} externally_initialized constant i32 5 + // CXX17: @_ZN1Q2k1E = {{.*}} constant i32 5 static constexpr int k1 = 5; static constexpr int k2 = 6; }; @@ -30,14 +30,14 @@ __constant__ const int &use_Q_k2 = Q::k2; template<typename T> struct X { // CXX14: @_ZN1XIiE1aE = available_externally {{.*}}constant i32 123 - // CXX17: @_ZN1XIiE1aE = {{.*}}externally_initialized constant i32 123 + // CXX17: @_ZN1XIiE1aE = {{.*}}constant i32 123 static constexpr int a = 123; }; __constant__ const int &use_X_a = X<int>::a; template <typename T, T a, T b> struct A { // CXX14: @_ZN1AIiLi1ELi2EE1xE = available_externally {{.*}}constant i32 2 - // CXX17: @_ZN1AIiLi1ELi2EE1xE = {{.*}}externally_initialized constant i32 2 + // CXX17: @_ZN1AIiLi1ELi2EE1xE = {{.*}}constant i32 2 constexpr static T x = a * b; }; __constant__ const int &y = A<int, 1, 2>::x; diff --git a/clang/test/CodeGenCUDA/host-used-device-var.cu b/clang/test/CodeGenCUDA/host-used-device-var.cu index 5328660c9dc9d..7470021c62bed 100644 --- a/clang/test/CodeGenCUDA/host-used-device-var.cu +++ b/clang/test/CodeGenCUDA/host-used-device-var.cu @@ -74,8 +74,8 @@ inline constexpr int constexpr_var1b = 1; // Check constant constexpr variables ODR-used by host code only. // Device-side constexpr variables accessed by host code should be externalized and kept. -// DEV-DAG: @_ZL15constexpr_var2a = addrspace(4) externally_initialized constant i32 2 -// DEV-DAG: @constexpr_var2b = linkonce_odr addrspace(4) externally_initialized constant i32 2 +// DEV-DAG: @_ZL15constexpr_var2a = addrspace(4) constant i32 2 +// DEV-DAG: @constexpr_var2b = linkonce_odr addrspace(4) constant i32 2 __constant__ constexpr int constexpr_var2a = 2; inline __constant__ constexpr int constexpr_var2b = 2; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
