[libclc] libclc: Implement get_global_linear_id and get_local_linear_id (PR #184800)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 6 06:52:37 PST 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/184800
>From cf0cda96f2b5ec64ee189b2ef48b1d80efd7569c Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 5 Mar 2026 13:26:35 +0100
Subject: [PATCH 1/2] libclc: Implement get_global_linear_id and
get_local_linear_id
---
libclc/opencl/lib/generic/SOURCES | 2 +
.../generic/workitem/get_global_linear_id.cl | 40 +++++++++++++++++++
.../generic/workitem/get_local_linear_id.cl | 15 +++++++
3 files changed, 57 insertions(+)
create mode 100644 libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
create mode 100644 libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 312657f3bf106..e4de6633d6030 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -204,4 +204,6 @@ subgroup/sub_group_broadcast.cl
synchronization/work_group_barrier.cl
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_linear_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
new file mode 100644
index 0000000000000..557c5ab3516ff
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
@@ -0,0 +1,40 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+static size_t get_global_id_no_offset(uint dim) {
+ return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
+}
+
+static size_t get_global_linear_id_1d() { return get_global_id_no_offset(0); }
+
+static size_t get_global_linear_id_2d() {
+ return get_global_id_no_offset(1) * get_global_size(0) +
+ get_global_linear_id_1d();
+}
+
+static size_t get_global_linear_id_3d() {
+ return (get_global_id_no_offset(2) * get_global_size(1) +
+ get_global_id_no_offset(1)) *
+ get_global_size(0) +
+ get_global_id_no_offset(0);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_global_linear_id() {
+ switch (get_work_dim()) {
+ case 1:
+ return get_global_linear_id_1d();
+ case 2:
+ return get_global_linear_id_2d();
+ case 3:
+ return get_global_linear_id_3d();
+ default:
+ return 0;
+ }
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
new file mode 100644
index 0000000000000..10daa3d7b6bd2
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
@@ -0,0 +1,15 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_OVERLOAD _CLC_DEF _CLC_CONST size_t get_local_linear_id() {
+ return (get_local_id(2) * get_local_size(1) + get_local_id(1)) *
+ get_local_size(0) +
+ get_local_id(0);
+}
>From 21eca3b4d48e97a15ff606251ede90a98a44cd0d Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Fri, 6 Mar 2026 15:44:49 +0100
Subject: [PATCH 2/2] Move to clc
---
.../clc/workitem/clc_get_global_linear_id.h | 16 +++++++
libclc/clc/lib/generic/SOURCES | 1 +
.../workitem/clc_get_global_linear_id.cl | 46 +++++++++++++++++++
.../workitem/clc_get_local_linear_id.cl | 12 ++---
.../generic/workitem/get_global_linear_id.cl | 31 +------------
.../generic/workitem/get_local_linear_id.cl | 6 +--
6 files changed, 73 insertions(+), 39 deletions(-)
create mode 100644 libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
create mode 100644 libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h b/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
new file mode 100644
index 0000000000000..fab5040e497ad
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__
+#define __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__
+
+#include "clc/internal/clc.h"
+
+_CLC_OVERLOAD _CLC_CONST _CLC_DECL size_t __clc_get_global_linear_id();
+
+#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index 421eb638720e3..b352382fbe6d9 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -176,6 +176,7 @@ shared/clc_min.cl
shared/clc_qualifier.cl
shared/clc_vload.cl
shared/clc_vstore.cl
+workitem/clc_get_global_linear_id.cl
workitem/clc_get_local_linear_id.cl
workitem/clc_get_num_sub_groups.cl
workitem/clc_get_sub_group_id.cl
diff --git a/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl b/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
new file mode 100644
index 0000000000000..ffd7e918a4084
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
@@ -0,0 +1,46 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_linear_id.h"
+#include "clc/workitem/clc_get_global_size.h"
+#include "clc/workitem/clc_get_group_id.h"
+#include "clc/workitem/clc_get_local_id.h"
+#include "clc/workitem/clc_get_local_size.h"
+#include "clc/workitem/clc_get_work_dim.h"
+
+static size_t get_global_id_no_offset(uint dim) {
+ return __clc_get_group_id(dim) * __clc_get_local_size(dim) +
+ __clc_get_local_id(dim);
+}
+
+static size_t get_global_linear_id_1d() { return get_global_id_no_offset(0); }
+
+static size_t get_global_linear_id_2d() {
+ return get_global_id_no_offset(1) * __clc_get_global_size(0) +
+ get_global_linear_id_1d();
+}
+
+static size_t get_global_linear_id_3d() {
+ return (get_global_id_no_offset(2) * __clc_get_global_size(1) +
+ get_global_id_no_offset(1)) *
+ __clc_get_global_size(0) +
+ get_global_id_no_offset(0);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t __clc_get_global_linear_id() {
+ switch (__clc_get_work_dim()) {
+ case 1:
+ return get_global_linear_id_1d();
+ case 2:
+ return get_global_linear_id_2d();
+ case 3:
+ return get_global_linear_id_3d();
+ default:
+ return 0;
+ }
+}
diff --git a/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl b/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl
index fd7905568d595..a8fcb0059f768 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl
@@ -6,13 +6,13 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/workitem/clc_get_local_id.h>
-#include <clc/workitem/clc_get_local_linear_id.h>
-#include <clc/workitem/clc_get_local_size.h>
+#include "clc/workitem/clc_get_local_id.h"
+#include "clc/workitem/clc_get_local_linear_id.h"
+#include "clc/workitem/clc_get_local_size.h"
-_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_linear_id() {
- return __clc_get_local_id(2) * __clc_get_local_size(1) *
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t __clc_get_local_linear_id() {
+ return (__clc_get_local_id(2) * __clc_get_local_size(1) +
+ __clc_get_local_id(1)) *
__clc_get_local_size(0) +
- __clc_get_local_id(1) * __clc_get_local_size(0) +
__clc_get_local_id(0);
}
diff --git a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
index 557c5ab3516ff..2d0f039d3a626 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
@@ -6,35 +6,8 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/opencl/opencl-base.h>
-
-static size_t get_global_id_no_offset(uint dim) {
- return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
-}
-
-static size_t get_global_linear_id_1d() { return get_global_id_no_offset(0); }
-
-static size_t get_global_linear_id_2d() {
- return get_global_id_no_offset(1) * get_global_size(0) +
- get_global_linear_id_1d();
-}
-
-static size_t get_global_linear_id_3d() {
- return (get_global_id_no_offset(2) * get_global_size(1) +
- get_global_id_no_offset(1)) *
- get_global_size(0) +
- get_global_id_no_offset(0);
-}
+#include "clc/workitem/clc_get_global_linear_id.h"
_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_global_linear_id() {
- switch (get_work_dim()) {
- case 1:
- return get_global_linear_id_1d();
- case 2:
- return get_global_linear_id_2d();
- case 3:
- return get_global_linear_id_3d();
- default:
- return 0;
- }
+ return __clc_get_global_linear_id();
}
diff --git a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
index 10daa3d7b6bd2..e7599fa8c75b2 100644
--- a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
@@ -6,10 +6,8 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/opencl/opencl-base.h>
+#include "clc/workitem/clc_get_local_linear_id.h"
_CLC_OVERLOAD _CLC_DEF _CLC_CONST size_t get_local_linear_id() {
- return (get_local_id(2) * get_local_size(1) + get_local_id(1)) *
- get_local_size(0) +
- get_local_id(0);
+ return __clc_get_local_linear_id();
}
More information about the cfe-commits
mailing list