[Clang] Emit noundef metadata next to range metadata

[nikic]

Summary

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.

Depends on D141386.

[nlopes]

LGTM

clang/lib/CodeGen/CGBuiltin.cpp

 static Value *emitRangedBuiltin(CodeGenFunction &CGF,
                                 unsigned IntrinsicID,
                                 int low, int high) {
     llvm::MDBuilder MDHelper(CGF.getLLVMContext());
     llvm::MDNode *RNode = MDHelper.createRange(APInt(32, low), APInt(32, high));
     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;
 }
 
 namespace {
   struct WidthAndSignedness {
     unsigned Width;
     bool Signed;
   };
auto *DstTy =
       CGF.Int16Ty->getPointerTo(GEP->getType()->getPointerAddressSpace());
   auto *Cast = CGF.Builder.CreateBitCast(GEP, DstTy);
   auto *LD = CGF.Builder.CreateLoad(
       Address(Cast, CGF.Int16Ty, CharUnits::fromQuantity(2)));
   llvm::MDBuilder MDHelper(CGF.getLLVMContext());
   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;
 }
 
 // \p Index is 0, 1, and 2 for x, y, and z dimension, respectively.
 Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) {
   const unsigned XOffset = 12;

clang/lib/CodeGen/CGExpr.cpp

llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile,
   }
 
   CGM.DecorateInstructionWithTBAA(Load, TBAAInfo);
 
   if (EmitScalarRangeCheck(Load, Ty, Loc)) {
     // 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);
 }
 
 llvm::Value *CodeGenFunction::EmitToMemory(llvm::Value *Value, QualType Ty) {
   // Bool has a different representation in memory than in registers.
   if (hasBooleanRepresentation(Ty)) {
     // This should really always be an i1, but sometimes it's already

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

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
 // RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
 // RUN:     | FileCheck -check-prefix=PRECOV5 %s
 
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
 // RUN:     -fcuda-is-device -mcode-object-version=5 -emit-llvm -o - -x hip %s \
 // RUN:     | FileCheck -check-prefix=COV5 %s
 
 #include "Inputs/cuda.h"
 
 // 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) {
   case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;
   case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
   case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
   default: *out = 0;
   }
 }
 
 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}

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

 // RUN: %clang_cc1 -O1 -emit-llvm %s -o - -triple=x86_64-linux-gnu | FileCheck -DLIKELY=2000 -DUNLIKELY=1 %s
 // RUN: %clang_cc1 -O1 -emit-llvm %s -triple=x86_64-linux-gnu -mllvm -likely-branch-weight=99 -mllvm -unlikely-branch-weight=42 -o - | FileCheck -DLIKELY=99 -DUNLIKELY=42 %s
 
 extern volatile bool b;
 extern volatile int i;
 extern bool A();
 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();
     }
   return B();
 }
 
 bool g() {
   // CHECK-LABEL: define{{.*}} zeroext i1 @_Z1gv
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY:[0-9]+]]
   if (b)
     [[unlikely]] {
       return A();
     }
 
   return B();
 }
 
 bool h() {
   // CHECK-LABEL: define{{.*}} zeroext i1 @_Z1hv
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]] return A();
 
   return B();
 }
 
 void NullStmt() {
   // CHECK-LABEL: define{{.*}}NullStmt
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]];
   else {
     // Make sure the branches aren't optimized away.
     b = true;
   }
 }
 
 void IfStmt() {
   // CHECK-LABEL: define{{.*}}IfStmt
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]] if (B()) {}
 
   // CHECK-NOT: br {{.*}} !prof
   // CHECK: br {{.*}} !prof
   if (b) {
     if (B())
       [[unlikely]] { b = false; }
   }
 }
 
 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())
       ;
 
   // CHECK-NOT: br {{.*}} %if.end{{.*}} !prof
   if (b)
     do
       [[unlikely]] {}
     while (B());
 }
 
 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 {
     // Make sure the branches aren't optimized away.
     b = true;
   }
 end:;
 }
 
 void ReturnStmt() {
   // CHECK-LABEL: define{{.*}}ReturnStmt
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]] return;
   else {
     // Make sure the branches aren't optimized away.
     b = true;
   }
 }
 
 void SwitchStmt() {
   // CHECK-LABEL: define{{.*}}SwitchStmt
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]] switch (i) {}
   else {
     // Make sure the branches aren't optimized away.
     b = true;
   }
   // CHECK-NOT: br {{.*}} %if.end{{.*}} !prof
   if (b)
     switch (i)
       [[unlikely]] {}
   else {
     // Make sure the branches aren't optimized away.
     b = true;
   }
 }
 
-// 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]]}

clang/test/CodeGenCXX/pr12251.cpp

 // RUN: %clang_cc1 %s -triple i386-unknown-unknown -emit-llvm -O1 -relaxed-aliasing -fstrict-enums -std=c++11 -o - | FileCheck %s
 // RUN: %clang_cc1 %s -triple i386-unknown-unknown -emit-llvm -O1 -relaxed-aliasing -std=c++11 -o - | FileCheck --check-prefix=NO-STRICT-ENUMS %s
 
 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.
 // NO-STRICT-ENUMS-LABEL: define{{.*}} zeroext i1 @_Z1fPb
 // NO-STRICT-ENUMS: load i8, ptr %{{[^ ]*}}, align 1, !range
 // NO-STRICT-ENUMS-NOT: !range
 
 enum e1 { };

 // CHECK-LABEL: define{{.*}} i32 @_Z2g2P2e2
 // CHECK: ret i32 %0
 
 enum e3 { e3_a = 16 };
 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) {
   return *x;
 }
 // CHECK-LABEL: define{{.*}} i64 @_Z3g12P3e12
 // CHECK: load i64, ptr %x, align {{[84]}}
 // CHECK-NOT: range

 }
 // CHECK-LABEL: define{{.*}} i32 @_Z3g16P3e16
 // CHECK: load i32, ptr %x, align 4
 // CHECK-NOT: range
 // CHECK: ret
 
 
 // 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}
 // CHECK: [[RANGE_i32_m1_1]] = !{i32 -1, i32 1}
 // CHECK: [[RANGE_i32_m64_64]] = !{i32 -64, i32 64}
 // CHECK: [[RANGE_i64_0_2pow33]] = !{i64 0, i64 8589934592}

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 void test_s_getreg(volatile global uint *out)
 {
   *out = __builtin_amdgcn_s_getreg(0);
   *out = __builtin_amdgcn_s_getreg(1);
   *out = __builtin_amdgcn_s_getreg(65535);
 }
 
 // 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) {
   case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
   case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
   case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
   default: *out = 0;
   }
 }
 
 // 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) {
   case 0: *out = __builtin_amdgcn_workgroup_size_x() + 1; break;
   case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
   case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
   default: *out = 0;
   }