https://github.com/wenju-he updated https://github.com/llvm/llvm-project/pull/153785
>From d0a9e7fa683d294aaabf24ccc34cea54a8a5eb1f Mon Sep 17 00:00:00 2001 From: Wenju He <wenju...@intel.com> Date: Fri, 15 Aug 2025 12:43:54 +0200 Subject: [PATCH 1/2] [libclc] Implement __clc_get_local_size/__clc_get_max_sub_group_size for amdgcn This simplifies downstream refactoring of libspirv workitem function in https://github.com/intel/llvm/tree/sycl/libclc/libspirv/lib/generic --- libclc/clc/lib/amdgcn/SOURCES | 2 ++ .../lib/amdgcn/workitem/clc_get_local_size.cl | 22 +++++++++++++++++++ .../workitem/clc_get_max_sub_group_size.cl | 13 +++++++++++ 3 files changed, 37 insertions(+) create mode 100644 libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl create mode 100644 libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl 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..6806d64fb7639 --- /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 0; + } +} 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(); +} >From a8fcf405007d05d123f397f66b4139d53586c1b5 Mon Sep 17 00:00:00 2001 From: Wenju He <wenju...@intel.com> Date: Fri, 15 Aug 2025 23:16:53 +0200 Subject: [PATCH 2/2] change out-of-bound get_local_size to 1 --- libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl index 6806d64fb7639..1e749404168d8 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl @@ -17,6 +17,6 @@ _CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) { case 2: return __builtin_amdgcn_workgroup_size_z(); default: - return 0; + return 1; } } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
