[PATCH] D63216: [AMDGPU] ImmArg and SourceOfDivergence for permlane/dpp

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec created this revision.
rampitec added reviewers: arsenm, kzhuravl, msearles.
Herald added subscribers: t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng, jvesely.
rampitec added parent revisions: D63203: [AMDGPU] gfx1010 dpp16 and dpp8, D63202: [AMDGPU] gfx1010 premlane instructions.

Added missing ImmArg and SourceOfDivergence to the crosslane
intrinsics


https://reviews.llvm.org/D63216

Files:
  include/llvm/IR/IntrinsicsAMDGPU.td
  lib/Target/AMDGPU/AMDGPUSearchableTables.td


Index: lib/Target/AMDGPU/AMDGPUSearchableTables.td
===================================================================
--- lib/Target/AMDGPU/AMDGPUSearchableTables.td
+++ lib/Target/AMDGPU/AMDGPUSearchableTables.td
@@ -97,6 +97,11 @@
 def : SourceOfDivergence<int_amdgcn_ds_swizzle>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_add>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>;
+def : SourceOfDivergence<int_amdgcn_permlane16>;
+def : SourceOfDivergence<int_amdgcn_permlanex16>;
+def : SourceOfDivergence<int_amdgcn_mov_dpp>;
+def : SourceOfDivergence<int_amdgcn_mov_dpp8>;
+def : SourceOfDivergence<int_amdgcn_update_dpp>;
 
 foreach intr = AMDGPUImageDimAtomicIntrinsics in
 def : SourceOfDivergence<intr>;
Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1441,13 +1441,13 @@
 def int_amdgcn_permlane16 :
   Intrinsic<[llvm_i32_ty],
             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
-            [IntrNoMem, IntrConvergent]>;
+            [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
 
 // llvm.amdgcn.permlanex16 <old> <src0> <src1> <src2> <fi> <bound_control>
 def int_amdgcn_permlanex16 :
   Intrinsic<[llvm_i32_ty],
             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
-            [IntrNoMem, IntrConvergent]>;
+            [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
 
 // llvm.amdgcn.mov.dpp8.i32 <src> <sel>
 // <sel> is a 32-bit constant whose high 8 bits must be zero which selects
@@ -1455,7 +1455,7 @@
 def int_amdgcn_mov_dpp8 :
   Intrinsic<[llvm_anyint_ty],
             [LLVMMatchType<0>, llvm_i32_ty],
-            [IntrNoMem, IntrConvergent]>;
+            [IntrNoMem, IntrConvergent, ImmArg<1>]>;
 
 def int_amdgcn_s_get_waveid_in_workgroup :
   GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D63216.204305.patch
Type: text/x-patch
Size: 1979 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20190612/15d947f3/attachment-0001.bin>