AMDGPU: export s_waitcnt builtin

[arsenm]

test/CodeGenOpenCL/builtins-amdgcn.cl

272

Should also have a test in builtins-amdgcn-error.cl that a non-constant input is rejected

[jvesely]

force parameter to be integer immediate

[arsenm]

LGTM

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_getreg, "UiIi", "n")
+BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
 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")

test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_class_f64
 // CHECK: call i1 @llvm.amdgcn.class.f64
 void test_class_f64(global double* out, double a, int b)
 {
   *out = __builtin_amdgcn_class(a, b);
 }
 
+// CHECK-LABEL: @test_s_waitcnt
+// CHECK: call void @llvm.amdgcn.s.waitcnt(
+void test_s_waitcnt()
+{
+  __builtin_amdgcn_s_waitcnt(0);
+}
+
 // 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

test/SemaOpenCL/builtins-amdgcn-error.cl

ulong test_s_memrealtime()
   return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
 }
 
 void test_s_sleep(int x)
 {
   __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
 }
 
+void test_s_waitcnt(int x)
+{
+  __builtin_amdgcn_s_waitcnt(x); // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}}
+}
+
 void test_s_incperflevel(int x)
 {
   __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}}
 }
 
 void test_s_decperflevel(int x)
 {
   __builtin_amdgcn_s_decperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}}