Author: Andreas Jonson Date: 2024-06-19T17:23:53+02:00 New Revision: 01ba3fa37b22828abca81011b0cb2e4cbbb8cfda
URL: https://github.com/llvm/llvm-project/commit/01ba3fa37b22828abca81011b0cb2e4cbbb8cfda DIFF: https://github.com/llvm/llvm-project/commit/01ba3fa37b22828abca81011b0cb2e4cbbb8cfda.diff LOG: [Clang] Swap range and noundef metadata to attribute for intrinsics. (#94851) Added: Modified: clang/lib/CodeGen/CGBuiltin.cpp clang/test/CodeGenOpenCL/builtins-amdgcn.cl clang/test/CodeGenOpenCL/builtins-r600.cl Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 08a89bd123d03..4f393ac1d1805 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -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 { diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index ffc190b76db98..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{{.*}} 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) } diff --git a/clang/test/CodeGenOpenCL/builtins-r600.cl b/clang/test/CodeGenOpenCL/builtins-r600.cl index c6b40f079b3f2..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 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} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
