[AMDGPU] Add wave barrier builtin

[rampitec]

Summary

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.

[rampitec]

Renamed builtin to remove s_ as there is no corresponding instruction.

[arsenm]

LGTM

cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
 
 //===----------------------------------------------------------------------===//
 // 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")
 BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
 BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
 BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc")
 BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc")

cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_s_barrier
 // CHECK: call void @llvm.amdgcn.s.barrier(
 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)
 {
   *out = __builtin_amdgcn_s_memtime();
 }
 
 // CHECK-LABEL: @test_s_sleep