[clang] 6a67e05 - [HIP] Add __builtin_amdgcn_groupstaticsize

[Aaron En Ye Shi via cfe-commits]

Author: Aaron En Ye Shi
Date: 2021-05-13T15:50:08Z
New Revision: 6a67e05a26eb5f58665bd6d063b9f389e7dd28a7

URL: https://github.com/llvm/llvm-project/commit/6a67e05a26eb5f58665bd6d063b9f389e7dd28a7
DIFF: https://github.com/llvm/llvm-project/commit/6a67e05a26eb5f58665bd6d063b9f389e7dd28a7.diff

LOG: [HIP] Add __builtin_amdgcn_groupstaticsize

Differential Revision: https://reviews.llvm.org/D102403

Added: 
    

Modified: 
    clang/include/clang/Basic/BuiltinsAMDGPU.def
    clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
    clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
    clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
    clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 7dcbf9a096961..e3782fc403bdd 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -65,6 +65,7 @@ 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")

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
index 56da7eceb6deb..15ddd17316a71 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -62,3 +62,9 @@ int test_is_shared_global(const global int* ptr) {
 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();
+}

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
index a9bbaa9c3f54f..3b21d41659bbd 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
@@ -30,3 +30,10 @@ 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();
+}

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
index 420506ec083c1..3bc1811a35b71 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
@@ -3,6 +3,7 @@
 // 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
@@ -18,3 +19,10 @@ 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();
+}

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index be594585fad6d..fe2c1692fbcc4 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -145,3 +145,10 @@ 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();
+}