[PATCH] D134872: AMDGPU: Add __builtin_amdgcn_permlane64

[Matt Arsenault via Phabricator via cfe-commits]

arsenm created this revision.
arsenm added reviewers: yaxunl, AMDGPU.
Herald added subscribers: kosarev, kerbowa, t-tye, tpr, dstuttard, jvesely, kzhuravl.
Herald added a project: All.
arsenm requested review of this revision.
Herald added subscribers: llvm-commits, wdng.
Herald added a project: LLVM.

https://reviews.llvm.org/D134872

Files:
  clang/include/clang/Basic/BuiltinsAMDGPU.def
  clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
  clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
  llvm/include/llvm/IR/IntrinsicsAMDGPU.td


Index: llvm/include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1991,6 +1991,7 @@
 
 // llvm.amdgcn.permlane64 <src0>
 def int_amdgcn_permlane64 :
+  ClangBuiltin<"__builtin_amdgcn_permlane64">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn]>;
 
Index: clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
===================================================================
--- clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
+++ clang/test/SemaOpenCL/builtins-amdgcn-gfx11.cl
@@ -13,4 +13,6 @@
 #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
+
+  *out1 = __builtin_amdgcn_permlane64(x); // GFX10-error {{'__builtin_amdgcn_permlane64' needs target feature gfx11-insts}}
 }
Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
===================================================================
--- clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
+++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
@@ -31,3 +31,9 @@
 {
   *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, 128);
 }
+
+// CHECK-LABEL: @test_permlane64(
+// CHECK: call i32 @llvm.amdgcn.permlane64(i32 %a)
+void test_permlane64(global uint* out, uint a) {
+  *out = __builtin_amdgcn_permlane64(a);
+}
Index: clang/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -259,6 +259,9 @@
 // GFX11+ only builtins.
 //===----------------------------------------------------------------------===//
 
+// TODO: This is a no-op in wave32. Should the builtin require wavefrontsize64?
+TARGET_BUILTIN(__builtin_amdgcn_permlane64, "UiUi", "nc", "gfx11-insts")
+
 //===----------------------------------------------------------------------===//
 // WMMA builtins.
 // Postfix w32 indicates the builtin requires wavefront size of 32.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D134872.463858.patch
Type: text/x-patch
Size: 2201 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220929/1e2e9d3b/attachment-0001.bin>