[AMDGPU] Add builtin s_sendmsg_rtn

[b-sumner]

Following existing naming, it might make sense to rename "rtn_b32" --> "rtn" and "rtn_b64" --> "rtnl".

[yaxunl]

In D132140#3732262, @b-sumner wrote:

Following existing naming, it might make sense to rename "rtn_b32" --> "rtn" and "rtn_b64" --> "rtnl".

will modify. thanks.

[yaxunl]

revised by Brian's comments

[b-sumner]

In D132140#3732337, @yaxunl wrote:

revised by Brian's comments

Thank you. Looks good to me.

[tra]

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl

24

Is there a particular reason for this test?

Argument and return value type checks should've been handled by Sema based on builtin definition.

Integer and float conversions should've been done automatically, too.

We already have tests for all of that.

[yaxunl]

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl

24

Will remove.

[yaxunl]

remove unnecessary tests

clang/include/clang/Basic/BuiltinsAMDGPU.def

 
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w64, "V4fV16hV16hV4f", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64, "V4fV16sV16sV4f", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w64, "V8hV16hV16hV8hIb", "nc", "gfx11-insts")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64, "V8sV16sV16sV8sIb", "nc", "gfx11-insts")
 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.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_read_exec, "WUi", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
 
 BUILTIN(__builtin_amdgcn_endpgm, "v", "nr")

clang/lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_atomic_dec64: {
 
     QualType PtrTy = E->getArg(0)->IgnoreImpCasts()->getType();
     bool Volatile =
       PtrTy->castAs<PointerType>()->getPointeeType().isVolatileQualified();
     Value *IsVolatile = Builder.getInt1(static_cast<bool>(Volatile));
 
     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;
   }
 }
 
 /// Handle a SystemZ function in which the final argument is a pointer
 /// to an int that receives the post-instruction CC value.  At the LLVM level
 /// this is represented as a function that returns a {result, cc} pair.

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl

+// 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);
+}

clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl

+// 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
+}