https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185166
Move opencl handling on top of clc into opencl generic, delete amdgpu implementations in opencl. >From 13e75dec4a6c320f4b8dd8faed6be05f209cb708 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Fri, 6 Mar 2026 17:55:54 +0100 Subject: [PATCH] libclc: Avoid duplicated get_local_size/get_global_size functions Move opencl handling on top of clc into opencl generic, delete amdgpu implementations in opencl. --- libclc/opencl/lib/amdgcn-amdhsa/SOURCES | 2 -- .../amdgcn-amdhsa/workitem/get_global_size.cl | 20 -------------- .../amdgcn-amdhsa/workitem/get_local_size.cl | 26 ------------------- libclc/opencl/lib/amdgcn/SOURCES | 2 -- .../lib/amdgcn/workitem/get_global_size.cl | 13 ---------- .../lib/amdgcn/workitem/get_local_size.cl | 22 ---------------- libclc/opencl/lib/generic/SOURCES | 1 + .../workitem/get_local_size.cl | 0 libclc/opencl/lib/ptx-nvidiacl/SOURCES | 1 - 9 files changed, 1 insertion(+), 86 deletions(-) delete mode 100644 libclc/opencl/lib/amdgcn-amdhsa/SOURCES delete mode 100644 libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl delete mode 100644 libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl delete mode 100644 libclc/opencl/lib/amdgcn/workitem/get_global_size.cl delete mode 100644 libclc/opencl/lib/amdgcn/workitem/get_local_size.cl rename libclc/opencl/lib/{ptx-nvidiacl => generic}/workitem/get_local_size.cl (100%) diff --git a/libclc/opencl/lib/amdgcn-amdhsa/SOURCES b/libclc/opencl/lib/amdgcn-amdhsa/SOURCES deleted file mode 100644 index ee3a48ce2c474..0000000000000 --- a/libclc/opencl/lib/amdgcn-amdhsa/SOURCES +++ /dev/null @@ -1,2 +0,0 @@ -workitem/get_global_size.cl -workitem/get_local_size.cl diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl deleted file mode 100644 index f21a060849dbe..0000000000000 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl +++ /dev/null @@ -1,20 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 <amdhsa_abi.h> -#include <clc/opencl/opencl-base.h> - -_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { - if (dim > 2) - return 1; - __constant amdhsa_implicit_kernarg_v5 *args = - (__constant amdhsa_implicit_kernarg_v5 *) - __builtin_amdgcn_implicitarg_ptr(); - return args->block_count[dim] * (uint)args->group_size[dim] + - (uint)args->remainder[dim]; -} diff --git a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl b/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl deleted file mode 100644 index ed1e17776361e..0000000000000 --- a/libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl +++ /dev/null @@ -1,26 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 <amdhsa_abi.h> -#include <clc/opencl/opencl-base.h> - -_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) { - if (dim > 2) - return 1; - - __constant amdhsa_implicit_kernarg_v5 *args = - (__constant amdhsa_implicit_kernarg_v5 *) - __builtin_amdgcn_implicitarg_ptr(); - - uint group_ids[3] = {__builtin_amdgcn_workgroup_id_x(), - __builtin_amdgcn_workgroup_id_y(), - __builtin_amdgcn_workgroup_id_z()}; - - return group_ids[dim] < args->block_count[dim] ? (size_t)args->group_size[dim] - : (size_t)args->remainder[dim]; -} diff --git a/libclc/opencl/lib/amdgcn/SOURCES b/libclc/opencl/lib/amdgcn/SOURCES index ac72d8a00c9d0..e52f54789bfab 100644 --- a/libclc/opencl/lib/amdgcn/SOURCES +++ b/libclc/opencl/lib/amdgcn/SOURCES @@ -3,7 +3,5 @@ subgroup/subgroup.cl synchronization/sub_group_barrier.cl workitem/get_global_offset.cl workitem/get_group_id.cl -workitem/get_global_size.cl workitem/get_local_id.cl -workitem/get_local_size.cl workitem/get_work_dim.cl diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl b/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl deleted file mode 100644 index eca7199a766fc..0000000000000 --- a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl +++ /dev/null @@ -1,13 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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_global_size.h> - -_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) { - return __clc_get_global_size(dim); -} diff --git a/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl b/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl deleted file mode 100644 index 34e4f2f1b4c19..0000000000000 --- a/libclc/opencl/lib/amdgcn/workitem/get_local_size.cl +++ /dev/null @@ -1,22 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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/opencl/opencl-base.h> - -_CLC_DEF _CLC_OVERLOAD size_t 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/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES index dd51d5c084d51..5add00020d6ca 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -205,4 +205,5 @@ synchronization/work_group_barrier.cl workitem/get_enqueued_local_size.cl workitem/get_global_id.cl workitem/get_global_size.cl +workitem/get_local_size.cl workitem/get_num_groups.cl diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl b/libclc/opencl/lib/generic/workitem/get_local_size.cl similarity index 100% rename from libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl rename to libclc/opencl/lib/generic/workitem/get_local_size.cl diff --git a/libclc/opencl/lib/ptx-nvidiacl/SOURCES b/libclc/opencl/lib/ptx-nvidiacl/SOURCES index b8e8f64b5802a..3ece564c9760e 100644 --- a/libclc/opencl/lib/ptx-nvidiacl/SOURCES +++ b/libclc/opencl/lib/ptx-nvidiacl/SOURCES @@ -3,7 +3,6 @@ workitem/get_global_id.cl workitem/get_group_id.cl 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_sub_groups.cl workitem/get_sub_group_id.cl _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
