shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[libclc] Implement __clc_get_local_size/__clc_get_max_sub_group_size for amdgcn (#153785)

Browse Source 
This simplifies downstream refactoring of libspirv workitem function in
https://github.com/intel/llvm/tree/sycl/libclc/libspirv/lib/generic
This commit is contained in:
Wenju He 2025-08-19 07:51:17 +08:00 committed by GitHub
parent 4d2288d318
commit a450dc80bf
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
3 changed files with 37 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
libclc/clc/lib/amdgcn/SOURCES
View File

@ -5,4 +5,6 @@ workitem/clc_get_global_offset.cl
workitem/clc_get_global_size.cl
workitem/clc_get_group_id.cl
workitem/clc_get_local_id.cl
workitem/clc_get_local_size.cl
workitem/clc_get_max_sub_group_size.cl
workitem/clc_get_work_dim.cl

22
libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl Normal file
View File

@ -0,0 +1,22 @@
//===----------------------------------------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#include <clc/workitem/clc_get_local_size.h>
_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) {
switch (dim) {
case 0:
return __builtin_amdgcn_workgroup_size_x();
case 1:
return __builtin_amdgcn_workgroup_size_y();
case 2:
return __builtin_amdgcn_workgroup_size_z();
default:
return 1;
}
}

13
libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl Normal file
View File

@ -0,0 +1,13 @@
//===----------------------------------------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#include <clc/workitem/clc_get_max_sub_group_size.h>
_CLC_OVERLOAD _CLC_DEF uint __clc_get_max_sub_group_size() {
return __builtin_amdgcn_wavefrontsize();
}