[llvm] r287007 - [AMDGPU] Add wave barrier builtin

[Stanislav Mekhanoshin via llvm-commits]

Author: rampitec
Date: Tue Nov 15 13:00:15 2016
New Revision: 287007

URL: http://llvm.org/viewvc/llvm-project?rev=287007&view=rev
Log:
[AMDGPU] Add wave barrier builtin

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.

Differential Revision: https://reviews.llvm.org/D26585

Added:
    llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll
Modified:
    llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
    llvm/trunk/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp
    llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp
    llvm/trunk/lib/Target/AMDGPU/SIInstructions.td

Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=287007&r1=287006&r2=287007&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Tue Nov 15 13:00:15 2016
@@ -107,6 +107,9 @@ def int_amdgcn_dispatch_id :
 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<

Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp?rev=287007&r1=287006&r2=287007&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp Tue Nov 15 13:00:15 2016
@@ -196,6 +196,12 @@ void AMDGPUAsmPrinter::EmitInstruction(c
       return;
     }
 
+    if (MI->getOpcode() == AMDGPU::WAVE_BARRIER) {
+      if (isVerbose())
+        OutStreamer->emitRawComment(" wave barrier");
+      return;
+    }
+
     MCInst TmpInst;
     MCInstLowering.lower(MI, TmpInst);
     EmitToStreamer(*OutStreamer, TmpInst);

Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp?rev=287007&r1=287006&r2=287007&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp Tue Nov 15 13:00:15 2016
@@ -3492,6 +3492,9 @@ unsigned SIInstrInfo::getInstSizeInBytes
   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)) {

Modified: llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstructions.td?rev=287007&r1=287006&r2=287007&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstructions.td Tue Nov 15 13:00:15 2016
@@ -137,6 +137,17 @@ def S_ANDN2_B64_term : PseudoInstSI<(out
   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.
 

Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll?rev=287007&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll Tue Nov 15 13:00:15 2016
@@ -0,0 +1,16 @@
+; 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 }