diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -278,6 +278,9 @@
 TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64, "V4iIbV4iIbV4iV4iIb", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64, "V4iIbV2iIbV2iV4iIb", "nc", "gfx11-insts")
 
+TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtn, "UiUIi", "n", "gfx11-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtnl, "UWiUIi", "n", "gfx11-insts")
+
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//
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
@@ -17054,6 +17054,15 @@
 
     return Builder.CreateCall(F, {Ptr, Val, MemOrder, MemScope, IsVolatile});
   }
+  case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:
+  case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {
+    llvm::Value *Arg = EmitScalarExpr(E->getArg(0));
+    llvm::Type *ResultType = ConvertType(E->getType());
+    // s_sendmsg_rtn is mangled using return type only.
+    Function *F =
+        CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});
+    return Builder.CreateCall(F, {Arg});
+  }
   default:
     return nullptr;
   }
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
new file mode 100644
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
@@ -0,0 +1,20 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1101 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1102 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1103 -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned int uint;
+typedef unsigned long ulong;
+
+// CHECK-LABEL: @test_s_sendmsg_rtn(
+// CHECK: call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 0)
+void test_s_sendmsg_rtn(global uint* out) {
+  *out = __builtin_amdgcn_s_sendmsg_rtn(0);
+}
+
+// CHECK-LABEL: @test_s_sendmsg_rtnl(
+// CHECK: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0)
+void test_s_sendmsg_rtnl(global ulong* out) {
+  *out = __builtin_amdgcn_s_sendmsg_rtnl(0);
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl b/clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
new file mode 100644
--- /dev/null
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1030 -verify=GFX10 -S -o - %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1100 -verify=GFX11 -S -o - %s
+
+typedef unsigned int uint;
+typedef unsigned long ulong;
+
+void test(global uint* out1, global ulong* out2, int x) {
+  *out1 = __builtin_amdgcn_s_sendmsg_rtn(0); // GFX10-error {{'__builtin_amdgcn_s_sendmsg_rtn' needs target feature gfx11-insts}}
+  *out2 = __builtin_amdgcn_s_sendmsg_rtnl(0); // GFX10-error {{'__builtin_amdgcn_s_sendmsg_rtnl' needs target feature gfx11-insts}}
+#if __has_builtin(__builtin_amdgcn_s_sendmsg_rtn)
+  *out1 = __builtin_amdgcn_s_sendmsg_rtn(x); // GFX11-error {{argument to '__builtin_amdgcn_s_sendmsg_rtn' must be a constant integer}}
+#endif
+#if __has_builtin(__builtin_amdgcn_s_sendmsg_rtnl)
+  *out2 = __builtin_amdgcn_s_sendmsg_rtnl(x); // GFX11-error {{argument to '__builtin_amdgcn_s_sendmsg_rtnl' must be a constant integer}}
+#endif
+}