[PATCH] D26584: [AMDGPU] Add wave barrier builtin
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Tue Nov 15 11:07:55 PST 2016
This revision was automatically updated to reflect the committed changes.
Closed by commit rL287006: [AMDGPU] Add wave barrier builtin (authored by rampitec).
Changed prior to commit:
https://reviews.llvm.org/D26584?vs=77753&id=78036#toc
Repository:
rL LLVM
https://reviews.llvm.org/D26584
Files:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
Index: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
@@ -36,6 +36,7 @@
// Instruction builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
+BUILTIN(__builtin_amdgcn_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")
Index: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -270,6 +270,13 @@
__builtin_amdgcn_s_barrier();
}
+// CHECK-LABEL: @test_wave_barrier
+// CHECK: call void @llvm.amdgcn.wave.barrier(
+void test_wave_barrier()
+{
+ __builtin_amdgcn_wave_barrier();
+}
+
// CHECK-LABEL: @test_s_memtime
// CHECK: call i64 @llvm.amdgcn.s.memtime()
void test_s_memtime(global ulong* out)
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D26584.78036.patch
Type: text/x-patch
Size: 1167 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20161115/b160e868/attachment.bin>
More information about the llvm-commits
mailing list