Skip to content

Commit 0285656

Browse files
nikicnikic
committed
[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
1 parent 84a5d93 commit 0285656

File tree

6 files changed

+46
-38
lines changed

6 files changed

+46
-38
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp

+4
Original file line numberDiff line numberDiff line change
@@ -687,6 +687,8 @@ static Value *emitRangedBuiltin(CodeGenFunction &CGF,
687687
Function *F = CGF.CGM.getIntrinsic(IntrinsicID, {});
688688
llvm::Instruction *Call = CGF.Builder.CreateCall(F);
689689
Call->setMetadata(llvm::LLVMContext::MD_range, RNode);
690+
Call->setMetadata(llvm::LLVMContext::MD_noundef,
691+
llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt));
690692
return Call;
691693
}
692694

@@ -16785,6 +16787,8 @@ Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
1678516787
llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1),
1678616788
APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));
1678716789
LD->setMetadata(llvm::LLVMContext::MD_range, RNode);
16790+
LD->setMetadata(llvm::LLVMContext::MD_noundef,
16791+
llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt));
1678816792
LD->setMetadata(llvm::LLVMContext::MD_invariant_load,
1678916793
llvm::MDNode::get(CGF.getLLVMContext(), std::nullopt));
1679016794
return LD;

‎clang/lib/CodeGen/CGExpr.cpp

+4-1
Original file line numberDiff line numberDiff line change
@@ -1751,8 +1751,11 @@ llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile,
17511751
// In order to prevent the optimizer from throwing away the check, don't
17521752
// attach range metadata to the load.
17531753
} else if (CGM.getCodeGenOpts().OptimizationLevel > 0)
1754-
if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty))
1754+
if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty)) {
17551755
Load->setMetadata(llvm::LLVMContext::MD_range, RangeInfo);
1756+
Load->setMetadata(llvm::LLVMContext::MD_noundef,
1757+
llvm::MDNode::get(getLLVMContext(), std::nullopt));
1758+
}
17561759

17571760
return EmitFromMemory(Load, Ty);
17581761
}

‎clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu

+6-6
Original file line numberDiff line numberDiff line change
@@ -12,20 +12,20 @@
1212
// PRECOV5-LABEL: test_get_workgroup_size
1313
// PRECOV5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
1414
// PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4
15-
// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
15+
// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
1616
// PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6
17-
// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
17+
// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
1818
// PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
19-
// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
19+
// PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
2020

2121
// COV5-LABEL: test_get_workgroup_size
2222
// COV5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
2323
// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12
24-
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
24+
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
2525
// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14
26-
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
26+
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
2727
// COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
28-
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
28+
// COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
2929
__device__ void test_get_workgroup_size(int d, int *out)
3030
{
3131
switch (d) {

‎clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp

+15-15
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@ extern bool B();
88

99
bool f() {
1010
// CHECK-LABEL: define{{.*}} zeroext i1 @_Z1fv
11-
// CHECK: br {{.*}} !prof !7
11+
// CHECK: br {{.*}} !prof ![[PROF_LIKELY:[0-9]+]]
1212
if (b)
1313
[[likely]] {
1414
return A();
@@ -18,7 +18,7 @@ bool f() {
1818

1919
bool g() {
2020
// CHECK-LABEL: define{{.*}} zeroext i1 @_Z1gv
21-
// CHECK: br {{.*}} !prof !8
21+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY:[0-9]+]]
2222
if (b)
2323
[[unlikely]] {
2424
return A();
@@ -29,7 +29,7 @@ bool g() {
2929

3030
bool h() {
3131
// CHECK-LABEL: define{{.*}} zeroext i1 @_Z1hv
32-
// CHECK: br {{.*}} !prof !8
32+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
3333
if (b)
3434
[[unlikely]] return A();
3535

@@ -38,7 +38,7 @@ bool h() {
3838

3939
void NullStmt() {
4040
// CHECK-LABEL: define{{.*}}NullStmt
41-
// CHECK: br {{.*}} !prof !8
41+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
4242
if (b)
4343
[[unlikely]];
4444
else {
@@ -49,7 +49,7 @@ void NullStmt() {
4949

5050
void IfStmt() {
5151
// CHECK-LABEL: define{{.*}}IfStmt
52-
// CHECK: br {{.*}} !prof !8
52+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
5353
if (b)
5454
[[unlikely]] if (B()) {}
5555

@@ -63,20 +63,20 @@ void IfStmt() {
6363

6464
void WhileStmt() {
6565
// CHECK-LABEL: define{{.*}}WhileStmt
66-
// CHECK: br {{.*}} !prof !8
66+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
6767
if (b)
6868
[[unlikely]] while (B()) {}
6969

7070
// CHECK-NOT: br {{.*}} %if.end{{.*}} !prof
7171
if (b)
72-
// CHECK: br {{.*}} !prof !7
72+
// CHECK: br {{.*}} !prof ![[PROF_LIKELY]]
7373
while (B())
7474
[[unlikely]] { b = false; }
7575
}
7676

7777
void DoStmt() {
7878
// CHECK-LABEL: define{{.*}}DoStmt
79-
// CHECK: br {{.*}} !prof !8
79+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
8080
if (b)
8181
[[unlikely]] do {}
8282
while (B())
@@ -91,20 +91,20 @@ void DoStmt() {
9191

9292
void ForStmt() {
9393
// CHECK-LABEL: define{{.*}}ForStmt
94-
// CHECK: br {{.*}} !prof !8
94+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
9595
if (b)
9696
[[unlikely]] for (; B();) {}
9797

9898
// CHECK-NOT: br {{.*}} %if.end{{.*}} !prof
9999
if (b)
100-
// CHECK: br {{.*}} !prof !7
100+
// CHECK: br {{.*}} !prof ![[PROF_LIKELY]]
101101
for (; B();)
102102
[[unlikely]] {}
103103
}
104104

105105
void GotoStmt() {
106106
// CHECK-LABEL: define{{.*}}GotoStmt
107-
// CHECK: br {{.*}} !prof !8
107+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
108108
if (b)
109109
[[unlikely]] goto end;
110110
else {
@@ -116,7 +116,7 @@ end:;
116116

117117
void ReturnStmt() {
118118
// CHECK-LABEL: define{{.*}}ReturnStmt
119-
// CHECK: br {{.*}} !prof !8
119+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
120120
if (b)
121121
[[unlikely]] return;
122122
else {
@@ -127,7 +127,7 @@ void ReturnStmt() {
127127

128128
void SwitchStmt() {
129129
// CHECK-LABEL: define{{.*}}SwitchStmt
130-
// CHECK: br {{.*}} !prof !8
130+
// CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
131131
if (b)
132132
[[unlikely]] switch (i) {}
133133
else {
@@ -144,5 +144,5 @@ void SwitchStmt() {
144144
}
145145
}
146146

147-
// CHECK: !7 = !{!"branch_weights", i32 [[UNLIKELY]], i32 [[LIKELY]]}
148-
// CHECK: !8 = !{!"branch_weights", i32 [[LIKELY]], i32 [[UNLIKELY]]}
147+
// CHECK: ![[PROF_LIKELY]] = !{!"branch_weights", i32 [[UNLIKELY]], i32 [[LIKELY]]}
148+
// CHECK: ![[PROF_UNLIKELY]] = !{!"branch_weights", i32 [[LIKELY]], i32 [[UNLIKELY]]}

‎clang/test/CodeGenCXX/pr12251.cpp

+11-10
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,7 @@ bool f(bool *x) {
55
return *x;
66
}
77
// CHECK-LABEL: define{{.*}} zeroext i1 @_Z1fPb
8-
// CHECK: load i8, ptr %{{[^ ]*}}, align 1, !range [[RANGE_i8_0_2:![^ ]*]]
8+
// CHECK: load i8, ptr %{{[^ ]*}}, align 1, !range [[RANGE_i8_0_2:![0-9]+]], !noundef [[NOUNDEF:![0-9]+]]
99

1010
// Only enum-tests follow. Ensure that after the bool test, no further range
1111
// metadata shows up when strict enums are disabled.
@@ -32,63 +32,63 @@ e3 g3(e3 *x) {
3232
return *x;
3333
}
3434
// CHECK-LABEL: define{{.*}} i32 @_Z2g3P2e3
35-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32:![^ ]*]]
35+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32:![0-9]+]], !noundef [[NOUNDEF]]
3636

3737
enum e4 { e4_a = -16};
3838
e4 g4(e4 *x) {
3939
return *x;
4040
}
4141
// CHECK-LABEL: define{{.*}} i32 @_Z2g4P2e4
42-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16:![^ ]*]]
42+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16:![0-9]+]], !noundef [[NOUNDEF]]
4343

4444
enum e5 { e5_a = -16, e5_b = 16};
4545
e5 g5(e5 *x) {
4646
return *x;
4747
}
4848
// CHECK-LABEL: define{{.*}} i32 @_Z2g5P2e5
49-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![^ ]*]]
49+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![0-9]+]], !noundef [[NOUNDEF]]
5050

5151
enum e6 { e6_a = -1 };
5252
e6 g6(e6 *x) {
5353
return *x;
5454
}
5555
// CHECK-LABEL: define{{.*}} i32 @_Z2g6P2e6
56-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m1_1:![^ ]*]]
56+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m1_1:![0-9]+]], !noundef [[NOUNDEF]]
5757

5858
enum e7 { e7_a = -16, e7_b = 2};
5959
e7 g7(e7 *x) {
6060
return *x;
6161
}
6262
// CHECK-LABEL: define{{.*}} i32 @_Z2g7P2e7
63-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16]]
63+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m16_16]], !noundef [[NOUNDEF]]
6464

6565
enum e8 { e8_a = -17};
6666
e8 g8(e8 *x) {
6767
return *x;
6868
}
6969
// CHECK-LABEL: define{{.*}} i32 @_Z2g8P2e8
70-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![^ ]*]]
70+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m32_32:![0-9]+]], !noundef [[NOUNDEF]]
7171

7272
enum e9 { e9_a = 17};
7373
e9 g9(e9 *x) {
7474
return *x;
7575
}
7676
// CHECK-LABEL: define{{.*}} i32 @_Z2g9P2e9
77-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32]]
77+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_0_32]], !noundef [[NOUNDEF]]
7878

7979
enum e10 { e10_a = -16, e10_b = 32};
8080
e10 g10(e10 *x) {
8181
return *x;
8282
}
8383
// CHECK-LABEL: define{{.*}} i32 @_Z3g10P3e10
84-
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m64_64:![^ ]*]]
84+
// CHECK: load i32, ptr %x, align 4, !range [[RANGE_i32_m64_64:![0-9]+]], !noundef [[NOUNDEF]]
8585

8686
enum e11 {e11_a = 4294967296 };
8787
enum e11 g11(enum e11 *x) {
8888
return *x;
8989
}
9090
// CHECK-LABEL: define{{.*}} i64 @_Z3g11P3e11
91-
// CHECK: load i64, ptr %x, align {{[84]}}, !range [[RANGE_i64_0_2pow33:![^ ]*]]
91+
// CHECK: load i64, ptr %x, align {{[84]}}, !range [[RANGE_i64_0_2pow33:![0-9]+]], !noundef [[NOUNDEF]]
9292

9393
enum e12 {e12_a = 9223372036854775808U };
9494
enum e12 g12(enum e12 *x) {
@@ -137,6 +137,7 @@ e16 g16(e16 *x) {
137137

138138

139139
// CHECK: [[RANGE_i8_0_2]] = !{i8 0, i8 2}
140+
// CHECK: [[NOUNDEF]] = !{}
140141
// CHECK: [[RANGE_i32_0_32]] = !{i32 0, i32 32}
141142
// CHECK: [[RANGE_i32_m16_16]] = !{i32 -16, i32 16}
142143
// CHECK: [[RANGE_i32_m32_32]] = !{i32 -32, i32 32}

‎clang/test/CodeGenOpenCL/builtins-amdgcn.cl

+6-6
Original file line numberDiff line numberDiff line change
@@ -569,9 +569,9 @@ void test_s_getreg(volatile global uint *out)
569569
}
570570

571571
// CHECK-LABEL: @test_get_local_id(
572-
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[$WI_RANGE:![0-9]*]]
573-
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[$WI_RANGE]]
574-
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[$WI_RANGE]]
572+
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[$WI_RANGE:![0-9]*]], !noundef
573+
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[$WI_RANGE]], !noundef
574+
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[$WI_RANGE]], !noundef
575575
void test_get_local_id(int d, global int *out)
576576
{
577577
switch (d) {
@@ -585,11 +585,11 @@ void test_get_local_id(int d, global int *out)
585585
// CHECK-LABEL: @test_get_workgroup_size(
586586
// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
587587
// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 4
588-
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load
588+
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
589589
// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 6
590-
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
590+
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
591591
// CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 8
592-
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load
592+
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
593593
void test_get_workgroup_size(int d, global int *out)
594594
{
595595
switch (d) {

0 commit comments

Comments
 (0)