caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[Clang] Swap range and noundef metadata to attribute for intrinsics. (#94851)

Browse Source
This commit is contained in:
Andreas Jonson
2024-06-19 17:23:53 +02:00
committed by GitHub
parent 881dc132a2
commit 01ba3fa37b
3 changed files with 13 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

17
clang/lib/CodeGen/CGBuiltin.cpp
View File

@@ -734,17 +734,14 @@ 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->addRetAttr(llvm::Attribute::AttrKind::NoUndef);
return Call;
}
namespace {

7
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@@ -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 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) {
@@ -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) }

7
clang/test/CodeGenOpenCL/builtins-r600.cl
View File

@@ -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 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) {
@@ -52,4 +52,3 @@ void test_get_local_id(int d, global int *out)
}
}
// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}