[clang] [Clang] Swap range metadata to attribute for intrinsics. (PR #94851)
Andreas Jonson via cfe-commits
cfe-commits at lists.llvm.org
Tue Jun 11 12:11:57 PDT 2024
https://github.com/andjo403 updated https://github.com/llvm/llvm-project/pull/94851
>From 6b4556ea373d90a780c132ab2c51ae46d40a3f49 Mon Sep 17 00:00:00 2001
From: Andreas Jonson <andjo403 at hotmail.com>
Date: Sun, 26 May 2024 10:02:25 +0200
Subject: [PATCH 1/2] [Clang] swap range metadata to attribute for Intrinsics.
---
clang/lib/CodeGen/CGBuiltin.cpp | 18 ++++++++----------
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 7 +++----
clang/test/CodeGenOpenCL/builtins-r600.cl | 7 +++----
3 files changed, 14 insertions(+), 18 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index c16b69ba87567..e053e1a46cbb1 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -734,17 +734,15 @@ static llvm::Value *EmitOverflowIntrinsic(CodeGenFunction &CGF,
return CGF.Builder.CreateExtractValue(Tmp, 0);
}
-static Value *emitRangedBuiltin(CodeGenFunction &CGF,
- unsigned IntrinsicID,
+static Value *emitRangedBuiltin(CodeGenFunction &CGF, unsigned IntrinsicID,
int low, int high) {
- llvm::MDBuilder MDHelper(CGF.getLLVMContext());
- llvm::MDNode *RNode = MDHelper.createRange(APInt(32, low), APInt(32, high));
- Function *F = CGF.CGM.getIntrinsic(IntrinsicID, {});
- llvm::Instruction *Call = CGF.Builder.CreateCall(F);
- Call->setMetadata(llvm::LLVMContext::MD_range, RNode);
- Call->setMetadata(llvm::LLVMContext::MD_noundef,
- llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt));
- return Call;
+ Function *F = CGF.CGM.getIntrinsic(IntrinsicID, {});
+ llvm::CallInst *Call = CGF.Builder.CreateCall(F);
+ llvm::ConstantRange CR(APInt(32, low), APInt(32, high));
+ Call->addRangeRetAttr(CR);
+ Call->setMetadata(llvm::LLVMContext::MD_noundef,
+ llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt));
+ return Call;
}
namespace {
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index ffc190b76db98..121d1f15449e3 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -605,9 +605,9 @@ void test_s_getreg(volatile global uint *out)
}
// CHECK-LABEL: @test_get_local_id(
-// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workitem.id.x(), !range [[$WI_RANGE:![0-9]*]], !noundef
-// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workitem.id.y(), !range [[$WI_RANGE]], !noundef
-// CHECK: tail call{{.*}} i32 @llvm.amdgcn.workitem.id.z(), !range [[$WI_RANGE]], !noundef
+// CHECK: tail call range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.x(), !noundef
+// CHECK: tail call range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.y(), !noundef
+// CHECK: tail call range(i32 0, 1024{{.*}}) i32 @llvm.amdgcn.workitem.id.z(), !noundef
void test_get_local_id(int d, global int *out)
{
switch (d) {
@@ -891,6 +891,5 @@ void test_set_fpenv(unsigned long env) {
__builtin_amdgcn_set_fpenv(env);
}
-// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
// CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }
diff --git a/clang/test/CodeGenOpenCL/builtins-r600.cl b/clang/test/CodeGenOpenCL/builtins-r600.cl
index c6b40f079b3f2..e9d60be2ee444 100644
--- a/clang/test/CodeGenOpenCL/builtins-r600.cl
+++ b/clang/test/CodeGenOpenCL/builtins-r600.cl
@@ -39,9 +39,9 @@ void test_get_group_id(int d, global int *out)
}
// CHECK-LABEL: @test_get_local_id(
-// CHECK: tail call i32 @llvm.r600.read.tidig.x(), !range [[WI_RANGE:![0-9]*]]
-// CHECK: tail call i32 @llvm.r600.read.tidig.y(), !range [[WI_RANGE]]
-// CHECK: tail call i32 @llvm.r600.read.tidig.z(), !range [[WI_RANGE]]
+// CHECK: tail call range(i32 0, 1024) i32 @llvm.r600.read.tidig.x()
+// CHECK: tail call range(i32 0, 1024) i32 @llvm.r600.read.tidig.y()
+// CHECK: tail call range(i32 0, 1024) i32 @llvm.r600.read.tidig.z()
void test_get_local_id(int d, global int *out)
{
switch (d) {
@@ -52,4 +52,3 @@ void test_get_local_id(int d, global int *out)
}
}
-// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
>From 10f061528a6338be14b6aaad5ab3a17f880415c0 Mon Sep 17 00:00:00 2001
From: Andreas Jonson <andjo403 at hotmail.com>
Date: Tue, 11 Jun 2024 20:38:27 +0200
Subject: [PATCH 2/2] [clang] Swap NoUndef metadata to return attribute.
---
clang/lib/CodeGen/CGBuiltin.cpp | 3 +--
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 6 +++---
clang/test/CodeGenOpenCL/builtins-r600.cl | 6 +++---
3 files changed, 7 insertions(+), 8 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index e053e1a46cbb1..f21a8f1e7014a 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -740,8 +740,7 @@ static Value *emitRangedBuiltin(CodeGenFunction &CGF, unsigned IntrinsicID,
llvm::CallInst *Call = CGF.Builder.CreateCall(F);
llvm::ConstantRange CR(APInt(32, low), APInt(32, high));
Call->addRangeRetAttr(CR);
- Call->setMetadata(llvm::LLVMContext::MD_noundef,
- llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt));
+ Call->addRetAttr(llvm::Attribute::AttrKind::NoUndef);
return Call;
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 121d1f15449e3..95daa2cdbc92c 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -605,9 +605,9 @@ void test_s_getreg(volatile global uint *out)
}
// CHECK-LABEL: @test_get_local_id(
-// CHECK: tail call range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.x(), !noundef
-// CHECK: tail call range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.y(), !noundef
-// CHECK: tail call range(i32 0, 1024{{.*}}) i32 @llvm.amdgcn.workitem.id.z(), !noundef
+// CHECK: tail call noundef range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.x()
+// CHECK: tail call noundef range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.y()
+// CHECK: tail call noundef range(i32 0, 1024){{.*}} i32 @llvm.amdgcn.workitem.id.z()
void test_get_local_id(int d, global int *out)
{
switch (d) {
diff --git a/clang/test/CodeGenOpenCL/builtins-r600.cl b/clang/test/CodeGenOpenCL/builtins-r600.cl
index e9d60be2ee444..a82c4fb90ec50 100644
--- a/clang/test/CodeGenOpenCL/builtins-r600.cl
+++ b/clang/test/CodeGenOpenCL/builtins-r600.cl
@@ -39,9 +39,9 @@ void test_get_group_id(int d, global int *out)
}
// CHECK-LABEL: @test_get_local_id(
-// CHECK: tail call range(i32 0, 1024) i32 @llvm.r600.read.tidig.x()
-// CHECK: tail call range(i32 0, 1024) i32 @llvm.r600.read.tidig.y()
-// CHECK: tail call range(i32 0, 1024) i32 @llvm.r600.read.tidig.z()
+// CHECK: tail call noundef range(i32 0, 1024) i32 @llvm.r600.read.tidig.x()
+// CHECK: tail call noundef range(i32 0, 1024) i32 @llvm.r600.read.tidig.y()
+// CHECK: tail call noundef range(i32 0, 1024) i32 @llvm.r600.read.tidig.z()
void test_get_local_id(int d, global int *out)
{
switch (d) {
More information about the cfe-commits
mailing list