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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 453681.
yaxunl added a comment.

revised by Brian's comments


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D132140/new/

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,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
+}
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(
+// 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);
+}
+
+// Test mismatched argument and return types are handled.
+
+// CHECK-LABEL: @test_s_sendmsg_rtnl_mismatch(
+// CHECK: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0)
+// CHECK:  uitofp i64 %{{.*}} to double
+void test_s_sendmsg_rtnl_mismatch(global double* out) {
+  *out = __builtin_amdgcn_s_sendmsg_rtnl((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:
+  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;
   }
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, "UiUIi", "n", "gfx11-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtnl, "UWiUIi", "n", "gfx11-insts")
+
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//


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