AMDGPU/SI: Implement DS_PERMUTE/DS_BPERMUTE Instruction Definitions and Intrinsics.

[cfang]

Summary

These are instructions introduced in VI+ Chips. We defined the instructions in this patch, and introduce intrinsics
llvm.amdgcn.ds.permute/llvm.amdgcn.ds.bpermute to expose them.

[arsenm]

lib/Target/AMDGPU/SIInstrInfo.cpp

230–231 ↗

(On Diff #49084)

Checking just offset0 should be sufficient, and can be moved above

lib/Target/AMDGPU/VIInstructions.td

131 ↗

(On Diff #49084)

Are we sure these don't real M0?

test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll

3–4 ↗

(On Diff #49084)

I would prefer splitting the 2 separate intrinsics into separate patches. These are also missing the readnone (which shoulda also use attribute groups)

[• tstellarAMD]

lib/Target/AMDGPU/VIInstructions.td

131 ↗

(On Diff #49084)

It reads M0, but it is supposed to ignore its value, so for our purposes we can treat it as if it doesn't read M0.

[cfang]

Update the patch based on Matt's Review:

1. Check only Offset0Imm and move the check one line ahead.
2. It is safe to remove M0 from the Uses list for ds_permute.ds_bpermute.
3. split the LIT test for ds_permute and ds_bpermute separately.

[cfang]

test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll

4–5 ↗

(On Diff #49202)

I just split the test case. If you want to split the intrinsics and/or instruction definitions, I can do it in the integration. Thanks.

[arsenm]

lib/Target/AMDGPU/VIInstructions.td

136–139 ↗

(On Diff #49202)

Why can't these be patterns on the instruction definition instead of standalone Pats?

[cfang]

Move the pattern for ds_permute intrinsic code generation into
the instruction definition, based on Matt's comment.

[arsenm]

LGTM

test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll

6 ↗

(On Diff #49427)

Might want to check that there are 2 VGPR operands

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

 
 def int_amdgcn_s_dcache_wb_vol :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_memrealtime :
   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
   Intrinsic<[llvm_i64_ty], [], []>;
+
+// llvm.amdgcn.ds.permute <index> <src>
+def int_amdgcn_ds_permute :
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
+
+// llvm.amdgcn.ds.bpermute <index> <src>
+def int_amdgcn_ds_bpermute :
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
+
 }

llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp

if (OffsetImm) {
       return true;
     }
 
     // The 2 offset instructions use offset0 and offset1 instead. We can treat
     // these as a load with a single offset if the 2 offsets are consecutive. We
     // will use this for some partially aligned loads.
     const MachineOperand *Offset0Imm = getNamedOperand(*LdSt,
                                                        AMDGPU::OpName::offset0);
+    // DS_PERMUTE does not have Offset0Imm (and Offset1Imm).
+    if (!Offset0Imm)
+      return false;
+
     const MachineOperand *Offset1Imm = getNamedOperand(*LdSt,
                                                        AMDGPU::OpName::offset1);
 
     uint8_t Offset0 = Offset0Imm->getImm();
     uint8_t Offset1 = Offset1Imm->getImm();
 
     if (Offset1 > Offset0 && Offset1 - Offset0 == 1) {
       // Each of these offsets is in element sized units, so we need to convert

llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td

let hasPostISelHook = 1 in {
 
     let data1 = 0 in {
       def _si : DS_Off16_Real_si <op, opName, outs, ins, asm>;
       def _vi : DS_Off16_Real_vi <op, opName, outs, ins, asm>;
     }
   }
 }
 
+multiclass DS_1A1D_PERMUTE <bits<8> op, string opName, RegisterClass rc,
+                            SDPatternOperator node = null_frag,
+  dag outs = (outs rc:$vdst),
+  dag ins = (ins VGPR_32:$addr, rc:$data0),
+  string asm = opName#" $vdst, $addr, $data0"> {
+
+  let mayLoad = 0, mayStore = 0, isConvergent = 1 in {
+    def "" : DS_Pseudo <opName, outs, ins,
+     [(set (i32 rc:$vdst),
+         (node (i32 VGPR_32:$addr), (i32 rc:$data0)))]>;
+
+    let data1 = 0, offset0 = 0, offset1 = 0, gds = 0  in {
+      def "_vi" : DS_Real_vi <op, opName, outs, ins, asm>;
+    }
+  }
+}
+
 multiclass DS_1A2D_RET_m <bits<8> op, string opName, RegisterClass rc,
                           string noRetOp = "", dag ins,
   dag outs = (outs rc:$vdst),
   string asm = opName#" $vdst, $addr, $data0, $data1"#"$offset"#"$gds"> {
 
   let hasPostISelHook = 1 in {
     def "" : DS_Pseudo <opName, outs, ins, []>,
              AtomicNoRet<noRetOp, 1>;

llvm/trunk/lib/Target/AMDGPU/VIInstructions.td

 // Misc Patterns
 //===----------------------------------------------------------------------===//
 
 def : Pat <
   (i64 (readcyclecounter)),
   (S_MEMREALTIME)
 >;
 
+//===----------------------------------------------------------------------===//
+// DS_PERMUTE/DS_BPERMUTE Instructions.
+//===----------------------------------------------------------------------===//
+
+let Uses = [EXEC] in {
+defm DS_PERMUTE_B32 : DS_1A1D_PERMUTE <0x3e, "ds_permute_b32", VGPR_32,
+                                       int_amdgcn_ds_permute>;
+defm DS_BPERMUTE_B32 : DS_1A1D_PERMUTE <0x3f, "ds_bpermute_b32", VGPR_32,
+                                       int_amdgcn_ds_bpermute>;
+}
+
 } // End Predicates = [isVI]

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
+
+declare i32 @llvm.amdgcn.ds.bpermute(i32, i32) #0
+
+; FUNC-LABEL: {{^}}ds_bpermute:
+; CHECK: ds_bpermute_b32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define void @ds_bpermute(i32 addrspace(1)* %out, i32 %index, i32 %src) nounwind {
+  %bpermute = call i32 @llvm.amdgcn.ds.bpermute(i32 %index, i32 %src) #0
+  store i32 %bpermute, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone convergent }

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
+
+declare i32 @llvm.amdgcn.ds.permute(i32, i32) #0
+
+; FUNC-LABEL: {{^}}ds_permute:
+; CHECK: ds_permute_b32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define void @ds_permute(i32 addrspace(1)* %out, i32 %index, i32 %src) nounwind {
+  %bpermute = call i32 @llvm.amdgcn.ds.permute(i32 %index, i32 %src) #0
+  store i32 %bpermute, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone convergent }