Skip to content

[Clang] Swap range metadata to attribute for intrinsics. #94851

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Jun 19, 2024
Merged

[Clang] Swap range metadata to attribute for intrinsics. #94851

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6b4556e
[Clang] swap range metadata to attribute for Intrinsics.
andjo403 May 26, 2024
10f0615
[clang] Swap NoUndef metadata to return attribute.
andjo403 Jun 11, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
17 changes: 7 additions & 10 deletions clang/lib/CodeGen/CGBuiltin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Expand Down
7 changes: 3 additions & 4 deletions clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down Expand Up @@ -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 changes: 3 additions & 4 deletions clang/test/CodeGenOpenCL/builtins-r600.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -52,4 +52,3 @@ void test_get_local_id(int d, global int *out)
}
}

// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
Loading