Author: rampitec Date: Tue Nov 15 12:58:03 2016 New Revision: 287006 URL: http://llvm.org/viewvc/llvm-project?rev=287006&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/D26584 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=287006&r1=287005&r2=287006&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Tue Nov 15 12:58:03 2016 @@ -36,6 +36,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, // 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") Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=287006&r1=287005&r2=287006&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Tue Nov 15 12:58:03 2016 @@ -270,6 +270,13 @@ void test_s_barrier() __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) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
