maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[HIP] Add __builtin_amdgcn_groupstaticsize 

Differential Revision: https://reviews.llvm.org/D102403
This commit is contained in:
Aaron En Ye Shi 2021-05-12 19:53:47 +00:00
parent 8fa168fc50
commit 6a67e05a26
5 changed files with 29 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -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")

6
clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
View File

@ -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();
}

7
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
View File

@ -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();
}

8
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
View File

@ -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();
}

7
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
View File

@ -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();
}