diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -687,6 +687,8 @@
     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 @@
   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
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -1751,8 +1751,11 @@
     // 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
--- 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
--- a/clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
+++ b/clang/test/CodeGenCXX/attr-likelihood-if-branch-weights.cpp
@@ -8,7 +8,7 @@
 
 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 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 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 @@
 
 void NullStmt() {
   // CHECK-LABEL: define{{.*}}NullStmt
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]];
   else {
@@ -49,7 +49,7 @@
 
 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 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 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 @@
 
 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 SwitchStmt() {
   // CHECK-LABEL: define{{.*}}SwitchStmt
-  // CHECK: br {{.*}} !prof !8
+  // CHECK: br {{.*}} !prof ![[PROF_UNLIKELY]]
   if (b)
     [[unlikely]] switch (i) {}
   else {
@@ -144,5 +144,5 @@
   }
 }
 
-// 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
--- a/clang/test/CodeGenCXX/pr12251.cpp
+++ b/clang/test/CodeGenCXX/pr12251.cpp
@@ -5,7 +5,7 @@
   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 @@
   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 @@
 
 
 // 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
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -569,9 +569,9 @@
 }
 
 // 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 @@
 // 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) {