Author: Wenju He Date: 2025-08-19T07:51:17+08:00 New Revision: a450dc80bf3fe8af7caf6f88265cf629fb261169
URL: https://github.com/llvm/llvm-project/commit/a450dc80bf3fe8af7caf6f88265cf629fb261169 DIFF: https://github.com/llvm/llvm-project/commit/a450dc80bf3fe8af7caf6f88265cf629fb261169.diff LOG: [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 Added: libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl Modified: libclc/clc/lib/amdgcn/SOURCES Removed: ################################################################################ diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES index 76c3266e3af7b..53bbe388f7dfc 100644 --- a/libclc/clc/lib/amdgcn/SOURCES +++ b/libclc/clc/lib/amdgcn/SOURCES @@ -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 diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl new file mode 100644 index 0000000000000..1e749404168d8 --- /dev/null +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl @@ -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; + } +} diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl new file mode 100644 index 0000000000000..cc56f8d9c325d --- /dev/null +++ b/libclc/clc/lib/amdgcn/workitem/clc_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_OVERLOAD _CLC_DEF uint __clc_get_max_sub_group_size() { + return __builtin_amdgcn_wavefrontsize(); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
