[libclc] libclc: Implement get_global_linear_id and get_local_linear_id (PR #184800)

[Matt Arsenault via cfe-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/184800

None

>From 4013343b2c623d74af9062face97422860a7cc56 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] 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 bb5e8ab08a711..85eb5ba40213f 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -200,4 +200,6 @@ shared/min.cl
 shared/vload.cl
 shared/vstore.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);
+}