Author: Nikita Popov Date: 2023-01-12T10:03:05+01:00 New Revision: 02856565ac12e21f14f2af64ce1135ecc3c2021f
URL: https://github.com/llvm/llvm-project/commit/02856565ac12e21f14f2af64ce1135ecc3c2021f DIFF: https://github.com/llvm/llvm-project/commit/02856565ac12e21f14f2af64ce1135ecc3c2021f.diff LOG: [Clang] Emit noundef metadata next to range metadata To preserve the previous semantics after D141386, adjust places that currently emit !range metadata to also emit !noundef metadata. This retains range violation as immediate undefined behavior, rather than just poison. Differential Revision: https://reviews.llvm.org/D141494 Added: Modified: clang/lib/CodeGen/CGBuiltin.cpp clang/lib/CodeGen/CGExpr.cpp clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp clang/test/CodeGenCXX/pr12251.cpp clang/test/CodeGenOpenCL/builtins-amdgcn.cl Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 430b5f43cdd5a..479e24245df4d 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -687,6 +687,8 @@ static Value *emitRangedBuiltin(CodeGenFunction &CGF, 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; } @@ -16785,6 +16787,8 @@ Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) { llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1), APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1)); LD->setMetadata(llvm::LLVMContext::MD_range, RNode); + LD->setMetadata(llvm::LLVMContext::MD_noundef, + llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt)); LD->setMetadata(llvm::LLVMContext::MD_invariant_load, llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt)); return LD; diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index 6d5e729b1eea9..34974c63984e6 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -1751,8 +1751,11 @@ llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile, // In order to prevent the optimizer from throwing away the check, don't // attach range metadata to the load. } else if (CGM.getCodeGenOpts().OptimizationLevel > 0) - if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty)) + if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty)) { Load->setMetadata(llvm::LLVMContext::MD_range, RangeInfo); + Load->setMetadata(llvm::LLVMContext::MD_noundef, + llvm::MDNode::get(getLLVMContext(), std::nullopt)); + } return EmitFromMemory(Load, Ty); } diff --git a/clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu b/clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu index e506b875b6748..c098917a0a0e2 100644 --- a/clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu +++ b/clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu @@ -12,20 +12,20 @@ // PRECOV5-LABEL: test_get_workgroup_size // PRECOV5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4 -// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6 -// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8 -// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // COV5-LABEL: test_get_workgroup_size // COV5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12 -// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14 -// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16 -// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef __device__ void test_get_workgroup_size(int d, int *out) { switch (d) { diff --git a/clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp b/clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp index 4f0009e290566..2583a4a84c3c3 100644 --- a/clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp +++ b/clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp @@ -8,7 +8,7 @@ extern bool B(); bool f() { // CHECK-LABEL: define{{.*}} zeroext i1 @_Z1fv - // CHECK: br {{.*}} !prof !7 + // CHECK: br {{.*}} !prof ![[PROF_LIKELY:[0-9]+]] if (b) [[likely]] { return A(); @@ -18,7 +18,7 @@ bool f() { bool g() { // CHECK-LABEL: define{{.*}} zeroext i1 @_Z1gv - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY:[0-9]+]] if (b) [[unlikely]] { return A(); @@ -29,7 +29,7 @@ bool g() { bool h() { // CHECK-LABEL: define{{.*}} zeroext i1 @_Z1hv - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] return A(); @@ -38,7 +38,7 @@ bool h() { void NullStmt() { // CHECK-LABEL: define{{.*}}NullStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]]; else { @@ -49,7 +49,7 @@ void NullStmt() { void IfStmt() { // CHECK-LABEL: define{{.*}}IfStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] if (B()) {} @@ -63,20 +63,20 @@ void IfStmt() { void WhileStmt() { // CHECK-LABEL: define{{.*}}WhileStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] while (B()) {} // CHECK-NOT: br {{.*}} %if.end{{.*}} !prof if (b) - // CHECK: br {{.*}} !prof !7 + // CHECK: br {{.*}} !prof ![[PROF_LIKELY]] while (B()) [[unlikely]] { b = false; } } void DoStmt() { // CHECK-LABEL: define{{.*}}DoStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] do {} while (B()) @@ -91,20 +91,20 @@ void DoStmt() { void ForStmt() { // CHECK-LABEL: define{{.*}}ForStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] for (; B();) {} // CHECK-NOT: br {{.*}} %if.end{{.*}} !prof if (b) - // CHECK: br {{.*}} !prof !7 + // CHECK: br {{.*}} !prof ![[PROF_LIKELY]] for (; B();) [[unlikely]] {} } void GotoStmt() { // CHECK-LABEL: define{{.*}}GotoStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] goto end; else { @@ -116,7 +116,7 @@ end:; void ReturnStmt() { // CHECK-LABEL: define{{.*}}ReturnStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] return; else { @@ -127,7 +127,7 @@ void ReturnStmt() { void SwitchStmt() { // CHECK-LABEL: define{{.*}}SwitchStmt - // CHECK: br {{.*}} !prof !8 + // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]] if (b) [[unlikely]] switch (i) {} else { @@ -144,5 +144,5 @@ void SwitchStmt() { } } -// CHECK: !7 = !{!"branch_weights", i32 [[UNLIKELY]], i32 [[LIKELY]]} -// CHECK: !8 = !{!"branch_weights", i32 [[LIKELY]], i32 [[UNLIKELY]]} +// CHECK: ![[PROF_LIKELY]] = !{!"branch_weights", i32 [[UNLIKELY]], i32 [[LIKELY]]} +// CHECK: ![[PROF_UNLIKELY]] = !{!"branch_weights", i32 [[LIKELY]], i32 [[UNLIKELY]]} diff --git a/clang/test/CodeGenCXX/pr12251.cpp b/clang/test/CodeGenCXX/pr12251.cpp index a267a3aed077d..bd5c85b83f2ca 100644 --- a/clang/test/CodeGenCXX/pr12251.cpp +++ b/clang/test/CodeGenCXX/pr12251.cpp @@ -5,7 +5,7 @@ bool f(bool *x) { return *x; } // CHECK-LABEL: define{{.*}} zeroext i1 @_Z1fPb -// CHECK: load i8, ptr %{{[^ ]*}}, align 1, !range [[RANGE_i8_0_2:![^ ]*]] +// CHECK: load i8, ptr %{{[^ ]*}}, align 1, !range [[RANGE_i8_0_2:![0-9]+]], !noundef [[NOUNDEF:![0-9]+]] // Only enum-tests follow. Ensure that after the bool test, no further range // metadata shows up when strict enums are disabled. @@ -32,63 +32,63 @@ e3 g3(e3 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g3P2e3 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32:![^ ]*]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32:![0-9]+]], !noundef [[NOUNDEF]] enum e4 { e4_a = -16}; e4 g4(e4 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g4P2e4 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16:![^ ]*]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16:![0-9]+]], !noundef [[NOUNDEF]] enum e5 { e5_a = -16, e5_b = 16}; e5 g5(e5 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g5P2e5 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![^ ]*]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![0-9]+]], !noundef [[NOUNDEF]] enum e6 { e6_a = -1 }; e6 g6(e6 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g6P2e6 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m1_1:![^ ]*]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m1_1:![0-9]+]], !noundef [[NOUNDEF]] enum e7 { e7_a = -16, e7_b = 2}; e7 g7(e7 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g7P2e7 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16]], !noundef [[NOUNDEF]] enum e8 { e8_a = -17}; e8 g8(e8 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g8P2e8 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![^ ]*]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![0-9]+]], !noundef [[NOUNDEF]] enum e9 { e9_a = 17}; e9 g9(e9 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z2g9P2e9 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32]], !noundef [[NOUNDEF]] enum e10 { e10_a = -16, e10_b = 32}; e10 g10(e10 *x) { return *x; } // CHECK-LABEL: define{{.*}} i32 @_Z3g10P3e10 -// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m64_64:![^ ]*]] +// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m64_64:![0-9]+]], !noundef [[NOUNDEF]] enum e11 {e11_a = 4294967296 }; enum e11 g11(enum e11 *x) { return *x; } // CHECK-LABEL: define{{.*}} i64 @_Z3g11P3e11 -// CHECK: load i64, ptr %x, align {{[84]}}, !range [[RANGE_i64_0_2pow33:![^ ]*]] +// CHECK: load i64, ptr %x, align {{[84]}}, !range [[RANGE_i64_0_2pow33:![0-9]+]], !noundef [[NOUNDEF]] enum e12 {e12_a = 9223372036854775808U }; enum e12 g12(enum e12 *x) { @@ -137,6 +137,7 @@ e16 g16(e16 *x) { // CHECK: [[RANGE_i8_0_2]] = !{i8 0, i8 2} +// CHECK: [[NOUNDEF]] = !{} // CHECK: [[RANGE_i32_0_32]] = !{i32 0, i32 32} // CHECK: [[RANGE_i32_m16_16]] = !{i32 -16, i32 16} // CHECK: [[RANGE_i32_m32_32]] = !{i32 -32, i32 32} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 82cd3177d6c6d..094851b218898 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -569,9 +569,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]*]] -// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[$WI_RANGE]] -// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[$WI_RANGE]] +// 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 void test_get_local_id(int d, global int *out) { switch (d) { @@ -585,11 +585,11 @@ void test_get_local_id(int d, global int *out) // CHECK-LABEL: @test_get_workgroup_size( // CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 4 -// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 6 -// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 8 -// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load +// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef void test_get_workgroup_size(int d, global int *out) { switch (d) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
