Skip to content

Commit

Permalink
[Clang] Emit noundef metadata next to range metadata
Browse files Browse the repository at this point in the history 
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
  • Loading branch information
@nikic
nikic committed Jan 12, 2023
1 parent 84a5d93 commit 0285656
Show file tree
Hide file tree
Showing 6 changed files with 46 additions and 38 deletions.
4 changes: 4 additions & 0 deletions clang/lib/CodeGen/CGBuiltin.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}

Expand Down Expand Up @@ -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;
Expand Down
5 changes: 4 additions & 1 deletion clang/lib/CodeGen/CGExpr.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Expand Down
12 changes: 6 additions & 6 deletions clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
30 changes: 15 additions & 15 deletions clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand All @@ -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();
Expand All @@ -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();

Expand All @@ -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 {
Expand All @@ -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()) {}

Expand All @@ -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())
Expand All @@ -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 {
Expand All @@ -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 {
Expand All @@ -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 {
Expand All @@ -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 changes: 11 additions & 10 deletions clang/test/CodeGenCXX/pr12251.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand All @@ -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) {
Expand Down Expand Up @@ -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}
Expand Down
12 changes: 6 additions & 6 deletions clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -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) {
Expand Down

0 comments on commit 0285656

Please sign in to comment.