[libclc] libclc: Remove target definitions of opencl workitem functions (PR #185206)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Sat Mar 7 09:25:57 PST 2026
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185206
These were just calling the __clc implementation, so move it
to generic.
>From 9dd021d271fd7727a7658e7bcfe6b34f21f7ca40 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sat, 7 Mar 2026 17:34:36 +0100
Subject: [PATCH] libclc: Remove target definitions of opencl workitem
functions
These were just calling the __clc implementation, so move it
to generic.
---
libclc/opencl/lib/amdgcn/SOURCES | 4 ----
libclc/opencl/lib/generic/SOURCES | 4 ++++
.../workitem/get_global_offset.cl | 0
.../{amdgcn => generic}/workitem/get_group_id.cl | 0
.../{amdgcn => generic}/workitem/get_local_id.cl | 0
.../{amdgcn => generic}/workitem/get_work_dim.cl | 0
libclc/opencl/lib/ptx-nvidiacl/SOURCES | 9 ---------
.../lib/ptx-nvidiacl/workitem/get_global_id.cl | 13 -------------
.../lib/ptx-nvidiacl/workitem/get_group_id.cl | 13 -------------
.../lib/ptx-nvidiacl/workitem/get_local_id.cl | 13 -------------
.../ptx-nvidiacl/workitem/get_local_linear_id.cl | 13 -------------
.../ptx-nvidiacl/workitem/get_max_sub_group_size.cl | 13 -------------
.../lib/ptx-nvidiacl/workitem/get_num_sub_groups.cl | 13 -------------
.../lib/ptx-nvidiacl/workitem/get_sub_group_id.cl | 13 -------------
.../ptx-nvidiacl/workitem/get_sub_group_local_id.cl | 13 -------------
.../lib/ptx-nvidiacl/workitem/get_sub_group_size.cl | 13 -------------
16 files changed, 4 insertions(+), 130 deletions(-)
rename libclc/opencl/lib/{amdgcn => generic}/workitem/get_global_offset.cl (100%)
rename libclc/opencl/lib/{amdgcn => generic}/workitem/get_group_id.cl (100%)
rename libclc/opencl/lib/{amdgcn => generic}/workitem/get_local_id.cl (100%)
rename libclc/opencl/lib/{amdgcn => generic}/workitem/get_work_dim.cl (100%)
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_linear_id.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_max_sub_group_size.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_sub_groups.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_id.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.cl
delete mode 100644 libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_size.cl
diff --git a/libclc/opencl/lib/amdgcn/SOURCES b/libclc/opencl/lib/amdgcn/SOURCES
index e52f54789bfab..ad2b4ca6589ed 100644
--- a/libclc/opencl/lib/amdgcn/SOURCES
+++ b/libclc/opencl/lib/amdgcn/SOURCES
@@ -1,7 +1,3 @@
mem_fence/fence.cl
subgroup/subgroup.cl
synchronization/sub_group_barrier.cl
-workitem/get_global_offset.cl
-workitem/get_group_id.cl
-workitem/get_local_id.cl
-workitem/get_work_dim.cl
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 18dd6fd3a10c2..398e326ffe482 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -206,6 +206,10 @@ workitem/get_enqueued_local_size.cl
workitem/get_global_id.cl
workitem/get_global_linear_id.cl
workitem/get_global_size.cl
+workitem/get_local_id.cl
workitem/get_local_linear_id.cl
workitem/get_local_size.cl
workitem/get_num_groups.cl
+workitem/get_global_offset.cl
+workitem/get_group_id.cl
+workitem/get_work_dim.cl
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl b/libclc/opencl/lib/generic/workitem/get_global_offset.cl
similarity index 100%
rename from libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl
rename to libclc/opencl/lib/generic/workitem/get_global_offset.cl
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl b/libclc/opencl/lib/generic/workitem/get_group_id.cl
similarity index 100%
rename from libclc/opencl/lib/amdgcn/workitem/get_group_id.cl
rename to libclc/opencl/lib/generic/workitem/get_group_id.cl
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl b/libclc/opencl/lib/generic/workitem/get_local_id.cl
similarity index 100%
rename from libclc/opencl/lib/amdgcn/workitem/get_local_id.cl
rename to libclc/opencl/lib/generic/workitem/get_local_id.cl
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl b/libclc/opencl/lib/generic/workitem/get_work_dim.cl
similarity index 100%
rename from libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl
rename to libclc/opencl/lib/generic/workitem/get_work_dim.cl
diff --git a/libclc/opencl/lib/ptx-nvidiacl/SOURCES b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
index 3ece564c9760e..f20917346a3bc 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
@@ -1,10 +1 @@
mem_fence/fence.cl
-workitem/get_global_id.cl
-workitem/get_group_id.cl
-workitem/get_local_id.cl
-workitem/get_local_linear_id.cl
-workitem/get_max_sub_group_size.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
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl
deleted file mode 100644
index c84991ade76fc..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.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_id.h>
-
-_CLC_OVERLOAD _CLC_DEF size_t get_global_id(uint dim) {
- return __clc_get_global_id(dim);
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl
deleted file mode 100644
index 6fc9832058736..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.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_group_id.h>
-
-_CLC_OVERLOAD _CLC_DEF size_t get_group_id(uint dim) {
- return __clc_get_group_id(dim);
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl
deleted file mode 100644
index a2b608baf7aeb..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.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_local_id.h>
-
-_CLC_OVERLOAD _CLC_DEF size_t get_local_id(uint dim) {
- return __clc_get_local_id(dim);
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_linear_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_linear_id.cl
deleted file mode 100644
index ccb301960726a..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_linear_id.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_local_linear_id.h>
-
-_CLC_OVERLOAD _CLC_DEF size_t get_local_linear_id() {
- return __clc_get_local_linear_id();
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_max_sub_group_size.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_max_sub_group_size.cl
deleted file mode 100644
index aa06c9561404b..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_max_sub_group_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_max_sub_group_size.h>
-
-_CLC_OVERLOAD _CLC_DEF uint get_max_sub_group_size() {
- return __clc_get_max_sub_group_size();
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_sub_groups.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_sub_groups.cl
deleted file mode 100644
index e23d78dafa5e8..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_sub_groups.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_num_sub_groups.h>
-
-_CLC_OVERLOAD _CLC_DEF uint get_num_sub_groups() {
- return __clc_get_num_sub_groups();
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_id.cl
deleted file mode 100644
index deb6134cd0704..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_id.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_sub_group_id.h>
-
-_CLC_OVERLOAD _CLC_DEF uint get_sub_group_id() {
- return __clc_get_sub_group_id();
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.cl
deleted file mode 100644
index 37487fbf3ece8..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.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_sub_group_local_id.h>
-
-_CLC_OVERLOAD _CLC_DEF uint get_sub_group_local_id() {
- return __clc_get_sub_group_local_id();
-}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_size.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_size.cl
deleted file mode 100644
index e6be6f75171d1..0000000000000
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_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_sub_group_size.h>
-
-_CLC_OVERLOAD _CLC_DEF uint get_sub_group_size() {
- return __clc_get_sub_group_size();
-}
More information about the cfe-commits
mailing list