Skip to content

Commit a450dc8

Browse files
authored
[libclc] Implement __clc_get_local_size/__clc_get_max_sub_group_size for amdgcn (#153785)
This simplifies downstream refactoring of libspirv workitem function in https://github.com/intel/llvm/tree/sycl/libclc/libspirv/lib/generic
1 parent 4d2288d commit a450dc8

File tree

3 files changed

+37
-0
lines changed

3 files changed

+37
-0
lines changed

libclc/clc/lib/amdgcn/SOURCES

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,4 +5,6 @@ workitem/clc_get_global_offset.cl
55
workitem/clc_get_global_size.cl
66
workitem/clc_get_group_id.cl
77
workitem/clc_get_local_id.cl
8+
workitem/clc_get_local_size.cl
9+
workitem/clc_get_max_sub_group_size.cl
810
workitem/clc_get_work_dim.cl
Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clc/workitem/clc_get_local_size.h>
10+
11+
_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) {
12+
switch (dim) {
13+
case 0:
14+
return __builtin_amdgcn_workgroup_size_x();
15+
case 1:
16+
return __builtin_amdgcn_workgroup_size_y();
17+
case 2:
18+
return __builtin_amdgcn_workgroup_size_z();
19+
default:
20+
return 1;
21+
}
22+
}
Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include <clc/workitem/clc_get_max_sub_group_size.h>
10+
11+
_CLC_OVERLOAD _CLC_DEF uint __clc_get_max_sub_group_size() {
12+
return __builtin_amdgcn_wavefrontsize();
13+
}

0 commit comments

Comments
 (0)