[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

diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 430b5f43cdd5ae3dcfa13237f8bd78da7289ce8a..479e24245df4df3f2534230d81254fbaeae1c188 100644 (file)
--- 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 6d5e729b1eea9ac9633c27471747edada75a65ff..34974c63984e626e425e0b6330c60fdff2ca0cb9 100644 (file)
--- 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 e506b875b674814e89f30f9b178e25bbcd540ebd..c098917a0a0e21906d81b3a27c71b64040643a30 100644 (file)
--- 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 4f0009e29056683434f4be940d1a57952f5fd1d8..2583a4a84c3c3e1cb22dadb2c3b0b161c0b74284 100644 (file)
--- 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 a267a3aed077de45e261b43ae92b19244ff0d6d3..bd5c85b83f2caf05b9c3d97d45a39c63d8160aaf 100644 (file)
--- 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 82cd3177d6c6d1f38416be235ae9d26fb934985b..094851b218898d4e7e0e7645cdc93df981620daf 100644 (file)
--- 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) {