[clang] 0285656 - [Clang] Emit noundef metadata next to range metadata
Nikita Popov via cfe-commits
cfe-commits at lists.llvm.org
Thu Jan 12 01:09:56 PST 2023
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) {
More information about the cfe-commits
mailing list