diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index ee877349a3314..3d3dc5572f053 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -139,6 +139,8 @@ def IntrNoReturn : IntrinsicProperty;
 // Applied by default.
 def IntrNoCallback : IntrinsicProperty<1>;
 
+def IntrNoRecurse : IntrinsicProperty;
+
 // IntrNoSync - Threads executing the intrinsic will not synchronize using
 // memory or other means. Applied by default.
 def IntrNoSync : IntrinsicProperty<1>;
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f721d5267cd2a..eb7bde6999491 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -237,7 +237,7 @@ def int_amdgcn_reloc_constant : DefaultAttrsIntrinsic<
 // the second one is copied to m0
 def int_amdgcn_s_sendmsg : ClangBuiltin<"__builtin_amdgcn_s_sendmsg">,
   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
-  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;
+  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
 def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
   [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;
@@ -246,7 +246,7 @@ def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
 // gfx11 intrinsic
 // The first parameter is s_sendmsg immediate (i16). Return type is i32 or i64.
 def int_amdgcn_s_sendmsg_rtn : Intrinsic <[llvm_anyint_ty], [llvm_i32_ty],
-  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;
+  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
 
 // Vanilla workgroup sync-barrier
 def int_amdgcn_s_barrier : ClangBuiltin<"__builtin_amdgcn_s_barrier">,
diff --git a/llvm/test/TableGen/intrinsic-attrs.td b/llvm/test/TableGen/intrinsic-attrs.td
index 579b5e8a21b86..70056d77ff9db 100644
--- a/llvm/test/TableGen/intrinsic-attrs.td
+++ b/llvm/test/TableGen/intrinsic-attrs.td
@@ -6,6 +6,8 @@ def int_random_gen   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrHasSideEffec
 
 def int_deref_ptr_ret : Intrinsic<[llvm_ptr_ty], [], [Dereferenceable<RetIndex, 16>]>;
 
+def int_no_recurse   : Intrinsic<[llvm_i32_ty], [], [IntrNoRecurse]>;
+
 // CHECK: static AttributeSet getIntrinsicArgAttributeSet(LLVMContext &C, unsigned ID) {
 // CHECK-NEXT:   switch (ID) {
 // CHECK-NEXT: default: llvm_unreachable("Invalid attribute set number");
@@ -21,11 +23,17 @@ def int_deref_ptr_ret : Intrinsic<[llvm_ptr_ty], [], [Dereferenceable<RetIndex,
 // CHECK-NEXT: return AttributeSet::get(C, {
 // CHECK-NEXT: Attribute::get(C, Attribute::NoUnwind),
 // CHECK-NEXT: });
+// CHECK: case 1:
+// CHECK-NEXT: return AttributeSet::get(C, {
+// CHECK-NEXT: Attribute::get(C, Attribute::NoUnwind),
+// CHECK-NEXT: Attribute::get(C, Attribute::NoRecurse),
+// CHECK-NEXT: });
 
 
 // CHECK: getAttributes(LLVMContext &C, ID id)
 // CHECK: 0 << 8 | 0, // llvm.deref.ptr.ret
-// CHECK: 1 << 8 | 1, // llvm.random.gen
+// CHECK: 1 << 8 | 1, // llvm.no.recurse
+// CHECK: 2 << 8 | 1, // llvm.random.gen
 
 // CHECK: case 1:
 // CHECK-NEXT: return AttributeList::get(C, {
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 0846f66ea6452..fca7038fa020f 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -388,6 +388,8 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
     isConvergent = true;
   else if (R->getName() == "IntrNoReturn")
     isNoReturn = true;
+  else if (R->getName() == "IntrNoRecurse")
+    isNoRecurse = true;
   else if (R->getName() == "IntrNoCallback")
     isNoCallback = true;
   else if (R->getName() == "IntrNoSync")
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index 8428d09a94009..90267cb1ed7da 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -89,6 +89,9 @@ struct CodeGenIntrinsic {
   /// True if the intrinsic is no-return.
   bool isNoReturn = false;
 
+  /// True if the intrinsic is norecurse.
+  bool isNoRecurse = false;
+
   /// True if the intrinsic is no-callback.
   bool isNoCallback = false;
 
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index 269d2a9544379..8362751846872 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -416,9 +416,9 @@ static bool compareFnAttributes(const CodeGenIntrinsic *L,
   auto TieBoolAttributes = [](const CodeGenIntrinsic *I) -> auto {
     // Sort throwing intrinsics after non-throwing intrinsics.
     return std::tie(I->canThrow, I->isNoDuplicate, I->isNoMerge, I->isNoReturn,
-                    I->isNoCallback, I->isNoSync, I->isNoFree, I->isWillReturn,
-                    I->isCold, I->isConvergent, I->isSpeculatable,
-                    I->hasSideEffects, I->isStrictFP);
+                    I->isNoRecurse, I->isNoCallback, I->isNoSync, I->isNoFree,
+                    I->isWillReturn, I->isCold, I->isConvergent,
+                    I->isSpeculatable, I->hasSideEffects, I->isStrictFP);
   };
 
   auto TieL = TieBoolAttributes(L);
@@ -440,10 +440,11 @@ static bool compareFnAttributes(const CodeGenIntrinsic *L,
 /// NoUnwind = !canThrow, so we need to negate it's sense to test if the
 // intrinsic has NoUnwind attribute.
 static bool hasFnAttributes(const CodeGenIntrinsic &Int) {
-  return !Int.canThrow || Int.isNoReturn || Int.isNoCallback || Int.isNoSync ||
-         Int.isNoFree || Int.isWillReturn || Int.isCold || Int.isNoDuplicate ||
-         Int.isNoMerge || Int.isConvergent || Int.isSpeculatable ||
-         Int.isStrictFP || getEffectiveME(Int) != MemoryEffects::unknown();
+  return !Int.canThrow || Int.isNoReturn || Int.isNoRecurse ||
+         Int.isNoCallback || Int.isNoSync || Int.isNoFree || Int.isWillReturn ||
+         Int.isCold || Int.isNoDuplicate || Int.isNoMerge || Int.isConvergent ||
+         Int.isSpeculatable || Int.isStrictFP ||
+         getEffectiveME(Int) != MemoryEffects::unknown();
 }
 
 namespace {
@@ -572,6 +573,8 @@ static AttributeSet getIntrinsicFnAttributeSet(LLVMContext &C, unsigned ID) {
       addAttribute("NoUnwind");
     if (Int.isNoReturn)
       addAttribute("NoReturn");
+    if (Int.isNoRecurse)
+      addAttribute("NoRecurse");
     if (Int.isNoCallback)
       addAttribute("NoCallback");
     if (Int.isNoSync)