[PATCH] D132140: [AMDGPU] Add builtin s_sendmsg_rtn_b{32|64}

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: b-sumner, arsenm, foad, kzhuravl, bcahoon.
Herald added subscribers: kosarev, kerbowa, t-tye, tpr, dstuttard, jvesely.
Herald added a project: All.
yaxunl requested review of this revision.
Herald added a subscriber: wdng.

https://reviews.llvm.org/D132140

Files:
  clang/include/clang/Basic/BuiltinsAMDGPU.def
  clang/lib/CodeGen/CGBuiltin.cpp
  clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
  clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl


Index: clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
===================================================================
--- /dev/null
+++ clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
@@ -0,0 +1,14 @@
+// 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
+}
Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
===================================================================
--- /dev/null
+++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
@@ -0,0 +1,30 @@
+// 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);
+}
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -17015,6 +17015,15 @@
 
     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;
   }
Index: clang/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -277,6 +277,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_b32, "UiUIi", "n", "gfx11-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtn_b64, "UWiUIi", "n", "gfx11-insts")
+
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D132140.453669.patch
Type: text/x-patch
Size: 4177 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220818/3af2f32e/attachment.bin>