https://github.com/keshavvinayak01 updated https://github.com/llvm/llvm-project/pull/218356
>From 99438e164b8977db5bbb0eed3077bf71290fcefd Mon Sep 17 00:00:00 2001 From: Keshav Vinayak Jha <[email protected]> Date: Mon, 24 Aug 2026 09:11:24 +0000 Subject: [PATCH 1/3] [Clang][OpenCL] Mark constant address space loads invariant --- clang/lib/CodeGen/CGExpr.cpp | 3 +++ clang/test/CodeGenOpenCL/invariant-load.cl | 24 ++++++++++++++++++++++ 2 files changed, 27 insertions(+) create mode 100644 clang/test/CodeGenOpenCL/invariant-load.cl diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index eff6a7de320d7..ef66b9f0d4af7 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -2248,6 +2248,9 @@ llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile, Addr.withElementType(convertTypeForLoadStore(Ty, Addr.getElementType())); llvm::LoadInst *Load = Builder.CreateLoad(Addr, Volatile); + if (Ty.getAddressSpace() == LangAS::opencl_constant) + Load->setMetadata(llvm::LLVMContext::MD_invariant_load, + llvm::MDNode::get(Load->getContext(), {})); if (isNontemporal) { llvm::MDNode *Node = llvm::MDNode::get( Load->getContext(), llvm::ConstantAsMetadata::get(Builder.getInt32(1))); diff --git a/clang/test/CodeGenOpenCL/invariant-load.cl b/clang/test/CodeGenOpenCL/invariant-load.cl new file mode 100644 index 0000000000000..717348398e28e --- /dev/null +++ b/clang/test/CodeGenOpenCL/invariant-load.cl @@ -0,0 +1,24 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -O0 -emit-llvm -o - %s | FileCheck %s --check-prefix=AMDGCN +// RUN: %clang_cc1 -triple spir64-unknown-unknown -cl-std=CL2.0 -O0 -emit-llvm -o - %s | FileCheck %s --check-prefix=SPIR + +kernel void constant_load(global int *out, constant int *in) { + out[0] = in[0]; +} + +// AMDGCN-LABEL: define{{.*}}@constant_load( +// AMDGCN: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load [[INVARIANT:![0-9]+]] +// SPIR-LABEL: define{{.*}}@constant_load( +// SPIR: load i32, ptr addrspace(2) %{{.*}}, align 4, !invariant.load [[INVARIANT:![0-9]+]] + +kernel void global_const_load(global int *out, global const int *in) { + out[0] = in[0]; +} + +// AMDGCN-LABEL: define{{.*}}@global_const_load( +// AMDGCN: load i32, ptr addrspace(1) %{{.*}}, align 4{{$}} +// SPIR-LABEL: define{{.*}}@global_const_load( +// SPIR: load i32, ptr addrspace(1) %{{.*}}, align 4{{$}} + +// AMDGCN: [[INVARIANT]] = !{} +// SPIR: [[INVARIANT]] = !{} >From c892aff09c611c92fb29b9a3eceb88a62c1ef818 Mon Sep 17 00:00:00 2001 From: Keshav Vinayak Jha <[email protected]> Date: Mon, 7 Sep 2026 08:17:13 +0000 Subject: [PATCH 2/3] [Clang][CUDA] Mark constant memory loads invariant Recover the underlying global's target address space for CUDA and HIP device loads, whose source types remain in the default address space. Add CUDA coverage and simplify the OpenCL test as requested in review. --- clang/lib/CodeGen/CGExpr.cpp | 19 +++++++++- clang/test/CodeGenCUDA/invariant-load.cu | 43 ++++++++++++++++++++++ clang/test/CodeGenOpenCL/invariant-load.cl | 19 +++------- 3 files changed, 67 insertions(+), 14 deletions(-) create mode 100644 clang/test/CodeGenCUDA/invariant-load.cu diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index ef66b9f0d4af7..845db06a4a0dd 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -42,6 +42,7 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/ScopeExit.h" #include "llvm/ADT/StringExtras.h" +#include "llvm/Analysis/ValueTracking.h" #include "llvm/IR/Constants.h" #include "llvm/IR/DataLayout.h" #include "llvm/IR/Intrinsics.h" @@ -2074,6 +2075,22 @@ llvm::Value *CodeGenFunction::emitScalarConstant( return Constant.getValue(); } +static bool isInvariantLoad(CodeGenFunction &CGF, LangAS TypeAS, + llvm::Value *Ptr) { + if (TypeAS == LangAS::opencl_constant) + return true; + + // CUDA represents constant memory as a declaration attribute, so the + // expression type uses the default address space. Recover the storage + // address space from the underlying global after the address-space cast. + if (!CGF.getLangOpts().CUDAIsDevice) + return false; + const auto *GV = + llvm::dyn_cast<llvm::GlobalValue>(llvm::getUnderlyingObject(Ptr)); + return GV && GV->getAddressSpace() == CGF.getContext().getTargetAddressSpace( + LangAS::cuda_constant); +} + llvm::Value *CodeGenFunction::EmitLoadOfScalar(LValue lvalue, SourceLocation Loc) { return EmitLoadOfScalar(lvalue.getAddress(), lvalue.isVolatile(), @@ -2248,7 +2265,7 @@ llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile, Addr.withElementType(convertTypeForLoadStore(Ty, Addr.getElementType())); llvm::LoadInst *Load = Builder.CreateLoad(Addr, Volatile); - if (Ty.getAddressSpace() == LangAS::opencl_constant) + if (isInvariantLoad(*this, Ty.getAddressSpace(), Addr.getBasePointer())) Load->setMetadata(llvm::LLVMContext::MD_invariant_load, llvm::MDNode::get(Load->getContext(), {})); if (isNontemporal) { diff --git a/clang/test/CodeGenCUDA/invariant-load.cu b/clang/test/CodeGenCUDA/invariant-load.cu new file mode 100644 index 0000000000000..12685405f9a2a --- /dev/null +++ b/clang/test/CodeGenCUDA/invariant-load.cu @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -triple amdgpu7.00-amd-amdhsa -fcuda-is-device -emit-llvm -o - %s | FileCheck %s + +#include "Inputs/cuda.h" + +__constant__ int constant_value; +__constant__ int constant_array[4]; +__device__ int device_value; + +struct S { + int member; +}; + +__constant__ S constant_struct; + +__device__ int constant_load() { + return constant_value; +} + +// CHECK-LABEL: define{{.*}}@_Z13constant_loadv( +// CHECK: load i32, ptr addrspacecast (ptr addrspace(4) @constant_value to ptr), align 4, !invariant.load [[INVARIANT:![0-9]+]] + +__device__ int constant_array_load(int index) { + return constant_array[index]; +} + +// CHECK-LABEL: define{{.*}}@_Z19constant_array_loadi( +// CHECK: load i32, ptr %{{.*}}, align 4, !invariant.load [[INVARIANT]] + +__device__ int constant_member_load() { + return constant_struct.member; +} + +// CHECK-LABEL: define{{.*}}@_Z20constant_member_loadv( +// CHECK: load i32, ptr addrspacecast (ptr addrspace(4) @constant_struct to ptr), align 4, !invariant.load [[INVARIANT]] + +__device__ int device_load() { + return device_value; +} + +// CHECK-LABEL: define{{.*}}@_Z11device_loadv( +// CHECK: load i32, ptr addrspacecast (ptr addrspace(1) @device_value to ptr), align 4{{$}} + +// CHECK: [[INVARIANT]] = !{} diff --git a/clang/test/CodeGenOpenCL/invariant-load.cl b/clang/test/CodeGenOpenCL/invariant-load.cl index 717348398e28e..df26f6c4aa495 100644 --- a/clang/test/CodeGenOpenCL/invariant-load.cl +++ b/clang/test/CodeGenOpenCL/invariant-load.cl @@ -1,24 +1,17 @@ -// REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -O0 -emit-llvm -o - %s | FileCheck %s --check-prefix=AMDGCN -// RUN: %clang_cc1 -triple spir64-unknown-unknown -cl-std=CL2.0 -O0 -emit-llvm -o - %s | FileCheck %s --check-prefix=SPIR +// RUN: %clang_cc1 -triple amdgpu7.00-amd-amdhsa -cl-std=CL2.0 -O0 -emit-llvm -o - %s | FileCheck %s kernel void constant_load(global int *out, constant int *in) { out[0] = in[0]; } -// AMDGCN-LABEL: define{{.*}}@constant_load( -// AMDGCN: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load [[INVARIANT:![0-9]+]] -// SPIR-LABEL: define{{.*}}@constant_load( -// SPIR: load i32, ptr addrspace(2) %{{.*}}, align 4, !invariant.load [[INVARIANT:![0-9]+]] +// CHECK-LABEL: define{{.*}}@constant_load( +// CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load [[INVARIANT:![0-9]+]] kernel void global_const_load(global int *out, global const int *in) { out[0] = in[0]; } -// AMDGCN-LABEL: define{{.*}}@global_const_load( -// AMDGCN: load i32, ptr addrspace(1) %{{.*}}, align 4{{$}} -// SPIR-LABEL: define{{.*}}@global_const_load( -// SPIR: load i32, ptr addrspace(1) %{{.*}}, align 4{{$}} +// CHECK-LABEL: define{{.*}}@global_const_load( +// CHECK: load i32, ptr addrspace(1) %{{.*}}, align 4{{$}} -// AMDGCN: [[INVARIANT]] = !{} -// SPIR: [[INVARIANT]] = !{} +// CHECK: [[INVARIANT]] = !{} >From b856f1a40188b4a5d2f472bd2068d964f6ec7bfe Mon Sep 17 00:00:00 2001 From: Keshav Vinayak Jha <[email protected]> Date: Thu, 10 Sep 2026 10:55:48 +0000 Subject: [PATCH 3/3] [Clang][CUDA] Track invariant loads from AST Record CUDA constant storage on LValueBaseInfo while the VarDecl is available, and propagate it through fields and conditional lvalues. This avoids recovering pointer provenance from LLVM IR. Co-authored-by: GPT-5 <[email protected]> Signed-off-by: Keshav Vinayak Jha <[email protected]> --- clang/lib/CodeGen/CGExpr.cpp | 42 ++++++++++-------------- clang/lib/CodeGen/CGValue.h | 9 ++++- clang/test/CodeGenCUDA/invariant-load.cu | 24 ++++++++++++++ 3 files changed, 49 insertions(+), 26 deletions(-) diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index 845db06a4a0dd..c49e5cf608e51 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -42,7 +42,6 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/ScopeExit.h" #include "llvm/ADT/StringExtras.h" -#include "llvm/Analysis/ValueTracking.h" #include "llvm/IR/Constants.h" #include "llvm/IR/DataLayout.h" #include "llvm/IR/Intrinsics.h" @@ -2075,22 +2074,6 @@ llvm::Value *CodeGenFunction::emitScalarConstant( return Constant.getValue(); } -static bool isInvariantLoad(CodeGenFunction &CGF, LangAS TypeAS, - llvm::Value *Ptr) { - if (TypeAS == LangAS::opencl_constant) - return true; - - // CUDA represents constant memory as a declaration attribute, so the - // expression type uses the default address space. Recover the storage - // address space from the underlying global after the address-space cast. - if (!CGF.getLangOpts().CUDAIsDevice) - return false; - const auto *GV = - llvm::dyn_cast<llvm::GlobalValue>(llvm::getUnderlyingObject(Ptr)); - return GV && GV->getAddressSpace() == CGF.getContext().getTargetAddressSpace( - LangAS::cuda_constant); -} - llvm::Value *CodeGenFunction::EmitLoadOfScalar(LValue lvalue, SourceLocation Loc) { return EmitLoadOfScalar(lvalue.getAddress(), lvalue.isVolatile(), @@ -2265,7 +2248,7 @@ llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile, Addr.withElementType(convertTypeForLoadStore(Ty, Addr.getElementType())); llvm::LoadInst *Load = Builder.CreateLoad(Addr, Volatile); - if (isInvariantLoad(*this, Ty.getAddressSpace(), Addr.getBasePointer())) + if (Ty.getAddressSpace() == LangAS::opencl_constant || BaseInfo.isInvariant()) Load->setMetadata(llvm::LLVMContext::MD_invariant_load, llvm::MDNode::get(Load->getContext(), {})); if (isNontemporal) { @@ -3520,10 +3503,14 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction &CGF, return EmitThreadPrivateVarDeclLValue(CGF, VD, T, Addr, RealVarTy, E->getExprLoc()); } - LValue LV = VD->getType()->isReferenceType() ? - CGF.EmitLoadOfReferenceLValue(Addr, VD->getType(), - AlignmentSource::Decl) : - CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl); + const bool IsReference = VD->getType()->isReferenceType(); + LValue LV = IsReference ? CGF.EmitLoadOfReferenceLValue(Addr, VD->getType(), + AlignmentSource::Decl) + : CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl); + // Preserve CUDA constant storage information from the AST declaration. + if (!IsReference && CGF.getLangOpts().CUDAIsDevice && + CGF.CGM.GetGlobalVarAddressSpace(VD) == LangAS::cuda_constant) + LV.setInvariant(true); setObjCGCLValueClass(CGF.getContext(), E, LV); return LV; } @@ -5875,7 +5862,9 @@ LValue CodeGenFunction::EmitLValueForField(LValue base, const FieldDecl *field, QualType FieldType = field->getType(); const RecordDecl *rec = field->getParent(); AlignmentSource BaseAlignSource = BaseInfo.getAlignmentSource(); - LValueBaseInfo FieldBaseInfo(getFieldAlignmentSource(BaseAlignSource)); + // A field inherits invariant storage from its base object. + LValueBaseInfo FieldBaseInfo = BaseInfo; + FieldBaseInfo.setAlignmentSource(getFieldAlignmentSource(BaseAlignSource)); TBAAAccessInfo FieldTBAAInfo; if (base.getTBAAInfo().isMayAlias() || rec->hasAttr<MayAliasAttr>() || FieldType->isVectorType()) { @@ -6191,8 +6180,11 @@ LValue CodeGenFunction::EmitConditionalOperatorLValue( Info.RHS->getBaseInfo().getAlignmentSource()); TBAAAccessInfo TBAAInfo = CGM.mergeTBAAInfoForConditionalOperator( Info.LHS->getTBAAInfo(), Info.RHS->getTBAAInfo()); - return MakeAddrLValue(result, expr->getType(), LValueBaseInfo(alignSource), - TBAAInfo); + LValueBaseInfo BaseInfo(alignSource); + // Both possible storage locations must be invariant. + BaseInfo.setInvariant(Info.LHS->getBaseInfo().isInvariant() && + Info.RHS->getBaseInfo().isInvariant()); + return MakeAddrLValue(result, expr->getType(), BaseInfo, TBAAInfo); } else { assert((Info.LHS || Info.RHS) && "both operands of glvalue conditional are throw-expressions?"); diff --git a/clang/lib/CodeGen/CGValue.h b/clang/lib/CodeGen/CGValue.h index 118ed6690627c..6eb9ba182fd3e 100644 --- a/clang/lib/CodeGen/CGValue.h +++ b/clang/lib/CodeGen/CGValue.h @@ -166,12 +166,18 @@ static inline AlignmentSource getFieldAlignmentSource(AlignmentSource Source) { class LValueBaseInfo { AlignmentSource AlignSource; + // Whether loads from the base object's storage are invariant. + bool IsInvariant : 1; + public: explicit LValueBaseInfo(AlignmentSource Source = AlignmentSource::Type) - : AlignSource(Source) {} + : AlignSource(Source), IsInvariant(false) {} AlignmentSource getAlignmentSource() const { return AlignSource; } void setAlignmentSource(AlignmentSource Source) { AlignSource = Source; } + bool isInvariant() const { return IsInvariant; } + void setInvariant(bool Value) { IsInvariant = Value; } + void mergeForCast(const LValueBaseInfo &Info) { setAlignmentSource(Info.getAlignmentSource()); } @@ -357,6 +363,7 @@ class LValue { LValueBaseInfo getBaseInfo() const { return BaseInfo; } void setBaseInfo(LValueBaseInfo Info) { BaseInfo = Info; } + void setInvariant(bool Value) { BaseInfo.setInvariant(Value); } KnownNonNull_t isKnownNonNull() const { return Addr.isKnownNonNull(); } LValue setKnownNonNull() { diff --git a/clang/test/CodeGenCUDA/invariant-load.cu b/clang/test/CodeGenCUDA/invariant-load.cu index 12685405f9a2a..308eb94615ea1 100644 --- a/clang/test/CodeGenCUDA/invariant-load.cu +++ b/clang/test/CodeGenCUDA/invariant-load.cu @@ -3,7 +3,9 @@ #include "Inputs/cuda.h" __constant__ int constant_value; +__constant__ int other_constant_value; __constant__ int constant_array[4]; +__constant__ int *constant_pointer; __device__ int device_value; struct S { @@ -40,4 +42,26 @@ __device__ int device_load() { // CHECK-LABEL: define{{.*}}@_Z11device_loadv( // CHECK: load i32, ptr addrspacecast (ptr addrspace(1) @device_value to ptr), align 4{{$}} +__device__ int constant_pointer_load() { + return *constant_pointer; +} + +// CHECK-LABEL: define{{.*}}@_Z21constant_pointer_loadv( +// CHECK: load ptr, ptr addrspacecast (ptr addrspace(4) @constant_pointer to ptr), align 8, !invariant.load [[INVARIANT]] +// CHECK-NEXT: load i32, ptr %{{.*}}, align 4{{$}} + +__device__ int mixed_conditional_load(bool select_constant) { + return *&(select_constant ? constant_value : device_value); +} + +// CHECK-LABEL: define{{.*}}@_Z22mixed_conditional_loadb( +// CHECK: load i32, ptr %{{.*}}, align 4{{$}} + +__device__ int constant_conditional_load(bool select_first) { + return *&(select_first ? constant_value : other_constant_value); +} + +// CHECK-LABEL: define{{.*}}@_Z25constant_conditional_loadb( +// CHECK: load i32, ptr %{{.*}}, align 4, !invariant.load [[INVARIANT]] + // CHECK: [[INVARIANT]] = !{} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
