[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_b32, "UiUIi", "n", "gfx11-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtn_b64, "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

 
     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_b32:
+  case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn_b64: {
+    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 ushort;
+typedef unsigned int uint;
+typedef unsigned long ulong;
+
+// CHECK-LABEL: @test_s_sendmsg_rtn_b32(
+// CHECK: call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 0)
+void test_s_sendmsg_rtn_b32(global uint* out) {
+  *out = __builtin_amdgcn_s_sendmsg_rtn_b32(0);
+}
+
+// CHECK-LABEL: @test_s_sendmsg_rtn_b64(
+// CHECK: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0)
+void test_s_sendmsg_rtn_b64(global ulong* out) {
+  *out = __builtin_amdgcn_s_sendmsg_rtn_b64(0);
+}
+
+// Test mismatched argument and return types are handled.
+
+// CHECK-LABEL: @test_s_sendmsg_rtn_b64_mismatch(
+// CHECK: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0)
+// CHECK:  uitofp i64 %{{.*}} to double
+void test_s_sendmsg_rtn_b64_mismatch(global double* out) {
+  *out = __builtin_amdgcn_s_sendmsg_rtn_b64((ushort)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_b32(0); // GFX10-error {{'__builtin_amdgcn_s_sendmsg_rtn_b32' needs target feature gfx11-insts}}
+  *out2 = __builtin_amdgcn_s_sendmsg_rtn_b64(0); // GFX10-error {{'__builtin_amdgcn_s_sendmsg_rtn_b64' needs target feature gfx11-insts}}
+#if __has_builtin(__builtin_amdgcn_s_sendmsg_rtn_b32)
+  *out1 = __builtin_amdgcn_s_sendmsg_rtn_b32(x); // GFX11-error {{argument to '__builtin_amdgcn_s_sendmsg_rtn_b32' must be a constant integer}}
+  *out2 = __builtin_amdgcn_s_sendmsg_rtn_b64(x); // GFX11-error {{argument to '__builtin_amdgcn_s_sendmsg_rtn_b64' must be a constant integer}}
+#endif
+}