[clang] [Clang][OpenCL] Mark scalar loads from constant address space invariant (PR #218356)
Keshav Vinayak Jha via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 7 01:23:49 PDT 2026
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 <keshavvinayakjha at gmail.com>
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 <keshavvinayakjha at gmail.com>
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]] = !{}
More information about the cfe-commits
mailing list