AMDGPU: Add cache invalidation instructions.

[arsenm]

Summary

These are necessary for implementing mem_fence for OpenCL 2.0.

The VI assembler tests are disabled since it seems to be using the wrong encoding or opcode.

[• tstellarAMD]

Is there a reason to split the assembler tests into 3 files?

I can look into the problem with VI.

[arsenm]

In D12055#230374, @tstellarAMD wrote:

Is there a reason to split the assembler tests into 3 files?

I can look into the problem with VI.

I think I was mostly just confused by always testing with not. I can probably merge them back in, although I think it would also make sense to split cases that are supposed to error into separate tests

[• tstellarAMD]

In D12055#230382, @arsenm wrote:

In D12055#230374, @tstellarAMD wrote:

Is there a reason to split the assembler tests into 3 files?

I can look into the problem with VI.

I think I was mostly just confused by always testing with not. I can probably merge them back in, although I think it would also make sense to split cases that are supposed to error into separate tests

I went back and forth on this a lot, but I think for testing just instructions that are legal on some targets, but not others, it's better to have all the tests in the same file. Otherwise, we will end up having to duplicate the same assembly string across multiple files.

[arsenm]

Keep assembler tests in same file. Assumes another patch that changes all of the check prefixes in the original test first

[• tstellarAMD]

Just one small change request, with that fixed, LGTM.

test/MC/AMDGPU/mubuf.s

366

Could you move the XVI: line into its own file and XFAIL it. I think this will make it easier to track.

[arsenm]

r248532

include/llvm/IR/IntrinsicsAMDGPU.td

 
 def int_AMDGPU_class : GCCBuiltin<"__builtin_amdgpu_class">,
   Intrinsic<[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]>;
 
 def int_AMDGPU_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
                                        "__builtin_amdgpu_read_workdim">;
 
 } // End TargetPrefix = "AMDGPU"
+
+let TargetPrefix = "amdgcn" in {
+
+// SI only
+def int_amdgcn_buffer_wbinvl1_sc :
+  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
+  Intrinsic<[], [], []>;
+
+// On CI+
+def int_amdgcn_buffer_wbinvl1_vol :
+  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
+  Intrinsic<[], [], []>;
+
+def int_amdgcn_buffer_wbinvl1 :
+  GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
+  Intrinsic<[], [], []>;
+
+}

lib/Target/AMDGPU/CIInstructions.td

 defm V_LOG_LEGACY_F32 : VOP1Inst <vop1<0x45, 0x4c>, "v_log_legacy_f32",
   VOP_F32_F32
 >;
 defm V_EXP_LEGACY_F32 : VOP1Inst <vop1<0x46, 0x4b>, "v_exp_legacy_f32",
   VOP_F32_F32
 >;
 
 //===----------------------------------------------------------------------===//
+// MUBUF Instructions
+//===----------------------------------------------------------------------===//
+
+defm BUFFER_WBINVL1_VOL : MUBUF_Invalidate <mubuf<0x70, 0x3f>,
+  "buffer_wbinvl1_vol", int_amdgcn_buffer_wbinvl1_vol
+>;
+
+//===----------------------------------------------------------------------===//
 // Flat Instructions
 //===----------------------------------------------------------------------===//
 
 def FLAT_LOAD_UBYTE : FLAT_Load_Helper <0x8, "flat_load_ubyte", VGPR_32>;
 def FLAT_LOAD_SBYTE : FLAT_Load_Helper <0x9, "flat_load_sbyte", VGPR_32>;
 def FLAT_LOAD_USHORT : FLAT_Load_Helper <0xa, "flat_load_ushort", VGPR_32>;
 def FLAT_LOAD_SSHORT : FLAT_Load_Helper <0xb, "flat_load_sshort", VGPR_32>;
 def FLAT_LOAD_DWORD : FLAT_Load_Helper <0xc, "flat_load_dword", VGPR_32>;

lib/Target/AMDGPU/SIInstrInfo.td

let offen = 0, idxen = 0 in {
                                     [(st store_vt:$vdata,
                                       (MUBUFAddr64 v4i32:$srsrc, i64:$vaddr,
                                                    i32:$soffset, i16:$offset,
                                                    i1:$glc, i1:$slc, i1:$tfe))]>;
     }
   } // End mayLoad = 0, mayStore = 1
 }
 
+// For cache invalidation instructions.
+multiclass MUBUF_Invalidate <mubuf op, string opName, SDPatternOperator node> {
+  let hasSideEffects = 1, mayStore = 1, AsmMatchConverter = "" in {
+    def "" : MUBUF_Pseudo <opName, (outs), (ins), [(node)]>;
+
+    // Set everything to 0.
+    let offset = 0, offen = 0, idxen = 0, glc = 0, vaddr = 0,
+        vdata = 0, srsrc = 0, slc = 0, tfe = 0, soffset = 0 in {
+      let addr64 = 0 in {
+        def _si : MUBUF_Real_si <op, opName, (outs), (ins), opName>;
+      }
+
+      def _vi : MUBUF_Real_vi <op, opName, (outs), (ins), opName>;
+    }
+  } // End hasSideEffects = 1, mayStore = 1, AsmMatchConverter = ""
+}
+
 class FLAT_Load_Helper <bits<7> op, string asm, RegisterClass regClass> :
       FLAT <op, (outs regClass:$vdst),
                 (ins VReg_64:$addr, glc_flat:$glc, slc_flat:$slc, tfe_flat:$tfe),
             asm#" $vdst, $addr"#"$glc"#"$slc"#"$tfe", []> {
   let data = 0;
   let mayLoad = 1;
 }
 

lib/Target/AMDGPU/SIInstructions.td

 def SendMsgImm : Operand<i32> {
   let PrintMethod = "printSendMsg";
 }
 
 def isGCN : Predicate<"Subtarget->getGeneration() "
                       ">= AMDGPUSubtarget::SOUTHERN_ISLANDS">,
             AssemblerPredicate<"FeatureGCN">;
 def isSI : Predicate<"Subtarget->getGeneration() "
-                      "== AMDGPUSubtarget::SOUTHERN_ISLANDS">;
+                      "== AMDGPUSubtarget::SOUTHERN_ISLANDS">,
+           AssemblerPredicate<"FeatureSouthernIslands">;
+
 
 def has16BankLDS : Predicate<"Subtarget->getLDSBankCount() == 16">;
 def has32BankLDS : Predicate<"Subtarget->getLDSBankCount() == 32">;
 
 def SWaitMatchClass : AsmOperandClass {
   let Name = "SWaitCnt";
   let RenderMethod = "addImmOperands";
   let ParserMethod = "parseSWaitCntOps";

 //def BUFFER_ATOMIC_AND_X2 : MUBUF_X2 <mubuf<0x59, 0x68>, "buffer_atomic_and_x2", []>;
 //def BUFFER_ATOMIC_OR_X2 : MUBUF_X2 <mubuf<0x5a, 0x69>, "buffer_atomic_or_x2", []>;
 //def BUFFER_ATOMIC_XOR_X2 : MUBUF_X2 <mubuf<0x5b, 0x6a>, "buffer_atomic_xor_x2", []>;
 //def BUFFER_ATOMIC_INC_X2 : MUBUF_X2 <mubuf<0x5c, 0x6b>, "buffer_atomic_inc_x2", []>;
 //def BUFFER_ATOMIC_DEC_X2 : MUBUF_X2 <mubuf<0x5d, 0x6c>, "buffer_atomic_dec_x2", []>;
 //def BUFFER_ATOMIC_FCMPSWAP_X2 : MUBUF_X2 <mubuf<0x5e>, "buffer_atomic_fcmpswap_x2", []>; // isn't on VI
 //def BUFFER_ATOMIC_FMIN_X2 : MUBUF_X2 <mubuf<0x5f>, "buffer_atomic_fmin_x2", []>; // isn't on VI
 //def BUFFER_ATOMIC_FMAX_X2 : MUBUF_X2 <mubuf<0x60>, "buffer_atomic_fmax_x2", []>; // isn't on VI
-//def BUFFER_WBINVL1_SC : MUBUF_WBINVL1 <mubuf<0x70>, "buffer_wbinvl1_sc", []>; // isn't on CI & VI
-//def BUFFER_WBINVL1_VOL : MUBUF_WBINVL1 <mubuf<0x70, 0x3f>, "buffer_wbinvl1_vol", []>; // isn't on SI
-//def BUFFER_WBINVL1 : MUBUF_WBINVL1 <mubuf<0x71, 0x3e>, "buffer_wbinvl1", []>;
+
+let SubtargetPredicate = isSI in {
+defm BUFFER_WBINVL1_SC : MUBUF_Invalidate <mubuf<0x70>, "buffer_wbinvl1_sc", int_amdgcn_buffer_wbinvl1_sc>; // isn't on CI & VI
+}
+
+defm BUFFER_WBINVL1 : MUBUF_Invalidate <mubuf<0x71, 0x3e>, "buffer_wbinvl1", int_amdgcn_buffer_wbinvl1>;
 
 //===----------------------------------------------------------------------===//
 // MTBUF Instructions
 //===----------------------------------------------------------------------===//
 
 //def TBUFFER_LOAD_FORMAT_X : MTBUF_ <0x00000000, "tbuffer_load_format_x", []>;
 //def TBUFFER_LOAD_FORMAT_XY : MTBUF_ <0x00000001, "tbuffer_load_format_xy", []>;
 //def TBUFFER_LOAD_FORMAT_XYZ : MTBUF_ <0x00000002, "tbuffer_load_format_xyz", []>;

test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.ll

+; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
+; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+
+declare void @llvm.amdgcn.buffer.wbinvl1() #0
+
+; GCN-LABEL: {{^}}test_buffer_wbinvl1:
+; GCN-NEXT: ; BB#0:
+; SI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
+; VI-NEXT: buffer_wbinvl1 ; encoding: [0x00,0x00,0xf8,0xe0,0x00,0x00,0x00,0x00]
+; GCN-NEXT: s_endpgm
+define void @test_buffer_wbinvl1() #0 {
+  call void @llvm.amdgcn.buffer.wbinvl1()
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.sc.ll

+; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
+
+declare void @llvm.amdgcn.buffer.wbinvl1.sc() #0
+
+; SI-LABEL: {{^}}test_buffer_wbinvl1_sc:
+; SI-NEXT: ; BB#0:
+; SI-NEXT: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+; SI-NEXT: s_endpgm
+define void @test_buffer_wbinvl1_sc() #0 {
+  call void @llvm.amdgcn.buffer.wbinvl1.sc()
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll

+; RUN: llc -march=amdgcn -mcpu=bonaire -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=CI %s
+; RUN: llc -march=amdgcn -mcpu=tonga -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+
+declare void @llvm.amdgcn.buffer.wbinvl1.vol() #0
+
+; GCN-LABEL: {{^}}test_buffer_wbinvl1_vol:
+; GCN-NEXT: ; BB#0:
+; CI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+; VI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
+; GCN-NEXT: s_endpgm
+define void @test_buffer_wbinvl1_vol() #0 {
+  call void @llvm.amdgcn.buffer.wbinvl1.vol()
+  ret void
+}
+
+attributes #0 = { nounwind }

test/MC/AMDGPU/mubuf-ci.s

+// RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck -check-prefix=CI %s
+// XUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck -check-prefix=VI %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck -check-prefix=NOSI %s
+
+buffer_wbinvl1_vol
+// CI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+// XVI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00]
+// NOSI: error: instruction not supported on this GPU

test/MC/AMDGPU/mubuf-si.s

+// RUN: llvm-mc -arch=amdgcn -mcpu=SI -show-encoding %s | FileCheck %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefix=NOCI %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga %s 2>&1 | FileCheck -check-prefix=NOVI %s
+
+//===----------------------------------------------------------------------===//
+// Cache invalidation
+//===----------------------------------------------------------------------===//
+
+buffer_wbinvl1_sc
+// CHECK: buffer_wbinvl1_sc ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00]
+// NOCI: error: instruction not supported on this GPU
+// NOVI: error: instruction not supported on this GPU
+

test/MC/AMDGPU/mubuf.s

 // CHECK: buffer_store_dword v1, s[4:7], s1 ; encoding: [0x00,0x00,0x70,0xe0,0x00,0x01,0x01,0x01]
 
 buffer_store_dwordx2 v[1:2], s[4:7], s1
 // CHECK: buffer_store_dwordx2 v[1:2], s[4:7], s1 ; encoding: [0x00,0x00,0x74,0xe0,0x00,0x01,0x01,0x01]
 
 buffer_store_dwordx4 v[1:4], s[4:7], s1
 // CHECK: buffer_store_dwordx4 v[1:4], s[4:7], s1 ; encoding: [0x00,0x00,0x78,0xe0,0x00,0x01,0x01,0x01]
 
+//===----------------------------------------------------------------------===//
+// Cache invalidation
+//===----------------------------------------------------------------------===//
+
+buffer_wbinvl1
+// CHECK: buffer_wbinvl1   ; encoding: [0x00,0x00,0xc4,0xe1,0x00,0x00,0x00,0x00]
+
 // TODO: Atomics