llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Wenju He (wenju-he) <details> <summary>Changes</summary> De-duplicate since some are already declaration in include/clc/workitem. Move subgroup workitem function implementations into individual files to align with other workitem functions. --- Full diff: https://github.com/llvm/llvm-project/pull/189328.diff 26 Files Affected: - (modified) libclc/clc/include/clc/subgroup/clc_subgroup.h (-6) - (added) libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h (+16) - (modified) libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h (+1-1) - (modified) libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h (+1-1) - (modified) libclc/clc/include/clc/workitem/clc_get_sub_group_id.h (+1-1) - (modified) libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h (+1-1) - (modified) libclc/clc/include/clc/workitem/clc_get_sub_group_size.h (+1-1) - (modified) libclc/clc/lib/amdgpu/CMakeLists.txt (+2) - (modified) libclc/clc/lib/amdgpu/subgroup/clc_sub_group_scan.cl (+1) - (modified) libclc/clc/lib/amdgpu/subgroup/clc_subgroup.cl (-10) - (added) libclc/clc/lib/amdgpu/workitem/clc_get_enqueued_num_sub_groups.cl (+16) - (modified) libclc/clc/lib/amdgpu/workitem/clc_get_num_sub_groups.cl (+1-1) - (added) libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_local_id.cl (+13) - (modified) libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_size.cl (+1) - (modified) libclc/clc/lib/generic/collective/clc_work_group_any_all.cl (+3) - (modified) libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl (+1) - (modified) libclc/clc/lib/generic/collective/clc_work_group_scan.cl (+1) - (modified) libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl (+1) - (modified) libclc/opencl/lib/generic/CMakeLists.txt (+6) - (modified) libclc/opencl/lib/generic/subgroup/subgroup.cl (-24) - (added) libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl (+13) - (added) libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl (+13) - (added) libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl (+13) - (added) libclc/opencl/lib/generic/workitem/get_sub_group_id.cl (+13) - (added) libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl (+13) - (added) libclc/opencl/lib/generic/workitem/get_sub_group_size.cl (+13) ``````````diff diff --git a/libclc/clc/include/clc/subgroup/clc_subgroup.h b/libclc/clc/include/clc/subgroup/clc_subgroup.h index f0a2a11d48445..133ba33644120 100644 --- a/libclc/clc/include/clc/subgroup/clc_subgroup.h +++ b/libclc/clc/include/clc/subgroup/clc_subgroup.h @@ -11,12 +11,6 @@ #include "clc/internal/clc.h" -_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_size(void); -_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_max_sub_group_size(void); -_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_num_sub_groups(void); -_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_enqueued_num_sub_groups(void); -_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_id(void); -_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_local_id(void); _CLC_DECL _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_all(int x); _CLC_DECL _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_any(int x); diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h new file mode 100644 index 0000000000000..14afdd80ca11f --- /dev/null +++ b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#ifndef __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__ +#define __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__ + +#include "clc/internal/clc.h" + +_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_enqueued_num_sub_groups(void); + +#endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__ diff --git a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h index a5c98aeba94b8..d5a3a13945e7b 100644 --- a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h +++ b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h @@ -11,6 +11,6 @@ #include "clc/internal/clc.h" -_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_max_sub_group_size(void); +_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_max_sub_group_size(void); #endif // __CLC_WORKITEM_CLC_GET_MAX_SUB_GROUP_SIZE_H__ diff --git a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h index b584df98e44a6..3d1da26e8a02a 100644 --- a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h +++ b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h @@ -11,6 +11,6 @@ #include "clc/internal/clc.h" -_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_num_sub_groups(); +_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_num_sub_groups(void); #endif // __CLC_WORKITEM_CLC_GET_NUM_SUB_GROUPS_H__ diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h index 44a4459aa48b4..b21e3c1d8df05 100644 --- a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h +++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h @@ -11,6 +11,6 @@ #include "clc/internal/clc.h" -_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_sub_group_id(); +_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_id(void); #endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_ID_H__ diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h index 52e4b3f28083a..0cf8890dd46b5 100644 --- a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h +++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h @@ -11,6 +11,6 @@ #include "clc/internal/clc.h" -_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_sub_group_local_id(); +_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_local_id(void); #endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_LOCAL_ID_H__ diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h index 4603bfdcbeb25..1dd857d16a2bf 100644 --- a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h +++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h @@ -11,6 +11,6 @@ #include "clc/internal/clc.h" -_CLC_OVERLOAD _CLC_CONST _CLC_DECL uint __clc_get_sub_group_size(); +_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_size(void); #endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_SIZE_H__ diff --git a/libclc/clc/lib/amdgpu/CMakeLists.txt b/libclc/clc/lib/amdgpu/CMakeLists.txt index a5cd47fab4462..69af2ebe525ad 100644 --- a/libclc/clc/lib/amdgpu/CMakeLists.txt +++ b/libclc/clc/lib/amdgpu/CMakeLists.txt @@ -35,6 +35,7 @@ libclc_configure_source_list(CLC_AMDGPU_SOURCES synchronization/clc_sub_group_barrier.cl synchronization/clc_work_group_barrier.cl workitem/clc_get_enqueued_local_size.cl + workitem/clc_get_enqueued_num_sub_groups.cl workitem/clc_get_global_offset.cl workitem/clc_get_global_size.cl workitem/clc_get_group_id.cl @@ -44,6 +45,7 @@ libclc_configure_source_list(CLC_AMDGPU_SOURCES workitem/clc_get_num_groups.cl workitem/clc_get_num_sub_groups.cl workitem/clc_get_sub_group_id.cl + workitem/clc_get_sub_group_local_id.cl workitem/clc_get_sub_group_size.cl workitem/clc_get_work_dim.cl) diff --git a/libclc/clc/lib/amdgpu/subgroup/clc_sub_group_scan.cl b/libclc/clc/lib/amdgpu/subgroup/clc_sub_group_scan.cl index 3ef735aac2aae..573866dee1fa0 100644 --- a/libclc/clc/lib/amdgpu/subgroup/clc_sub_group_scan.cl +++ b/libclc/clc/lib/amdgpu/subgroup/clc_sub_group_scan.cl @@ -14,6 +14,7 @@ #include "clc/subgroup/clc_sub_group_broadcast.h" #include "clc/subgroup/clc_sub_group_scan.h" #include "clc/subgroup/clc_subgroup.h" +#include "clc/workitem/clc_get_sub_group_local_id.h" #define QUAD_PERM (1 << 15) diff --git a/libclc/clc/lib/amdgpu/subgroup/clc_subgroup.cl b/libclc/clc/lib/amdgpu/subgroup/clc_subgroup.cl index 71f4abc42e895..eda7ca2aff394 100644 --- a/libclc/clc/lib/amdgpu/subgroup/clc_subgroup.cl +++ b/libclc/clc/lib/amdgpu/subgroup/clc_subgroup.cl @@ -9,16 +9,6 @@ #include "clc/amdgpu/amdgpu_utils.h" #include "clc/subgroup/clc_subgroup.h" -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_enqueued_num_sub_groups(void) { - return (__clc_amdgpu_enqueued_workgroup_size() + - __builtin_amdgcn_wavefrontsize() - 1) >> - __clc_amdgpu_wavesize_log2(); -} - -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_local_id(void) { - return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u)); -} - _CLC_DEF _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_all(int x) { return __builtin_amdgcn_ballot_w64(x) == __builtin_amdgcn_read_exec(); } diff --git a/libclc/clc/lib/amdgpu/workitem/clc_get_enqueued_num_sub_groups.cl b/libclc/clc/lib/amdgpu/workitem/clc_get_enqueued_num_sub_groups.cl new file mode 100644 index 0000000000000..bb702da96f0a1 --- /dev/null +++ b/libclc/clc/lib/amdgpu/workitem/clc_get_enqueued_num_sub_groups.cl @@ -0,0 +1,16 @@ +//===----------------------------------------------------------------------===// +// +// 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/amdgpu/amdgpu_utils.h" +#include "clc/workitem/clc_get_sub_group_local_id.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_enqueued_num_sub_groups(void) { + return (__clc_amdgpu_enqueued_workgroup_size() + + __builtin_amdgcn_wavefrontsize() - 1) >> + __clc_amdgpu_wavesize_log2(); +} diff --git a/libclc/clc/lib/amdgpu/workitem/clc_get_num_sub_groups.cl b/libclc/clc/lib/amdgpu/workitem/clc_get_num_sub_groups.cl index cb71ef282466b..5dcd3a57b4a4c 100644 --- a/libclc/clc/lib/amdgpu/workitem/clc_get_num_sub_groups.cl +++ b/libclc/clc/lib/amdgpu/workitem/clc_get_num_sub_groups.cl @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// #include "clc/amdgpu/amdgpu_utils.h" -#include "clc/subgroup/clc_subgroup.h" +#include "clc/workitem/clc_get_num_sub_groups.h" _CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_num_sub_groups(void) { uint group_size = __clc_amdgpu_workgroup_size(); diff --git a/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_local_id.cl b/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_local_id.cl new file mode 100644 index 0000000000000..2493cca0c365c --- /dev/null +++ b/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_local_id.cl @@ -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_sub_group_local_id.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_local_id(void) { + return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u)); +} diff --git a/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_size.cl b/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_size.cl index 77c9f8e91d8ee..7ee264f94b0d0 100644 --- a/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_size.cl +++ b/libclc/clc/lib/amdgpu/workitem/clc_get_sub_group_size.cl @@ -9,6 +9,7 @@ #include "clc/amdgpu/amdgpu_utils.h" #include "clc/shared/clc_min.h" #include "clc/workitem/clc_get_local_linear_id.h" +#include "clc/workitem/clc_get_sub_group_size.h" _CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_size(void) { uint wavesize = __builtin_amdgcn_wavefrontsize(); diff --git a/libclc/clc/lib/generic/collective/clc_work_group_any_all.cl b/libclc/clc/lib/generic/collective/clc_work_group_any_all.cl index 4c79ef1f73eba..33fe5f7bd7ddb 100644 --- a/libclc/clc/lib/generic/collective/clc_work_group_any_all.cl +++ b/libclc/clc/lib/generic/collective/clc_work_group_any_all.cl @@ -13,6 +13,9 @@ #include "clc/collective/clc_work_group_any_all.h" #include "clc/subgroup/clc_subgroup.h" #include "clc/synchronization/clc_work_group_barrier.h" +#include "clc/workitem/clc_get_num_sub_groups.h" +#include "clc/workitem/clc_get_sub_group_id.h" +#include "clc/workitem/clc_get_sub_group_local_id.h" #pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : enable diff --git a/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl index ebf2d2eb1710f..cdecc39725647 100644 --- a/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl +++ b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl @@ -13,6 +13,7 @@ #include "clc/subgroup/clc_subgroup.h" #include "clc/synchronization/clc_work_group_barrier.h" #include "clc/workitem/clc_get_local_id.h" +#include "clc/workitem/clc_get_num_sub_groups.h" #pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : enable diff --git a/libclc/clc/lib/generic/collective/clc_work_group_scan.cl b/libclc/clc/lib/generic/collective/clc_work_group_scan.cl index ae333cd9b8cdf..a4d377c0be964 100644 --- a/libclc/clc/lib/generic/collective/clc_work_group_scan.cl +++ b/libclc/clc/lib/generic/collective/clc_work_group_scan.cl @@ -22,6 +22,7 @@ #include "clc/workitem/clc_get_num_sub_groups.h" #include "clc/workitem/clc_get_sub_group_id.h" #include "clc/workitem/clc_get_sub_group_local_id.h" +#include "clc/workitem/clc_get_sub_group_size.h" #pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : enable diff --git a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl index 02391c52ca813..67b008c312f29 100644 --- a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl +++ b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl @@ -8,6 +8,7 @@ #include "clc/workitem/clc_get_local_linear_id.h" #include "clc/workitem/clc_get_max_sub_group_size.h" +#include "clc/workitem/clc_get_sub_group_id.h" _CLC_OVERLOAD _CLC_DEF uint __clc_get_sub_group_id(void) { return __clc_get_local_linear_id() / __clc_get_max_sub_group_size(); diff --git a/libclc/opencl/lib/generic/CMakeLists.txt b/libclc/opencl/lib/generic/CMakeLists.txt index 4ad60248139ae..1d0d7ddd705e8 100644 --- a/libclc/opencl/lib/generic/CMakeLists.txt +++ b/libclc/opencl/lib/generic/CMakeLists.txt @@ -215,6 +215,7 @@ libclc_configure_source_list(OPENCL_GENERIC_SOURCES synchronization/sub_group_barrier.cl synchronization/work_group_barrier.cl workitem/get_enqueued_local_size.cl + workitem/get_enqueued_num_sub_groups.cl workitem/get_global_id.cl workitem/get_global_linear_id.cl workitem/get_global_offset.cl @@ -223,7 +224,12 @@ libclc_configure_source_list(OPENCL_GENERIC_SOURCES workitem/get_local_id.cl workitem/get_local_linear_id.cl workitem/get_local_size.cl + workitem/get_max_sub_group_size.cl workitem/get_num_groups.cl + workitem/get_num_sub_groups.cl + workitem/get_sub_group_id.cl + workitem/get_sub_group_local_id.cl + workitem/get_sub_group_size.cl workitem/get_work_dim.cl ) diff --git a/libclc/opencl/lib/generic/subgroup/subgroup.cl b/libclc/opencl/lib/generic/subgroup/subgroup.cl index fd552ada4afaf..dfe9867fd0801 100644 --- a/libclc/opencl/lib/generic/subgroup/subgroup.cl +++ b/libclc/opencl/lib/generic/subgroup/subgroup.cl @@ -8,30 +8,6 @@ #include "clc/subgroup/clc_subgroup.h" -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_size(void) { - return __clc_get_sub_group_size(); -} - -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_max_sub_group_size(void) { - return __clc_get_max_sub_group_size(); -} - -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_num_sub_groups(void) { - return __clc_get_num_sub_groups(); -} - -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_enqueued_num_sub_groups(void) { - return __clc_get_enqueued_num_sub_groups(); -} - -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_id(void) { - return __clc_get_sub_group_id(); -} - -_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_local_id(void) { - return __clc_get_sub_group_local_id(); -} - _CLC_DEF _CLC_OVERLOAD _CLC_CONST int sub_group_all(int x) { return __clc_sub_group_all(x); } diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl b/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl new file mode 100644 index 0000000000000..fee3a588c2bbf --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl @@ -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_enqueued_num_sub_groups.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_enqueued_num_sub_groups(void) { + return __clc_get_enqueued_num_sub_groups(); +} diff --git a/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl b/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl new file mode 100644 index 0000000000000..bbd19a88a0165 --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl @@ -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_DEF _CLC_OVERLOAD _CLC_CONST uint get_max_sub_group_size(void) { + return __clc_get_max_sub_group_size(); +} diff --git a/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl b/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl new file mode 100644 index 0000000000000..77163234fe54d --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl @@ -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_num_sub_groups.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_num_sub_groups(void) { + return __clc_get_num_sub_groups(); +} diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl new file mode 100644 index 0000000000000..a1ad6adb4e2cb --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl @@ -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_sub_group_id.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_id(void) { + return __clc_get_sub_group_id(); +} diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl new file mode 100644 index 0000000000000..33164282165b3 --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl @@ -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_sub_group_local_id.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_local_id(void) { + return __clc_get_sub_group_local_id(); +} diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl new file mode 100644 index 0000000000000..62f3382b6d7df --- /dev/null +++ b/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl @@ -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_sub_group_size.h" + +_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_size(void) { + return __clc_get_sub_group_size(); +} `````````` </details> https://github.com/llvm/llvm-project/pull/189328 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
