AMDGPU/SI: Add llvm.amdgcn.s.waitcnt.all intrinsic

[nhaehnle]

Summary

So it appears that to guarantee some of the ordering requirements of a GLSL
memoryBarrier() executed in the shader, we need to emit an s_waitcnt.

(We can't use an s_barrier, because memoryBarrier() may appear anywhere in
the shader, in particular it may appear in non-uniform control flow.)

[arsenm]

What are the memory barrier requirements for GLSL? Do you need consistency between multiple workgroups?

[• tstellarAMD]

include/llvm/IR/IntrinsicsAMDGPU.td

71

I think the intrinsic should be int_amdgpu_s_waitcnt, and we should expose the configuration bits with an input argument.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

71

I'm not sure waitcnt should be directly exposed, and it should probably be a memfence intrinsic. However, I'm not clear if what is really wanted here is the cache flush intrinsics like is necessary for OpenCL 2.0

This also doesn't need to be convergent

[nhaehnle]

Yes, we need consistency between all shader invocations, which can span all the CUs and SEs on the chip. There isn't really a notion of workgroups for GLSL graphics shaders. Basically, the instruction needs to make sure that all past memory writes by the shader (actually, only 'coherent' and 'volatile' ones) are visible to all other shaders. I'm not sure about what OpenCL needs.

With this patch, the idea is to implement this by setting glc=1 on the coherent/volatile writes and using a wait. I believe (but have not tried) that an alternative would be to always use glc=0 and wait + explicitly request an L1 cache flush at the memory barrier.

Tom, do you want the numeric counts as input, or just bits that indicate whether to wait for vm/exp/lgkm?

[• tstellarAMD]

In D19203#404014, @nhaehnle wrote:

Tom, do you want the numeric counts as input, or just bits that indicate whether to wait for vm/exp/lgkm?

Just the bits. So, the input to the intrinsic and the instruction are the same.

[nhaehnle]

Changed the intrinsic to take a single argument.

Also changed the SIInsertWaits logic so that the intrinsic will be delayed
up to the next counter-incrementing instruction or the next "natural" wait
(and then merged with it).

[• tstellarAMD]

LGTM.

include/llvm/IR/IntrinsicsAMDGPU.td

 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
   "__builtin_amdgcn_workitem_id">;
 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
   "__builtin_amdgcn_workgroup_id">;
 
 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
   Intrinsic<[], [], [IntrConvergent]>;
 
+def int_amdgcn_s_waitcnt_all : Intrinsic<[], [], [IntrConvergent]>;
+
 def int_amdgcn_div_scale : Intrinsic<
   // 1st parameter: Numerator
   // 2nd parameter: Denominator
   // 3rd parameter: Constant to select select between first and
   //                second. (0 = first, 1 = second).
   [llvm_anyfloat_ty, llvm_i1_ty],
   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
   [IntrNoMem]

lib/Target/AMDGPU/SIInsertWaits.cpp

bool SIInsertWaits::runOnMachineFunction(MachineFunction &MF) {
   LastIssued = ZeroCounts;
   LastOpcodeType = OTHER;
   LastInstWritesM0 = false;
   ReturnsVoid = MF.getInfo<SIMachineFunctionInfo>()->returnsVoid();
 
   memset(&UsedRegs, 0, sizeof(UsedRegs));
   memset(&DefinedRegs, 0, sizeof(DefinedRegs));
 
+  SmallVector<MachineInstr *, 4> RemoveMI;
+
   for (MachineFunction::iterator BI = MF.begin(), BE = MF.end();
        BI != BE; ++BI) {
 
     MachineBasicBlock &MBB = *BI;
     for (MachineBasicBlock::iterator I = MBB.begin(), E = MBB.end();
          I != E; ++I) {
 
       if (ST.getGeneration() <= AMDGPUSubtarget::SEA_ISLANDS) {
for (MachineBasicBlock::iterator I = MBB.begin(), E = MBB.end();
       // instruction.
       if (ST.getGeneration() <= AMDGPUSubtarget::SOUTHERN_ISLANDS &&
           I->getOpcode() == AMDGPU::V_READFIRSTLANE_B32)
         TII->insertWaitStates(MBB, std::next(I), 4);
 
       // Wait for everything before a barrier.
       if (I->getOpcode() == AMDGPU::S_BARRIER)
         Changes |= insertWait(MBB, I, LastIssued);
-      else
+      else if (I->getOpcode() == AMDGPU::S_WAITCNT) {
+        RemoveMI.push_back(I);
+        insertWait(MBB, I, LastIssued);
+        Changes = true;
+      } else
         Changes |= insertWait(MBB, I, handleOperands(*I));
 
       pushInstruction(MBB, I);
       handleSendMsg(MBB, I);
     }
 
     // Wait for everything at the end of the MBB
     Changes |= insertWait(MBB, MBB.getFirstTerminator(), LastIssued);
   }
 
+  for (MachineInstr *I : RemoveMI)
+    I->eraseFromParent();
+
   return Changes;
 }

lib/Target/AMDGPU/SIInstructions.td

 > {
   let SchedRW = [WriteBarrier];
   let simm16 = 0;
   let mayLoad = 1;
   let mayStore = 1;
   let isConvergent = 1;
 }
 
+let mayLoad = 1, mayStore = 1, hasSideEffects = 1 in
 def S_WAITCNT : SOPP <0x0000000c, (ins WAIT_FLAG:$simm16), "s_waitcnt $simm16">;
 def S_SETHALT : SOPP <0x0000000d, (ins i16imm:$simm16), "s_sethalt $simm16">;
 
 // On SI the documentation says sleep for approximately 64 * low 2
 // bits, consistent with the reported maximum of 448. On VI the
 // maximum reported is 960 cycles, so 960 / 64 = 15 max, so is the
 // maximum really 15 on VI?
 def S_SLEEP : SOPP <0x0000000e, (ins i32imm:$simm16),
def : Pat <
   (i32 (addc i32:$src0, i32:$src1)),
   (S_ADD_U32 $src0, $src1)
 >;
 
 //===----------------------------------------------------------------------===//
 // SOPP Patterns
 //===----------------------------------------------------------------------===//
 
+def : Pat <
+  (int_amdgcn_s_waitcnt_all),
+  (S_WAITCNT)
+>;
+
 // FIXME: These should be removed eventually
 def : Pat <
   (int_AMDGPU_barrier_global),
   (S_BARRIER)
 >;
 
 def : Pat <
   (int_AMDGPU_barrier_local),

test/CodeGen/AMDGPU/llvm.amdgcn.s.waitcnt.ll

+; RUN: llc -march=amdgcn -mcpu=SI -verify-machineinstrs < %s | FileCheck -check-prefix=CHECK %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=CHECK %s
+
+; CHECK-LABEL: {{^}}test1:
+; CHECK: image_store
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0){{$}}
+; CHECK-NEXT: image_store
+; CHECK-NEXT: s_endpgm
+define amdgpu_ps void @test1(<8 x i32> inreg %rsrc, <4 x float> %d0, <4 x float> %d1, i32 %c0, i32 %c1) {
+  call void @llvm.amdgcn.image.store.i32(<4 x float> %d0, i32 %c0, <8 x i32> %rsrc, i32 15, i1 0, i1 0, i1 1, i1 0)
+  call void @llvm.amdgcn.s.waitcnt.all()
+  call void @llvm.amdgcn.image.store.i32(<4 x float> %d1, i32 %c1, <8 x i32> %rsrc, i32 15, i1 0, i1 0, i1 1, i1 0)
+  ret void
+}
+
+declare void @llvm.amdgcn.s.waitcnt.all() #0
+
+declare void @llvm.amdgcn.image.store.i32(<4 x float>, i32, <8 x i32>, i32, i1, i1, i1, i1) #0
+
+attributes #0 = { nounwind }