[HIP] Add __builtin_amdgcn_groupstaticsize

[yaxunl]

LGTM. Thanks.

clang/include/clang/Basic/BuiltinsAMDGPU.def

 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")
 BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n")
 BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n")
 BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
+BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n")
 
 BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n")
 BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n")
 
 BUILTIN(__builtin_amdgcn_atomic_dec32, "UZiUZiD*UZiUicC*", "n")
 BUILTIN(__builtin_amdgcn_atomic_dec64, "UWiUWiD*UWiUicC*", "n")
 
 // FIXME: Need to disallow constant address space.

clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl

 }
 
 // CHECK-LABEL: @test_is_private_global(
 // CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8*
 // CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]]
 int test_is_private_global(const global int* ptr) {
   return __builtin_amdgcn_is_private(ptr);
 }
+
+// CHECK-LABEL: @test_groupstaticsize
+// CHECK: call i32 @llvm.amdgcn.groupstaticsize()
+void test_groupstaticsize(global uint* out) {
+  *out = __builtin_amdgcn_groupstaticsize();
+}

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl

 }
 
 // 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_groupstaticsize
+// CHECK: call i32 @llvm.amdgcn.groupstaticsize()
+void test_groupstaticsize(global uint* out)
+{
+  *out = __builtin_amdgcn_groupstaticsize();
+}

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
+typedef unsigned int uint;
 typedef unsigned long ulong;
 
 // CHECK-LABEL: @test_fmed3_f16
 // CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c)
 void test_fmed3_f16(global half* out, half a, half b, half c)
 {
   *out = __builtin_amdgcn_fmed3h(a, b, c);
 }
 
 // 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_groupstaticsize
+// CHECK: call i32 @llvm.amdgcn.groupstaticsize()
+void test_groupstaticsize(global uint* out)
+{
+  *out = __builtin_amdgcn_groupstaticsize();
+}

clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

 }
 
 // CHECK-LABEL: @test_perm
 // CHECK: call i32 @llvm.amdgcn.perm(i32 %a, i32 %b, i32 %s)
 void test_perm(global uint* out, uint a, uint b, uint s)
 {
   *out = __builtin_amdgcn_perm(a, b, s);
 }
+
+// CHECK-LABEL: @test_groupstaticsize
+// CHECK: call i32 @llvm.amdgcn.groupstaticsize()
+void test_groupstaticsize(global uint* out)
+{
+  *out = __builtin_amdgcn_groupstaticsize();
+}