diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 1271868a53b86..26da82843c512 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -231,7 +231,7 @@ extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
 
 // CHECK-LABEL: @test_abs(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[TMP0:%.*]] = tail call noundef i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call noundef range(i32 0, -2147483648) i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
 // CHECK-NEXT:    ret i32 [[TMP0]]
 //
 extern "C" __device__ int test_abs(int x) {
@@ -240,7 +240,7 @@ extern "C" __device__ int test_abs(int x) {
 
 // CHECK-LABEL: @test_labs(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[TMP0:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call noundef range(i64 0, -9223372036854775808) i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
 // CHECK-NEXT:    ret i64 [[TMP0]]
 //
 extern "C" __device__ long test_labs(long x) {
@@ -249,7 +249,7 @@ extern "C" __device__ long test_labs(long x) {
 
 // CHECK-LABEL: @test_llabs(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[TMP0:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call noundef range(i64 0, -9223372036854775808) i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
 // CHECK-NEXT:    ret i64 [[TMP0]]
 //
 extern "C" __device__ long long test_llabs(long x) {
diff --git a/llvm/lib/Transforms/Utils/InlineFunction.cpp b/llvm/lib/Transforms/Utils/InlineFunction.cpp
index 7b846f2d2d72d..eb471b259c7d4 100644
--- a/llvm/lib/Transforms/Utils/InlineFunction.cpp
+++ b/llvm/lib/Transforms/Utils/InlineFunction.cpp
@@ -30,11 +30,12 @@
 #include "llvm/Analysis/ProfileSummaryInfo.h"
 #include "llvm/Analysis/ValueTracking.h"
 #include "llvm/Analysis/VectorUtils.h"
-#include "llvm/IR/AttributeMask.h"
 #include "llvm/IR/Argument.h"
+#include "llvm/IR/AttributeMask.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/CFG.h"
 #include "llvm/IR/Constant.h"
+#include "llvm/IR/ConstantRange.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/DebugInfo.h"
@@ -1450,6 +1451,8 @@ static AttrBuilder IdentifyValidPoisonGeneratingAttributes(CallBase &CB) {
     Valid.addAttribute(Attribute::NonNull);
   if (CB.hasRetAttr(Attribute::Alignment))
     Valid.addAlignmentAttr(CB.getRetAlign());
+  if (std::optional<ConstantRange> Range = CB.getRange())
+    Valid.addRangeAttr(*Range);
   return Valid;
 }
 
@@ -1541,6 +1544,14 @@ static void AddReturnAttributes(CallBase &CB, ValueToValueMapTy &VMap) {
     if (ValidPG.getAlignment().valueOrOne() < AL.getRetAlignment().valueOrOne())
       ValidPG.removeAttribute(Attribute::Alignment);
     if (ValidPG.hasAttributes()) {
+      Attribute CBRange = ValidPG.getAttribute(Attribute::Range);
+      if (CBRange.isValid()) {
+        Attribute NewRange = AL.getRetAttr(Attribute::Range);
+        if (NewRange.isValid()) {
+          ValidPG.addRangeAttr(
+              CBRange.getRange().intersectWith(NewRange.getRange()));
+        }
+      }
       // Three checks.
       // If the callsite has `noundef`, then a poison due to violating the
       // return attribute will create UB anyways so we can always propagate.
diff --git a/llvm/test/Transforms/Inline/ret_attr_align_and_noundef.ll b/llvm/test/Transforms/Inline/ret_attr_align_and_noundef.ll
index c038ffccf3e96..f4cebf1fcb5da 100644
--- a/llvm/test/Transforms/Inline/ret_attr_align_and_noundef.ll
+++ b/llvm/test/Transforms/Inline/ret_attr_align_and_noundef.ll
@@ -5,10 +5,12 @@
 
 declare ptr @foo()
 declare void @use.ptr(ptr) willreturn nounwind
+declare void @use.val(i8) willreturn nounwind
 declare void @bar()
 declare void @baz()
 declare ptr @llvm.ptrmask.p0.i64(ptr, i64)
 declare i1 @val()
+declare i8 @val8()
 
 define ptr @callee0123() {
 ; CHECK-LABEL: define ptr @callee0123() {
@@ -337,3 +339,74 @@ define ptr @caller12_todo() {
   %r = call nonnull ptr @callee12()
   ret ptr %r
 }
+
+define i8 @callee13() {
+; CHECK-LABEL: define i8 @callee13() {
+; CHECK-NEXT:    [[R:%.*]] = call i8 @val8()
+; CHECK-NEXT:    ret i8 [[R]]
+;
+  %r = call i8 @val8()
+  ret i8 %r
+}
+
+define i8 @caller13_okay_use_after_poison_anyways() {
+; CHECK-LABEL: define i8 @caller13_okay_use_after_poison_anyways() {
+; CHECK-NEXT:    [[R_I:%.*]] = call range(i8 0, 10) i8 @val8()
+; CHECK-NEXT:    call void @use.val(i8 [[R_I]])
+; CHECK-NEXT:    ret i8 [[R_I]]
+;
+  %r = call range(i8 0, 10) i8 @callee13()
+  call void @use.val(i8 %r)
+  ret i8 %r
+}
+
+define i8 @callee14() {
+; CHECK-LABEL: define i8 @callee14() {
+; CHECK-NEXT:    [[R:%.*]] = call noundef i8 @val8()
+; CHECK-NEXT:    ret i8 [[R]]
+;
+  %r = call noundef i8 @val8()
+  ret i8 %r
+}
+
+define i8 @caller14_fail_creates_ub() {
+; CHECK-LABEL: define i8 @caller14_fail_creates_ub() {
+; CHECK-NEXT:    [[R_I:%.*]] = call noundef i8 @val8()
+; CHECK-NEXT:    call void @use.val(i8 [[R_I]])
+; CHECK-NEXT:    ret i8 [[R_I]]
+;
+  %r = call range(i8 0, 10) i8 @callee14()
+  call void @use.val(i8 %r)
+  ret i8 %r
+}
+
+define i8 @caller14_okay_is_ub_anyways() {
+; CHECK-LABEL: define i8 @caller14_okay_is_ub_anyways() {
+; CHECK-NEXT:    [[R_I:%.*]] = call noundef range(i8 0, 10) i8 @val8()
+; CHECK-NEXT:    call void @use.val(i8 [[R_I]])
+; CHECK-NEXT:    ret i8 [[R_I]]
+;
+  %r = call noundef range(i8 0, 10) i8 @callee14()
+  call void @use.val(i8 %r)
+  ret i8 %r
+}
+
+define i8 @callee15() {
+; CHECK-LABEL: define i8 @callee15() {
+; CHECK-NEXT:    [[R:%.*]] = call range(i8 5, 10) i8 @val8()
+; CHECK-NEXT:    ret i8 [[R]]
+;
+  %r = call range(i8 5, 10) i8 @val8()
+  ret i8 %r
+}
+
+define i8 @caller15_okay_intersect_ranges() {
+; CHECK-LABEL: define i8 @caller15_okay_intersect_ranges() {
+; CHECK-NEXT:    [[R_I:%.*]] = call range(i8 5, 7) i8 @val8()
+; CHECK-NEXT:    call void @use.val(i8 [[R_I]])
+; CHECK-NEXT:    ret i8 [[R_I]]
+;
+  %r = call range(i8 0, 7) i8 @callee15()
+  call void @use.val(i8 %r)
+  ret i8 %r
+}