[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

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

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

143

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

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

include/llvm/IR/IntrinsicsAMDGPU.td

 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<

lib/Target/AMDGPU/AMDGPUMCInstLower.cpp

       return;
     }
 
+    if (MI->getOpcode() == AMDGPU::WAVE_BARRIER) {
+      if (isVerbose())
+        OutStreamer->emitRawComment(" wave barrier");
+      return;
+    }
+
     MCInst TmpInst;
     MCInstLowering.lower(MI, TmpInst);
     EmitToStreamer(*OutStreamer, TmpInst);

lib/Target/AMDGPU/SIInstructions.td

   let isTerminator = 1;
 }
 
+def WAVE_BARRIER : PseudoInstSI<(outs), (ins),
+  [(int_amdgcn_wave_barrier)]> {
+  let SALU = 1;
+  let SchedRW = [WriteBarrier];
+  let hasSideEffects = 1;
+  let mayLoad = 1;
+  let mayStore = 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.
 

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 }