AMDGPU: Add GWS instruction builtins

[rampitec]

LGTM

[arsenm]

r363684

[uweigand]

I'm now getting test suite failures like this:

/home/uweigand/llvm/llvm-head/tools/clang/test/CodeGenOpenCL/builtins-amdgcn.cl:554:3: error: cannot compile this builtin function yet
  __builtin_amdgcn_ds_gws_init(value, id);

Is this supposed to work? I'm not seeing any LLVM support for this builtin anywhere ...

include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_s_getpc, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
 BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n")
 BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n")
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
 BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
+BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n")
+BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n")
 
 // FIXME: Need to disallow constant address space.
 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")

test/CodeGenOpenCL/builtins-amdgcn.cl

 }
 
 // CHECK-LABEL: @test_ds_consume_lds(
 // CHECK: call i32 @llvm.amdgcn.ds.consume.p3i32(i32 addrspace(3)* %ptr, i1 false)
 kernel void test_ds_consume_lds(global int* out, local int* ptr) {
   *out = __builtin_amdgcn_ds_consume(ptr);
 }
 
+// CHECK-LABEL: @test_gws_init(
+// CHECK: call void @llvm.amdgcn.ds.gws.init(i32 %value, i32 %id)
+kernel void test_gws_init(uint value, uint id) {
+  __builtin_amdgcn_ds_gws_init(value, id);
+}
+
+// CHECK-LABEL: @test_gws_barrier(
+// CHECK: call void @llvm.amdgcn.ds.gws.barrier(i32 %value, i32 %id)
+kernel void test_gws_barrier(uint value, uint id) {
+  __builtin_amdgcn_ds_gws_barrier(value, id);
+}
+
 // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
 // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }
 // CHECK-DAG: ![[$EXEC]] = !{!"exec"}
 // CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"}
 // CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"}