This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG02856565ac12: [Clang] Emit noundef metadata next to range
metadata (authored by nikic).
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D141494/new/
https://reviews.llvm.org/D141494
Files:
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
Index: clang/test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -569,9 +569,9 @@
}
// 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 @@
// 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) {
Index: clang/test/CodeGenCXX/pr12251.cpp
===================================================================
--- clang/test/CodeGenCXX/pr12251.cpp
+++ clang/test/CodeGenCXX/pr12251.cpp
@@ -5,7 +5,7 @@
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 @@
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 @@
// 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}
Index: clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
===================================================================
--- clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
+++ clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
@@ -8,7 +8,7 @@
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 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 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 @@
void NullStmt() {
// CHECK-LABEL: define{{.*}}NullStmt
- // CHECK: br {{.*}} !prof !8
+ // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
if (b)
[[unlikely]];
else {
@@ -49,7 +49,7 @@
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 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 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 @@
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 SwitchStmt() {
// CHECK-LABEL: define{{.*}}SwitchStmt
- // CHECK: br {{.*}} !prof !8
+ // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
if (b)
[[unlikely]] switch (i) {}
else {
@@ -144,5 +144,5 @@
}
}
-// 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]]}
Index: clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
+++ 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) {
Index: clang/lib/CodeGen/CGExpr.cpp
===================================================================
--- clang/lib/CodeGen/CGExpr.cpp
+++ clang/lib/CodeGen/CGExpr.cpp
@@ -1751,8 +1751,11 @@
// 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);
}
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -687,6 +687,8 @@
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 @@
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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
