https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/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. >From 932e13d1b8eceff2b9523ea54955972891664d80 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Wed, 18 Feb 2026 15:51:25 -0600 Subject: [PATCH] [HIP] Do not apply 'externally_initialized' to constant device variables 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. --- clang/lib/CodeGen/CodeGenModule.cpp | 3 ++- clang/test/CodeGenCUDA/const-var.cu | 4 ++-- clang/test/CodeGenCUDA/constexpr-variables.cu | 8 ++++---- clang/test/CodeGenCUDA/host-used-device-var.cu | 4 ++-- 4 files changed, 10 insertions(+), 9 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 43b8af0b2156a..b7127f4f65e22 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6226,7 +6226,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
