[libclc] 10c0e1f - libclc: Avoid duplicated get_local_size/get_global_size functions (#185166)
via cfe-commits
cfe-commits at lists.llvm.org
Sat Mar 7 02:37:04 PST 2026
Author: Matt Arsenault
Date: 2026-03-07T11:37:00+01:00
New Revision: 10c0e1fc86bcb2a55d4ef5037b055d5a9808dba4
URL: https://github.com/llvm/llvm-project/commit/10c0e1fc86bcb2a55d4ef5037b055d5a9808dba4
DIFF: https://github.com/llvm/llvm-project/commit/10c0e1fc86bcb2a55d4ef5037b055d5a9808dba4.diff
LOG: libclc: Avoid duplicated get_local_size/get_global_size functions (#185166)
Move opencl handling on top of clc into opencl generic, delete
amdgpu implementations in opencl.
Added:
libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
libclc/opencl/lib/generic/workitem/get_local_size.cl
Modified:
libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl
libclc/clc/lib/ptx-nvidiacl/SOURCES
libclc/opencl/lib/amdgcn/SOURCES
libclc/opencl/lib/generic/SOURCES
libclc/opencl/lib/generic/workitem/get_global_size.cl
libclc/opencl/lib/ptx-nvidiacl/SOURCES
Removed:
libclc/opencl/lib/amdgcn-amdhsa/SOURCES
libclc/opencl/lib/amdgcn-amdhsa/workitem/get_global_size.cl
libclc/opencl/lib/amdgcn-amdhsa/workitem/get_local_size.cl
libclc/opencl/lib/amdgcn/workitem/get_global_size.cl
libclc/opencl/lib/amdgcn/workitem/get_local_size.cl
libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl
################################################################################
diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
index b1d8f27dc68c8..1886ab6d1a1a0 100644
--- a/libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
@@ -6,17 +6,15 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/workitem/clc_get_global_size.h>
+#include "clc/workitem/clc_get_global_size.h"
+#include <amdhsa_abi.h>
_CLC_DEF _CLC_OVERLOAD size_t __clc_get_global_size(uint dim) {
- switch (dim) {
- case 0:
- return __builtin_amdgcn_grid_size_x();
- case 1:
- return __builtin_amdgcn_grid_size_y();
- case 2:
- return __builtin_amdgcn_grid_size_z();
- default:
+ 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/clc/lib/amdgcn/workitem/clc_get_local_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl
index 1e749404168d8..57d801a2d02d4 100644
--- a/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_local_size.cl
@@ -6,17 +6,21 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/workitem/clc_get_local_size.h>
+#include "clc/workitem/clc_get_local_size.h"
+#include <amdhsa_abi.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:
+ 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/clc/lib/ptx-nvidiacl/SOURCES b/libclc/clc/lib/ptx-nvidiacl/SOURCES
index cafd90943f22e..9ed25c71a3f35 100644
--- a/libclc/clc/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/clc/lib/ptx-nvidiacl/SOURCES
@@ -5,6 +5,7 @@ math/clc_sqrt.cl
mem_fence/clc_mem_fence.cl
synchronization/clc_work_group_barrier.cl
workitem/clc_get_global_id.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
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
similarity index 58%
rename from libclc/opencl/lib/amdgcn/workitem/get_global_size.cl
rename to libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
index eca7199a766fc..262fd3b1b43a3 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
@@ -6,8 +6,10 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/workitem/clc_get_global_size.h>
+#include "clc/workitem/clc_get_global_size.h"
+#include "clc/workitem/clc_get_local_size.h"
+#include "clc/workitem/clc_get_num_groups.h"
-_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
- return __clc_get_global_size(dim);
+_CLC_DEF _CLC_OVERLOAD size_t __clc_get_global_size(uint dim) {
+ return __clc_get_num_groups(dim) * __clc_get_local_size(dim);
}
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_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 cdc1d8321dfeb..18dd6fd3a10c2 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -207,4 +207,5 @@ workitem/get_global_id.cl
workitem/get_global_linear_id.cl
workitem/get_global_size.cl
workitem/get_local_linear_id.cl
+workitem/get_local_size.cl
workitem/get_num_groups.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_global_size.cl b/libclc/opencl/lib/generic/workitem/get_global_size.cl
index 34d00f8fa809b..2fe343ca48c41 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_size.cl
@@ -6,8 +6,8 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/opencl/opencl-base.h>
+#include "clc/workitem/clc_get_global_size.h"
_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
- return get_num_groups(dim) * get_local_size(dim);
+ return __clc_get_global_size(dim);
}
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
More information about the cfe-commits
mailing list