[AMDGPU] Add wave barrier builtin

[rampitec]

Summary

The wave barrier represents the discardable barrier. Its main purpose is to carry convergent attribute, thus preventing illegal CFG optimizations.
All lanes in a wave come to convergence point simultaneously with SIMT, thus no special instruction is needed in the ISA.
The barrier is discarded during code generation.

[arsenm]

This also needs a corresponding convergent pseudoinstruction otherwise machine CFG changes could do the same things

include/llvm/IR/IntrinsicsAMDGPU.td

110 ↗

(On Diff #77739)

There is no scalar instruction, so this should drop the _s_

[rampitec]

Renamed builtin to remove s_ as there is no corresponding instruction.

[rampitec]

In D26585#593881, @arsenm wrote:

This also needs a corresponding convergent pseudoinstruction otherwise machine CFG changes could do the same things

I though this must the next step. One thing at a time.

[rampitec]

Added pseudo instruction WAVE_BARRIER to prevent MachineInst optimizations as well.

[• tstellarAMD]

lib/Target/AMDGPU/SIInstrInfo.cpp

903–906 ↗

(On Diff #77851)

Does it matter that Post-RA scheduler runs after the Post-RA Pseudo Expansion pass?

[rampitec]

lib/Target/AMDGPU/SIInstrInfo.cpp

903–906 ↗

(On Diff #77851)

I do not think that is important. Scheduler will not move an instruction to a non-equivalent control flow block.

[arsenm]

lib/Target/AMDGPU/SIInstrInfo.cpp

903–906 ↗

(On Diff #77851)

I think this should be emitted in the asm printer as a comment rather than deleting it during codegen passes

[rampitec]

Wave barrier replaced with comment at emission.

[arsenm]

lib/Target/AMDGPU/SIInstructions.td

142 ↗

(On Diff #77894)

I think this shouldn't be set. I think this will end up counting this as 4 bytes instead of 0

143 ↗

(On Diff #77894)

This should probably be empty and set to hasNoSchedulingInfo like si_mask_branch

[rampitec]

Fixed size calculation for new instructions.
Refined attributes.

[rampitec]

lib/Target/AMDGPU/SIInstructions.td

142 ↗

(On Diff #77894)

If it does not have SALU, then it is VALU, and that is what happens:

• Bad machine code: VALU instruction does not implicitly read exec mask ***

I have added it to SIInstrInfo::getInstSizeInBytes() instead.

[arsenm]

LGTM

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

 
 //===----------------------------------------------------------------------===//
 // Instruction Intrinsics
 //===----------------------------------------------------------------------===//
 
 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
   Intrinsic<[], [], [IntrConvergent]>;
 
+def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">,
+  Intrinsic<[], [], [IntrConvergent]>;
+
 def int_amdgcn_s_waitcnt : Intrinsic<[], [llvm_i32_ty], []>;
 
 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],

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

if (MI->isBundle()) {
     }
 
     if (MI->getOpcode() == AMDGPU::SI_RETURN) {
       if (isVerbose())
         OutStreamer->emitRawComment(" return");
       return;
     }
 
+    if (MI->getOpcode() == AMDGPU::WAVE_BARRIER) {
+      if (isVerbose())
+        OutStreamer->emitRawComment(" wave barrier");
+      return;
+    }
+
     MCInst TmpInst;
     MCInstLowering.lower(MI, TmpInst);
     EmitToStreamer(*OutStreamer, TmpInst);
 
     if (STI.dumpCode()) {
       // Disassemble instruction/operands to text.
       DisasmLines.resize(DisasmLines.size() + 1);
       std::string &DisasmLine = DisasmLines.back();

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

unsigned SIInstrInfo::getInstSizeInBytes(const MachineInstr &MI) const {
   // If we have a definitive size, we can use it. Otherwise we need to inspect
   // the operands to know the size.
   //
   // FIXME: Instructions that have a base 32-bit encoding report their size as
   // 4, even though they are really 8 bytes if they have a literal operand.
   if (DescSize != 0 && DescSize != 4)
     return DescSize;
 
+  if (Opc == AMDGPU::WAVE_BARRIER)
+    return 0;
+
   // 4-byte instructions may have a 32-bit literal encoded after them. Check
   // operands that coud ever be literals.
   if (isVALU(MI) || isSALU(MI)) {
     if (isFixedSize(MI)) {
       assert(DescSize == 4);
       return DescSize;
     }
 

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

 
 def S_ANDN2_B64_term : PseudoInstSI<(outs SReg_64:$dst),
    (ins SSrc_b64:$src0, SSrc_b64:$src1)> {
   let SALU = 1;
   let isAsCheapAsAMove = 1;
   let isTerminator = 1;
 }
 
+def WAVE_BARRIER : SPseudoInstSI<(outs), (ins),
+  [(int_amdgcn_wave_barrier)]> {
+  let SchedRW = [];
+  let hasNoSchedulingInfo = 1;
+  let hasSideEffects = 1;
+  let mayLoad = 1;
+  let mayStore = 1;
+  let isBarrier = 1;
+  let isConvergent = 1;
+}
+
 // SI pseudo instructions. These are used by the CFG structurizer pass
 // and should be lowered to ISA instructions prior to codegen.
 
 // Dummy terminator instruction to use after control flow instructions
 // replaced with exec mask operations.
 def SI_MASK_BRANCH : PseudoInstSI <
   (outs), (ins brtarget:$target)> {
   let isBranch = 0;

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll

+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_wave_barrier:
+; GCN-DAG: ; wave barrier
+; GCN-NOT: s_barrier
+
+define void @test_wave_barrier() #0 {
+entry:
+  call void @llvm.amdgcn.wave.barrier() #1
+  ret void
+}
+
+declare void @llvm.amdgcn.wave.barrier() #1
+
+attributes #0 = { nounwind }
+attributes #1 = { convergent nounwind }