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

[Stanislav Mekhanoshin via llvm-commits]

rampitec created this revision.
rampitec added a reviewer: vpykhtin.
rampitec added a subscriber: llvm-commits.
rampitec set the repository for this revision to rL LLVM.
Herald added a reviewer: tstellarAMD.
Herald added subscribers: tony-tye, yaxunl, nhaehnle, wdng, kzhuravl.

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 lowering to machine instructions.


Repository:
  rL LLVM

https://reviews.llvm.org/D26584

Files:
  include/clang/Basic/BuiltinsAMDGPU.def
  test/CodeGenOpenCL/builtins-amdgcn.cl


Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -270,6 +270,13 @@
   __builtin_amdgcn_s_barrier();
 }
 
+// CHECK-LABEL: @test_s_wave_barrier
+// CHECK: call void @llvm.amdgcn.s.wave.barrier(
+void test_s_wave_barrier()
+{
+  __builtin_amdgcn_s_wave_barrier();
+}
+
 // CHECK-LABEL: @test_s_memtime
 // CHECK: call i64 @llvm.amdgcn.s.memtime()
 void test_s_memtime(global ulong* out)
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -36,6 +36,7 @@
 // Instruction builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
+BUILTIN(__builtin_amdgcn_s_wave_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D26584.77738.patch
Type: text/x-patch
Size: 1117 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20161113/9a0ae545/attachment.bin>