[PATCH] D26585: [AMDGPU] Add wave barrier builtin

[Stanislav Mekhanoshin via llvm-commits]

rampitec updated the summary for this revision.
rampitec updated this revision to Diff 77851.
rampitec added a comment.

Added pseudo instruction WAVE_BARRIER to prevent MachineInst optimizations as well.


Repository:
  rL LLVM

https://reviews.llvm.org/D26585

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


Index: test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.wave.barrier.ll
@@ -0,0 +1,15 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test_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 }
Index: lib/Target/AMDGPU/SIInstructions.td
===================================================================
--- lib/Target/AMDGPU/SIInstructions.td
+++ lib/Target/AMDGPU/SIInstructions.td
@@ -137,6 +137,16 @@
   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.
 
Index: lib/Target/AMDGPU/SIInstrInfo.cpp
===================================================================
--- lib/Target/AMDGPU/SIInstrInfo.cpp
+++ lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -900,6 +900,10 @@
     MI.eraseFromParent();
     break;
   }
+  case AMDGPU::WAVE_BARRIER: {
+    MI.eraseFromParent();
+    break;
+  }
   }
   return true;
 }
Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -107,6 +107,9 @@
 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<


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D26585.77851.patch
Type: text/x-patch
Size: 2086 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20161114/ec6d993b/attachment.bin>