intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-25 10:55:58 +08:00
Code Issues Packages Projects Releases Wiki Activity

[Clang] Emit noundef metadata next to range metadata

Browse Source 
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
This commit is contained in:
Nikita Popov
2023-01-11 15:19:57 +01:00
parent 84a5d93f43
commit 02856565ac
6 changed files with 46 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
clang/lib/CodeGen/CGBuiltin.cpp
View File

@@ -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;

5
clang/lib/CodeGen/CGExpr.cpp
View File

@@ -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);
}

12
clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
View File

@@ -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) {

30
clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
View File

@@ -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]]}

21
clang/test/CodeGenCXX/pr12251.cpp
View File

@@ -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}

12
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@@ -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) {