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/2] [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/2] [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]] = !{} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
