AMDGPU/SI: add llvm.amdgcn.buffer.load/store.format intrinsics

[nhaehnle]

Summary

They correspond to BUFFER_LOAD/STORE_FORMAT_XYZW and will be used by Mesa
to implement the GL_ARB_shader_image_load_store extension.

The intention is that for llvm.amdgcn.buffer.load.format, LLVM will decide
whether one of the _X/_XY/_XYZ opcodes can be used (similar to image sampling
and loads). However, this is not currently implemented.

For llvm.amdgcn.buffer.store, LLVM cannot decide to use one of the "smaller"
opcodes and therefore the intrinsic is overloaded. Currently, only the v4f32
is actually implemented since GLSL also only has a vec4 variant of the store
instructions, although it's conceivable that Mesa will want to be smarter
about this in the future.

BUFFER_LOAD_FORMAT_XYZW is already exposed via llvm.SI.vs.load.input, which
has a legacy name, pretends not to access memory, and does not capture the
full flexibility of the instruction.

The situation with the intrinsic attributes for buffer.store.format is
unfortunate. Specifying only IntrReadWriteArgMem will eventually be correct,
but right now LLVM doesn't understand buffer resources as pointers, so might
possibly miscompile something in an optimization pass. On the other hand,
not specifying any attributes at all (which is conservative and correct)
clashes in TableGen pattern generation with the fact that BUFFER_STORE_*
instruction have hasSideEffects = 0 (quite reasonably), and I'm hesitant to
change that since those instructions are already used by compute.

[• tstellarAMD]

LGTM.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

183–190 ↗

(On Diff #48013)

I don't think it makes sense to directly expose all fields. TFE for example changes the register class of the result. The address fields also seem like they shouldn't be directly exposed and should use normal looking addressing mode matching.

204 ↗

(On Diff #48013)

I think this will break because the intrinsic doesn't have any pointer arguments

[nhaehnle]

Removed the tfe argument and use more conservative memory access attributes
by overriding mayLoad and hasSideEffects only for buffer_store_format_*.

About the index and offset arguments: index can't be merged since it uses
the stride in the buffer resource. The various offsets could be merged, but
it seems a lot of fighting with TableGen for little benefit.

If anything, I would try to unify the patterns for load and store before that,
but I couldn't figure out how to do that. The natural idea is to have a
multiclass BufferLoadStorePatterns and use !cast<MUBUF>(name#_OFFSET) etc.
to get the instructions. The problem is that load and store have different
numbers of arguments, and I couldn't figure out how to concat the operand
lists of dag nodes.

Unless there are suggestions for this or other strong objections, I'd rather
commit this as is.

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

def int_amdgcn_image_atomic_cmpswap : Intrinsic <
    llvm_i32_ty,       // cmp(VGPR)
    llvm_anyint_ty,    // vaddr(VGPR)
    llvm_v8i32_ty,     // rsrc(SGPR)
    llvm_i1_ty,        // r128(imm)
    llvm_i1_ty,        // da(imm)
    llvm_i1_ty],       // slc(imm)
   []>;
 
+def int_amdgcn_buffer_load_format : Intrinsic <
+  [llvm_v4f32_ty],
+  [llvm_v4i32_ty,     // rsrc(SGPR)
+   llvm_i32_ty,       // soffset(SGPR)
+   llvm_i32_ty,       // offset(imm)
+   llvm_i32_ty,       // vindex(VGPR)
+   llvm_i32_ty,       // voffset(VGPR)
+   llvm_i1_ty,        // glc(imm)
+   llvm_i1_ty],       // slc(imm)
+  [IntrReadMem]>;
+
+def int_amdgcn_buffer_store_format : Intrinsic <
+  [],
+  [llvm_anyfloat_ty,  // vdata(VGPR) -- can currently only select v4f32
+   llvm_v4i32_ty,     // rsrc(SGPR)
+   llvm_i32_ty,       // soffset(SGPR)
+   llvm_i32_ty,       // offset(imm)
+   llvm_i32_ty,       // vindex(VGPR)
+   llvm_i32_ty,       // voffset(VGPR)
+   llvm_i1_ty,        // glc(imm)
+   llvm_i1_ty],       // slc(imm)
+  []>;
 
 def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
   "__builtin_amdgcn_read_workdim">;
 
 
 def int_amdgcn_buffer_wbinvl1_sc :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
   Intrinsic<[], [], []>;

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

defm BUFFER_LOAD_FORMAT_XY : MUBUF_Load_Helper <
   mubuf<0x01>, "buffer_load_format_xy", VReg_64
 >;
 defm BUFFER_LOAD_FORMAT_XYZ : MUBUF_Load_Helper <
   mubuf<0x02>, "buffer_load_format_xyz", VReg_96
 >;
 defm BUFFER_LOAD_FORMAT_XYZW : MUBUF_Load_Helper <
   mubuf<0x03>, "buffer_load_format_xyzw", VReg_128
 >;
+// Without mayLoad and hasSideEffects, TableGen complains about the pattern
+// matching llvm.amdgcn.buffer.store.format. Eventually, we'll need a way
+// to express the effects of the intrinsic more precisely.
+let mayLoad = 1, hasSideEffects = 1 in {
   defm BUFFER_STORE_FORMAT_X : MUBUF_Store_Helper <
     mubuf<0x04>, "buffer_store_format_x", VGPR_32
   >;
   defm BUFFER_STORE_FORMAT_XY : MUBUF_Store_Helper <
     mubuf<0x05>, "buffer_store_format_xy", VReg_64
   >;
   defm BUFFER_STORE_FORMAT_XYZ : MUBUF_Store_Helper <
     mubuf<0x06>, "buffer_store_format_xyz", VReg_96
   >;
   defm BUFFER_STORE_FORMAT_XYZW : MUBUF_Store_Helper <
     mubuf<0x07>, "buffer_store_format_xyzw", VReg_128
   >;
+}
 defm BUFFER_LOAD_UBYTE : MUBUF_Load_Helper <
   mubuf<0x08, 0x10>, "buffer_load_ubyte", VGPR_32, i32, az_extloadi8_global
 >;
 defm BUFFER_LOAD_SBYTE : MUBUF_Load_Helper <
   mubuf<0x09, 0x11>, "buffer_load_sbyte", VGPR_32, i32, sextloadi8_global
 >;
 defm BUFFER_LOAD_USHORT : MUBUF_Load_Helper <
   mubuf<0x0a, 0x12>, "buffer_load_ushort", VGPR_32, i32, az_extloadi16_global

 def : Pat <
   (int_SI_export imm:$en, imm:$vm, imm:$done, imm:$tgt, imm:$compr,
                  f32:$src0, f32:$src1, f32:$src2, f32:$src3),
   (EXP imm:$en, imm:$tgt, imm:$compr, imm:$done, imm:$vm,
        $src0, $src1, $src2, $src3)
 >;
 
 //===----------------------------------------------------------------------===//
+// buffer_load/store_format patterns
+//===----------------------------------------------------------------------===//
+def : Pat<
+  (int_amdgcn_buffer_load_format v4i32:$rsrc, i32:$soffset, imm:$offset, 0, 0,
+                                 imm:$glc, imm:$slc),
+  (BUFFER_LOAD_FORMAT_XYZW_OFFSET $rsrc, $soffset, (as_i16imm $offset),
+    (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_load_format v4i32:$rsrc, i32:$soffset, imm:$offset, i32:$vindex, 0,
+                                 imm:$glc, imm:$slc),
+  (BUFFER_LOAD_FORMAT_XYZW_IDXEN $vindex, $rsrc, $soffset, (as_i16imm $offset),
+    (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_load_format v4i32:$rsrc, i32:$soffset, imm:$offset, 0, i32:$voffset,
+                                 imm:$glc, imm:$slc),
+  (BUFFER_LOAD_FORMAT_XYZW_OFFEN $voffset, $rsrc, $soffset, (as_i16imm $offset),
+    (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_load_format v4i32:$rsrc, i32:$soffset, imm:$offset, i32:$vindex, i32:$voffset,
+                                 imm:$glc, imm:$slc),
+  (BUFFER_LOAD_FORMAT_XYZW_BOTHEN
+    (REG_SEQUENCE VReg_64, $vindex, sub0, $voffset, sub1),
+    $rsrc, $soffset, (as_i16imm $offset),
+    (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_store_format v4f32:$vdata, v4i32:$rsrc,
+                                  i32:$soffset, imm:$offset, 0, 0,
+                                  imm:$glc, imm:$slc),
+  (BUFFER_STORE_FORMAT_XYZW_OFFSET $vdata, $rsrc, $soffset, (as_i16imm $offset),
+    (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_store_format v4f32:$vdata, v4i32:$rsrc,
+                                 i32:$soffset, imm:$offset, i32:$vindex, 0,
+                                 imm:$glc, imm:$slc),
+  (BUFFER_STORE_FORMAT_XYZW_IDXEN $vdata, $vindex, $rsrc, $soffset,
+    (as_i16imm $offset), (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_store_format v4f32:$vdata, v4i32:$rsrc,
+                                  i32:$soffset, imm:$offset, 0, i32:$voffset,
+                                  imm:$glc, imm:$slc),
+  (BUFFER_STORE_FORMAT_XYZW_OFFEN $vdata, $voffset, $rsrc, $soffset,
+    (as_i16imm $offset), (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+def : Pat<
+  (int_amdgcn_buffer_store_format v4f32:$vdata, v4i32:$rsrc, i32:$soffset,
+                                  imm:$offset, i32:$vindex, i32:$voffset,
+                                  imm:$glc, imm:$slc),
+  (BUFFER_STORE_FORMAT_XYZW_BOTHEN
+    $vdata,
+    (REG_SEQUENCE VReg_64, $vindex, sub0, $voffset, sub1),
+    $rsrc, $soffset, (as_i16imm $offset),
+    (as_i1imm $glc), (as_i1imm $slc), 0)
+>;
+
+//===----------------------------------------------------------------------===//
 // S_GETREG_B32 Intrinsic Pattern.
 //===----------------------------------------------------------------------===//
 def : Pat <
   (int_amdgcn_s_getreg imm:$simm16),
   (S_GETREG_B32 (as_i16imm $simm16))
 >;
 
 //===----------------------------------------------------------------------===//

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

void SIRegisterInfo::buildScratchLoadStore(MachineBasicBlock::iterator MI,
 
   MachineBasicBlock *MBB = MI->getParent();
   MachineFunction *MF = MI->getParent()->getParent();
   MachineRegisterInfo &MRI = MF->getRegInfo();
   const SIInstrInfo *TII =
       static_cast<const SIInstrInfo *>(MF->getSubtarget().getInstrInfo());
   LLVMContext &Ctx = MF->getFunction()->getContext();
   DebugLoc DL = MI->getDebugLoc();
-  bool IsLoad = TII->get(LoadStoreOp).mayLoad();
+  bool IsStore = TII->get(LoadStoreOp).mayStore();
 
   bool RanOutOfSGPRs = false;
   bool Scavenged = false;
   unsigned SOffset = ScratchOffset;
 
   unsigned NumSubRegs = getNumSubRegsForSpillOp(MI->getOpcode());
   unsigned Size = NumSubRegs * 4;
 
unsigned SubReg = NumSubRegs > 1 ?
         getPhysRegSubReg(Value, &AMDGPU::VGPR_32RegClass, i) :
         Value;
 
     unsigned SOffsetRegState = 0;
     if (i + 1 == e && Scavenged)
       SOffsetRegState |= RegState::Kill;
 
     BuildMI(*MBB, MI, DL, TII->get(LoadStoreOp))
-      .addReg(SubReg, getDefRegState(IsLoad))
+      .addReg(SubReg, getDefRegState(!IsStore))
       .addReg(ScratchRsrcReg)
       .addReg(SOffset, SOffsetRegState)
       .addImm(Offset)
       .addImm(0) // glc
       .addImm(0) // slc
       .addImm(0) // tfe
-      .addReg(Value, RegState::Implicit | getDefRegState(IsLoad))
+      .addReg(Value, RegState::Implicit | getDefRegState(!IsStore))
       .setMemRefs(MI->memoperands_begin(), MI->memoperands_end());
   }
 }
 
 void SIRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator MI,
                                         int SPAdj, unsigned FIOperandNum,
                                         RegScavenger *RS) const {
   MachineFunction *MF = MI->getParent()->getParent();

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.load.format.ll

+;RUN: llc < %s -march=amdgcn -mcpu=verde -verify-machineinstrs | FileCheck %s
+;RUN: llc < %s -march=amdgcn -mcpu=tonga -verify-machineinstrs | FileCheck %s
+
+;CHECK-LABEL: {{^}}buffer_load:
+;CHECK: buffer_load_format_xyzw v[0:3], s[0:3], s4
+;CHECK: buffer_load_format_xyzw v[4:7], s[0:3], s4 glc
+;CHECK: buffer_load_format_xyzw v[8:11], s[0:3], s4 slc
+;CHECK: s_waitcnt
+define {<4 x float>, <4 x float>, <4 x float>} @buffer_load(<4 x i32> inreg, i32 inreg) #0 {
+main_body:
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 %1, i32 0, i32 0, i32 0, i1 0, i1 0)
+  %data_glc = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 %1, i32 0, i32 0, i32 0, i1 1, i1 0)
+  %data_slc = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 %1, i32 0, i32 0, i32 0, i1 0, i1 1)
+  %r0 = insertvalue {<4 x float>, <4 x float>, <4 x float>} undef, <4 x float> %data, 0
+  %r1 = insertvalue {<4 x float>, <4 x float>, <4 x float>} %r0, <4 x float> %data_glc, 1
+  %r2 = insertvalue {<4 x float>, <4 x float>, <4 x float>} %r1, <4 x float> %data_slc, 2
+  ret {<4 x float>, <4 x float>, <4 x float>} %r2
+}
+
+;CHECK-LABEL: {{^}}buffer_load_immoffs:
+;CHECK: buffer_load_format_xyzw v[0:3], s[0:3], s4 offset:42
+;CHECK: s_waitcnt
+define <4 x float> @buffer_load_immoffs(<4 x i32> inreg, i32 inreg) #0 {
+main_body:
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 %1, i32 42, i32 0, i32 0, i1 0, i1 0)
+  ret <4 x float> %data
+}
+
+;CHECK-LABEL: {{^}}buffer_load_idx:
+;CHECK: buffer_load_format_xyzw v[0:3], v0, s[0:3], 0 idxen
+;CHECK: s_waitcnt
+define <4 x float> @buffer_load_idx(<4 x i32> inreg, i32) #0 {
+main_body:
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 0, i32 0, i32 %1, i32 0, i1 0, i1 0)
+  ret <4 x float> %data
+}
+
+;CHECK-LABEL: {{^}}buffer_load_ofs:
+;CHECK: buffer_load_format_xyzw v[0:3], v0, s[0:3], 0 offen
+;CHECK: s_waitcnt
+define <4 x float> @buffer_load_ofs(<4 x i32> inreg, i32) #0 {
+main_body:
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 0, i32 0, i32 0, i32 %1, i1 0, i1 0)
+  ret <4 x float> %data
+}
+
+;CHECK-LABEL: {{^}}buffer_load_both:
+;CHECK: buffer_load_format_xyzw v[0:3], v[0:1], s[0:3], 0 idxen offen
+;CHECK: s_waitcnt
+define <4 x float> @buffer_load_both(<4 x i32> inreg, i32, i32) #0 {
+main_body:
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 0, i32 0, i32 %1, i32 %2, i1 0, i1 0)
+  ret <4 x float> %data
+}
+
+;CHECK-LABEL: {{^}}buffer_load_both_reversed:
+;CHECK: v_mov_b32_e32 v2, v0
+;CHECK: buffer_load_format_xyzw v[0:3], v[1:2], s[0:3], 0 idxen offen
+;CHECK: s_waitcnt
+define <4 x float> @buffer_load_both_reversed(<4 x i32> inreg, i32, i32) #0 {
+main_body:
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 0, i32 0, i32 %2, i32 %1, i1 0, i1 0)
+  ret <4 x float> %data
+}
+
+declare <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32>, i32, i32, i32, i32, i1, i1) #1
+
+attributes #0 = { "ShaderType"="0" }
+attributes #1 = { nounwind readonly }

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.store.format.ll

+;RUN: llc < %s -march=amdgcn -mcpu=verde -verify-machineinstrs | FileCheck %s
+;RUN: llc < %s -march=amdgcn -mcpu=tonga -verify-machineinstrs | FileCheck %s
+
+;CHECK-LABEL: {{^}}buffer_store:
+;CHECK: buffer_store_format_xyzw v[0:3], s[0:3], s4
+;CHECK: buffer_store_format_xyzw v[4:7], s[0:3], s4 glc
+;CHECK: buffer_store_format_xyzw v[8:11], s[0:3], s4 slc
+define void @buffer_store(<4 x i32> inreg, i32 inreg, <4 x float>, <4 x float>, <4 x float>) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 %1, i32 0, i32 0, i32 0, i1 0, i1 0)
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %3, <4 x i32> %0, i32 %1, i32 0, i32 0, i32 0, i1 1, i1 0)
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %4, <4 x i32> %0, i32 %1, i32 0, i32 0, i32 0, i1 0, i1 1)
+  ret void
+}
+
+;CHECK-LABEL: {{^}}buffer_store_immoffs:
+;CHECK: buffer_store_format_xyzw v[0:3], s[0:3], s4 offset:42
+define void @buffer_store_immoffs(<4 x i32> inreg, i32 inreg, <4 x float>) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 %1, i32 42, i32 0, i32 0, i1 0, i1 0)
+  ret void
+}
+
+;CHECK-LABEL: {{^}}buffer_store_idx:
+;CHECK: buffer_store_format_xyzw v[0:3], v4, s[0:3], 0 idxen
+define void @buffer_store_idx(<4 x i32> inreg, i32 inreg, <4 x float>, i32) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 0, i32 0, i32 %3, i32 0, i1 0, i1 0)
+  ret void
+}
+
+;CHECK-LABEL: {{^}}buffer_store_ofs:
+;CHECK: buffer_store_format_xyzw v[0:3], v4, s[0:3], 0 offen
+define void @buffer_store_ofs(<4 x i32> inreg, i32 inreg, <4 x float>, i32) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 0, i32 0, i32 0, i32 %3, i1 0, i1 0)
+  ret void
+}
+
+;CHECK-LABEL: {{^}}buffer_store_both:
+;CHECK: buffer_store_format_xyzw v[0:3], v[4:5], s[0:3], 0 idxen offen
+define void @buffer_store_both(<4 x i32> inreg, i32 inreg, <4 x float>, i32, i32) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 0, i32 0, i32 %3, i32 %4, i1 0, i1 0)
+  ret void
+}
+
+;CHECK-LABEL: {{^}}buffer_store_both_reversed:
+;CHECK: v_mov_b32_e32 v6, v4
+;CHECK: buffer_store_format_xyzw v[0:3], v[5:6], s[0:3], 0 idxen offen
+define void @buffer_store_both_reversed(<4 x i32> inreg, i32 inreg, <4 x float>, i32, i32) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 0, i32 0, i32 %4, i32 %3, i1 0, i1 0)
+  ret void
+}
+
+; Ideally, the register allocator would avoid the wait here
+;
+;CHECK-LABEL: {{^}}buffer_store_wait:
+;CHECK: buffer_store_format_xyzw v[0:3], v4, s[0:3], 0 idxen
+;CHECK: s_waitcnt vmcnt(0) expcnt(0)
+;CHECK: buffer_load_format_xyzw v[0:3], v5, s[0:3], 0 idxen
+;CHECK: s_waitcnt vmcnt(0)
+;CHECK: buffer_store_format_xyzw v[0:3], v6, s[0:3], 0 idxen
+define void @buffer_store_wait(<4 x i32> inreg, i32 inreg, <4 x float>, i32, i32, i32) #0 {
+main_body:
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %2, <4 x i32> %0, i32 0, i32 0, i32 %3, i32 0, i1 0, i1 0)
+  %data = call <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32> %0, i32 0, i32 0, i32 %4, i32 0, i1 0, i1 0)
+  call void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float> %data, <4 x i32> %0, i32 0, i32 0, i32 %5, i32 0, i1 0, i1 0)
+  ret void
+}
+
+declare void @llvm.amdgcn.buffer.store.format.v4f32(<4 x float>, <4 x i32>, i32, i32, i32, i32, i1, i1) #1
+declare <4 x float> @llvm.amdgcn.buffer.load.format(<4 x i32>, i32, i32, i32, i32, i1, i1) #2
+
+attributes #0 = { "ShaderType"="0" }
+attributes #1 = { nounwind }
+attributes #2 = { nounwind readonly }