[libclc] [libclc] Declare workitem built-ins in clc, move ptx-nvidiacl workitem built-ins into clc (PR #144333)

Wenju He via cfe-commits cfe-commits at lists.llvm.org
Wed Jul 2 20:02:35 PDT 2025


https://github.com/wenju-he updated https://github.com/llvm/llvm-project/pull/144333

>From 4ef7a1c91eb33e72cc6c91725d59d6e5b6791ab9 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Mon, 16 Jun 2025 13:22:29 +0200
Subject: [PATCH 1/4] [libclc] Declare workitem built-ins in clc, move
 ptx-nvidiacl workitem built-ins into clc

Changes in this PR:
* Declare all workitem functions in clc and opencl folders.
* Call clc workitem function in corresponding OpenCL workitem function.
* Move ptx-nvidiacl workitem built-in implementations into clc.
* clc_get_global_offset, clc_get_enqueued_num_sub_groups,
  clc_get_enqueued_local_size and clc_get_work_dim are not implemented
  because it is not possible to provide generic implementation.
* Include only needed headers in OpenCL workitem functions.
* Add a FIXME to OpenCL get_global_id.cl.

Note opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl isn't changed
because it has different implementation from generic.
llvm-diff shows this PR only adds a dozen new symbols to
amdgcn--amdhsa.bc and nvptx64--nvidiacl.bc.
---
 .../workitem/clc_get_enqueued_local_size.h    | 16 ++++++++++
 .../clc_get_enqueued_num_sub_groups.h         | 16 ++++++++++
 .../include/clc/workitem/clc_get_global_id.h  | 16 ++++++++++
 .../clc/workitem/clc_get_global_linear_id.h   | 16 ++++++++++
 .../clc/workitem/clc_get_global_offset.h      | 16 ++++++++++
 .../clc/workitem/clc_get_global_size.h        | 16 ++++++++++
 .../include/clc/workitem/clc_get_group_id.h   | 16 ++++++++++
 .../include/clc/workitem/clc_get_local_id.h   | 16 ++++++++++
 .../clc/workitem/clc_get_local_linear_id.h    | 16 ++++++++++
 .../include/clc/workitem/clc_get_local_size.h | 16 ++++++++++
 .../clc/workitem/clc_get_max_sub_group_size.h | 16 ++++++++++
 .../include/clc/workitem/clc_get_num_groups.h | 16 ++++++++++
 .../clc/workitem/clc_get_num_sub_groups.h     | 16 ++++++++++
 .../clc/workitem/clc_get_sub_group_id.h       | 16 ++++++++++
 .../clc/workitem/clc_get_sub_group_local_id.h | 16 ++++++++++
 .../clc/workitem/clc_get_sub_group_size.h     | 16 ++++++++++
 .../include/clc/workitem/clc_get_work_dim.h   | 16 ++++++++++
 libclc/clc/lib/generic/SOURCES                |  6 ++++
 .../lib/generic/workitem/clc_get_global_id.cl | 18 +++++++++++
 .../workitem/clc_get_global_linear_id.cl      | 32 +++++++++++++++++++
 .../workitem/clc_get_local_linear_id.cl       | 16 ++++++++++
 .../workitem/clc_get_num_sub_groups.cl        | 18 +++++++++++
 .../generic/workitem/clc_get_sub_group_id.cl  | 25 +++++++++++++++
 .../workitem/clc_get_sub_group_size.cl        | 26 +++++++++++++++
 libclc/clc/lib/ptx-nvidiacl/SOURCES           |  7 ++++
 .../workitem/clc_get_global_size.cl           | 15 +++++++++
 .../workitem/clc_get_group_id.cl}             |  4 +--
 .../workitem/clc_get_local_id.cl}             |  4 +--
 .../workitem/clc_get_local_size.cl}           |  4 +--
 .../workitem/clc_get_max_sub_group_size.cl    | 13 ++++++++
 .../workitem/clc_get_num_groups.cl}           |  4 +--
 .../workitem/clc_get_sub_group_local_id.cl    | 13 ++++++++
 libclc/opencl/include/clc/opencl/clc.h        |  9 ++++++
 .../opencl/workitem/get_enqueued_local_size.h | 16 ++++++++++
 .../workitem/get_enqueued_num_sub_groups.h    | 16 ++++++++++
 .../clc/opencl/workitem/get_global_id.h       |  9 +++++-
 .../opencl/workitem/get_global_linear_id.h    | 16 ++++++++++
 .../clc/opencl/workitem/get_global_offset.h   |  9 +++++-
 .../clc/opencl/workitem/get_global_size.h     |  9 +++++-
 .../clc/opencl/workitem/get_group_id.h        |  9 +++++-
 .../clc/opencl/workitem/get_local_id.h        |  9 +++++-
 .../clc/opencl/workitem/get_local_linear_id.h | 16 ++++++++++
 .../clc/opencl/workitem/get_local_size.h      |  9 +++++-
 .../opencl/workitem/get_max_sub_group_size.h  | 16 ++++++++++
 .../clc/opencl/workitem/get_num_groups.h      |  9 +++++-
 .../clc/opencl/workitem/get_num_sub_groups.h  | 16 ++++++++++
 .../clc/opencl/workitem/get_sub_group_id.h    | 16 ++++++++++
 .../opencl/workitem/get_sub_group_local_id.h  | 16 ++++++++++
 .../clc/opencl/workitem/get_sub_group_size.h  | 16 ++++++++++
 .../clc/opencl/workitem/get_work_dim.h        |  9 +++++-
 libclc/opencl/lib/generic/SOURCES             | 15 +++++++++
 .../workitem/get_enqueued_local_size.cl       | 14 ++++++++
 .../workitem/get_enqueued_num_sub_groups.cl   | 14 ++++++++
 .../lib/generic/workitem/get_global_id.cl     |  2 ++
 .../generic/workitem/get_global_linear_id.cl  | 14 ++++++++
 .../lib/generic/workitem/get_global_offset.cl | 14 ++++++++
 .../lib/generic/workitem/get_global_size.cl   |  7 ++--
 .../lib/generic/workitem/get_group_id.cl      | 14 ++++++++
 .../lib/generic/workitem/get_local_id.cl      | 14 ++++++++
 .../generic/workitem/get_local_linear_id.cl   | 14 ++++++++
 .../lib/generic/workitem/get_local_size.cl    | 14 ++++++++
 .../workitem/get_max_sub_group_size.cl        | 14 ++++++++
 .../lib/generic/workitem/get_num_groups.cl    | 14 ++++++++
 .../generic/workitem/get_num_sub_groups.cl    | 14 ++++++++
 .../lib/generic/workitem/get_sub_group_id.cl  | 14 ++++++++
 .../workitem/get_sub_group_local_id.cl        | 14 ++++++++
 .../generic/workitem/get_sub_group_size.cl    | 14 ++++++++
 .../lib/generic/workitem/get_work_dim.cl      | 12 +++++++
 libclc/opencl/lib/ptx-nvidiacl/SOURCES        |  4 ---
 69 files changed, 915 insertions(+), 23 deletions(-)
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_global_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_global_offset.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_global_size.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_group_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_local_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_local_size.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_num_groups.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
 create mode 100644 libclc/clc/include/clc/workitem/clc_get_work_dim.h
 create mode 100644 libclc/clc/lib/generic/workitem/clc_get_global_id.cl
 create mode 100644 libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
 create mode 100644 libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl
 create mode 100644 libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl
 create mode 100644 libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
 create mode 100644 libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl
 create mode 100644 libclc/clc/lib/ptx-nvidiacl/SOURCES
 create mode 100644 libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
 rename libclc/{opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl => clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl} (85%)
 rename libclc/{opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl => clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl} (84%)
 rename libclc/{opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl => clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl} (84%)
 create mode 100644 libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl
 rename libclc/{opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl => clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl} (84%)
 create mode 100644 libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h
 create mode 100644 libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h
 create mode 100644 libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_global_offset.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_group_id.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_local_id.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_local_size.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_num_groups.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
 create mode 100644 libclc/opencl/lib/generic/workitem/get_work_dim.cl

diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
new file mode 100644
index 0000000000000..83aaed72c5036
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.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_ENQUEUED_LOCAL_SIZE_H__
+#define __CLC_WORKITEM_CLC_GET_ENQUEUED_LOCAL_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_enqueued_local_size(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_LOCAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
new file mode 100644
index 0000000000000..2a5af05f3f2d6
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.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_ENQUEUED_NUM_SUB_GROUPS_H__
+#define __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_DEF _CLC_OVERLOAD uint clc_get_enqueued_num_sub_groups();
+
+#endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_id.h b/libclc/clc/include/clc/workitem/clc_get_global_id.h
new file mode 100644
index 0000000000000..697d7a629d794
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_global_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_ID_H__
+#define __CLC_WORKITEM_CLC_GET_GLOBAL_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_id(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_ID_H__
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..ac5a73ecd9e6e
--- /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_DECL size_t clc_get_global_linear_id();
+
+#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_offset.h b/libclc/clc/include/clc/workitem/clc_get_global_offset.h
new file mode 100644
index 0000000000000..ae51d88fea6a4
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_global_offset.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_OFFSET_H__
+#define __CLC_WORKITEM_CLC_GET_GLOBAL_OFFSET_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_offset(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_OFFSET_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_size.h b/libclc/clc/include/clc/workitem/clc_get_global_size.h
new file mode 100644
index 0000000000000..67c315b1b6c38
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_global_size.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_SIZE_H__
+#define __CLC_WORKITEM_CLC_GET_GLOBAL_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_size(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_group_id.h b/libclc/clc/include/clc/workitem/clc_get_group_id.h
new file mode 100644
index 0000000000000..a562fa175bf59
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_group_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_GROUP_ID_H__
+#define __CLC_WORKITEM_CLC_GET_GROUP_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_group_id(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_GROUP_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_id.h b/libclc/clc/include/clc/workitem/clc_get_local_id.h
new file mode 100644
index 0000000000000..482bb890d2159
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_local_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_LOCAL_ID_H__
+#define __CLC_WORKITEM_CLC_GET_LOCAL_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_local_id(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_LOCAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h b/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
new file mode 100644
index 0000000000000..ef27982cb176c
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_local_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_LOCAL_LINEAR_ID_H__
+#define __CLC_WORKITEM_CLC_GET_LOCAL_LINEAR_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_local_linear_id();
+
+#endif // __CLC_WORKITEM_CLC_GET_LOCAL_LINEAR_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_size.h b/libclc/clc/include/clc/workitem/clc_get_local_size.h
new file mode 100644
index 0000000000000..93fa68562f288
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_local_size.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_LOCAL_SIZE_H__
+#define __CLC_WORKITEM_CLC_GET_LOCAL_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_local_size(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_LOCAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
new file mode 100644
index 0000000000000..7593869650be4
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.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_MAX_SUB_GROUP_SIZE_H__
+#define __CLC_WORKITEM_CLC_GET_MAX_SUB_GROUP_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_DEF _CLC_OVERLOAD uint clc_get_max_sub_group_size();
+
+#endif // __CLC_WORKITEM_CLC_GET_MAX_SUB_GROUP_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_num_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_groups.h
new file mode 100644
index 0000000000000..03e0abb48ec33
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_num_groups.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_NUM_GROUPS_H__
+#define __CLC_WORKITEM_CLC_GET_NUM_GROUPS_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t clc_get_num_groups(uint dim);
+
+#endif // __CLC_WORKITEM_CLC_GET_NUM_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
new file mode 100644
index 0000000000000..c17ebea1146c3
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.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_NUM_SUB_GROUPS_H__
+#define __CLC_WORKITEM_CLC_GET_NUM_SUB_GROUPS_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_DEF _CLC_OVERLOAD uint clc_get_num_sub_groups();
+
+#endif // __CLC_WORKITEM_CLC_GET_NUM_SUB_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
new file mode 100644
index 0000000000000..eadf08cb66f6b
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_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_SUB_GROUP_ID_H__
+#define __CLC_WORKITEM_CLC_GET_SUB_GROUP_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_DEF _CLC_OVERLOAD uint clc_get_sub_group_id();
+
+#endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
new file mode 100644
index 0000000000000..d631f37847a8a
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_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_SUB_GROUP_LOCAL_ID_H__
+#define __CLC_WORKITEM_CLC_GET_SUB_GROUP_LOCAL_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_DEF _CLC_OVERLOAD uint clc_get_sub_group_local_id();
+
+#endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_LOCAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
new file mode 100644
index 0000000000000..56b736759a455
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.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_SUB_GROUP_SIZE_H__
+#define __CLC_WORKITEM_CLC_GET_SUB_GROUP_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_DEF _CLC_OVERLOAD uint clc_get_sub_group_size();
+
+#endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_work_dim.h b/libclc/clc/include/clc/workitem/clc_get_work_dim.h
new file mode 100644
index 0000000000000..1df6c4e819902
--- /dev/null
+++ b/libclc/clc/include/clc/workitem/clc_get_work_dim.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_WORK_DIM_H__
+#define __CLC_WORKITEM_CLC_GET_WORK_DIM_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint clc_get_work_dim();
+
+#endif // __CLC_WORKITEM_CLC_GET_WORK_DIM_H__
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index d285bbba3dd26..d840c54d2af20 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -152,3 +152,9 @@ shared/clc_max.cl
 shared/clc_min.cl
 shared/clc_vload.cl
 shared/clc_vstore.cl
+workitem/clc_get_sub_group_size.cl
+workitem/clc_get_global_id.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_id.cl b/libclc/clc/lib/generic/workitem/clc_get_global_id.cl
new file mode 100644
index 0000000000000..bae19352a3e11
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_global_id.cl
@@ -0,0 +1,18 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_enqueued_local_size.h>
+#include <clc/workitem/clc_get_global_id.h>
+#include <clc/workitem/clc_get_global_offset.h>
+#include <clc/workitem/clc_get_group_id.h>
+#include <clc/workitem/clc_get_local_id.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t clc_get_global_id(uint dim) {
+  return clc_get_group_id(dim) * clc_get_enqueued_local_size(dim) +
+         clc_get_local_id(dim) + clc_get_global_offset(dim);
+}
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..850c33f901d86
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
@@ -0,0 +1,32 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+#include <clc/workitem/clc_get_global_linear_id.h>
+#include <clc/workitem/clc_get_global_offset.h>
+#include <clc/workitem/clc_get_global_size.h>
+#include <clc/workitem/clc_get_work_dim.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t clc_get_global_linear_id() {
+  uint dim = clc_get_work_dim();
+  switch (dim) {
+  default:
+  case 1:
+    return clc_get_global_id(0) - clc_get_global_offset(0);
+  case 2:
+    return (clc_get_global_id(1) - clc_get_global_offset(1)) *
+               clc_get_global_size(0) +
+           (clc_get_global_id(0) - clc_get_global_offset(0));
+  case 3:
+    return ((clc_get_global_id(2) - clc_get_global_offset(2)) *
+            clc_get_global_size(1) * clc_get_global_size(0)) +
+           ((clc_get_global_id(1) - clc_get_global_offset(1)) *
+            clc_get_global_size(0)) +
+           (clc_get_global_id(0) - clc_get_global_offset(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
new file mode 100644
index 0000000000000..0ef7a238af606
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_local_linear_id.cl
@@ -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
+//
+//===----------------------------------------------------------------------===//
+
+#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_get_local_size(0) +
+         clc_get_local_id(1) * clc_get_local_size(0) + clc_get_local_id(0);
+}
diff --git a/libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl b/libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl
new file mode 100644
index 0000000000000..317619de91178
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl
@@ -0,0 +1,18 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_size.h>
+#include <clc/workitem/clc_get_max_sub_group_size.h>
+#include <clc/workitem/clc_get_num_sub_groups.h>
+
+_CLC_OVERLOAD _CLC_DEF uint clc_get_num_sub_groups() {
+  size_t linear_size =
+      clc_get_local_size(0) * clc_get_local_size(1) * clc_get_local_size(2);
+  uint sg_size = clc_get_max_sub_group_size();
+  return (uint)((linear_size + sg_size - 1) / sg_size);
+}
diff --git a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
new file mode 100644
index 0000000000000..59ce1a9c367bd
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
@@ -0,0 +1,25 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+#include <clc/workitem/clc_get_local_size.h>
+#include <clc/workitem/clc_get_max_sub_group_size.h>
+#include <clc/workitem/clc_get_sub_group_id.h>
+
+_CLC_OVERLOAD _CLC_DEF uint clc_get_sub_group_id() {
+  // sreg.warpid is volatile and doesn't represent virtual warp index
+  // see https://docs.nvidia.com/cuda/parallel-thread-execution/index.html
+  size_t id_x = clc_get_local_id(0);
+  size_t id_y = clc_get_local_id(1);
+  size_t id_z = clc_get_local_id(2);
+  size_t size_x = clc_get_local_size(0);
+  size_t size_y = clc_get_local_size(1);
+  size_t size_z = clc_get_local_size(2);
+  uint sg_size = clc_get_max_sub_group_size();
+  return (id_z * size_y * size_x + id_y * size_x + id_x) / sg_size;
+}
diff --git a/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl b/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl
new file mode 100644
index 0000000000000..1457bc67a4176
--- /dev/null
+++ b/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl
@@ -0,0 +1,26 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_size.h>
+#include <clc/workitem/clc_get_max_sub_group_size.h>
+#include <clc/workitem/clc_get_num_sub_groups.h>
+#include <clc/workitem/clc_get_sub_group_id.h>
+#include <clc/workitem/clc_get_sub_group_size.h>
+
+_CLC_OVERLOAD _CLC_DEF uint clc_get_sub_group_size() {
+  if (clc_get_sub_group_id() != clc_get_num_sub_groups() - 1) {
+    return clc_get_max_sub_group_size();
+  }
+  size_t size_x = clc_get_local_size(0);
+  size_t size_y = clc_get_local_size(1);
+  size_t size_z = clc_get_local_size(2);
+  size_t linear_size = size_z * size_y * size_x;
+  size_t uniform_groups = clc_get_num_sub_groups() - 1;
+  size_t uniform_size = clc_get_max_sub_group_size() * uniform_groups;
+  return linear_size - uniform_size;
+}
diff --git a/libclc/clc/lib/ptx-nvidiacl/SOURCES b/libclc/clc/lib/ptx-nvidiacl/SOURCES
new file mode 100644
index 0000000000000..53b1c32de73c6
--- /dev/null
+++ b/libclc/clc/lib/ptx-nvidiacl/SOURCES
@@ -0,0 +1,7 @@
+workitem/clc_get_num_groups.cl
+workitem/clc_get_max_sub_group_size.cl
+workitem/clc_get_local_size.cl
+workitem/clc_get_local_id.cl
+workitem/clc_get_group_id.cl
+workitem/clc_get_sub_group_local_id.cl
+workitem/clc_get_global_size.cl
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
new file mode 100644
index 0000000000000..780341d3edb2a
--- /dev/null
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.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/workitem/clc_get_global_size.h>
+#include <clc/workitem/clc_get_local_size.h>
+#include <clc/workitem/clc_get_num_groups.h>
+
+_CLC_OVERLOAD _CLC_DEF 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/ptx-nvidiacl/workitem/get_group_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl
similarity index 85%
rename from libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl
rename to libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl
index 0dad4c2061fe6..38e46d9a8998c 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl
@@ -6,9 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_group_id.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t clc_get_group_id(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_ctaid_x();
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl
similarity index 84%
rename from libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl
rename to libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl
index 199b4610bdb7b..bcb408c7d8641 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl
@@ -6,9 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_local_id.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t clc_get_local_id(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_tid_x();
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl
similarity index 84%
rename from libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl
rename to libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl
index a93fa0a3c9649..e9342b4e0138b 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl
@@ -6,9 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_local_size.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t clc_get_local_size(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_ntid_x();
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl
new file mode 100644
index 0000000000000..f776301db6d1d
--- /dev/null
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl
@@ -0,0 +1,13 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 clc_get_max_sub_group_size() {
+  return __nvvm_read_ptx_sreg_warpsize();
+}
diff --git a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl
similarity index 84%
rename from libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl
rename to libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl
index 4c934968df865..ce6b9b99ac2fb 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl
@@ -6,9 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_num_groups.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t clc_get_num_groups(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_nctaid_x();
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl
new file mode 100644
index 0000000000000..481e030fa7af6
--- /dev/null
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl
@@ -0,0 +1,13 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 clc_get_sub_group_local_id() {
+  return __nvvm_read_ptx_sreg_laneid();
+}
diff --git a/libclc/opencl/include/clc/opencl/clc.h b/libclc/opencl/include/clc/opencl/clc.h
index 5859a00c3158b..520ccf21df3d6 100644
--- a/libclc/opencl/include/clc/opencl/clc.h
+++ b/libclc/opencl/include/clc/opencl/clc.h
@@ -36,13 +36,22 @@
 #include <clc/opencl/as_type.h>
 
 /* 6.11.1 Work-Item Functions */
+#include <clc/opencl/workitem/get_enqueued_local_size.h>
+#include <clc/opencl/workitem/get_enqueued_num_sub_groups.h>
 #include <clc/opencl/workitem/get_global_id.h>
+#include <clc/opencl/workitem/get_global_linear_id.h>
 #include <clc/opencl/workitem/get_global_offset.h>
 #include <clc/opencl/workitem/get_global_size.h>
 #include <clc/opencl/workitem/get_group_id.h>
 #include <clc/opencl/workitem/get_local_id.h>
+#include <clc/opencl/workitem/get_local_linear_id.h>
 #include <clc/opencl/workitem/get_local_size.h>
+#include <clc/opencl/workitem/get_max_sub_group_size.h>
 #include <clc/opencl/workitem/get_num_groups.h>
+#include <clc/opencl/workitem/get_num_sub_groups.h>
+#include <clc/opencl/workitem/get_sub_group_id.h>
+#include <clc/opencl/workitem/get_sub_group_local_id.h>
+#include <clc/opencl/workitem/get_sub_group_size.h>
 #include <clc/opencl/workitem/get_work_dim.h>
 
 /* 6.11.2 Math Functions */
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h b/libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h
new file mode 100644
index 0000000000000..58ac921de6203
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.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_OPENCL_WORKITEM_GET_ENQUEUED_LOCAL_SIZE_H__
+#define __CLC_OPENCL_WORKITEM_GET_ENQUEUED_LOCAL_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_enqueued_local_size(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_ENQUEUED_LOCAL_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h
new file mode 100644
index 0000000000000..5c2e2737063bf
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.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_OPENCL_WORKITEM_GET_ENQUEUED_NUM_SUB_GROUPS_H__
+#define __CLC_OPENCL_WORKITEM_GET_ENQUEUED_NUM_SUB_GROUPS_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_enqueued_num_sub_groups();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_ENQUEUED_NUM_SUB_GROUPS_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_id.h b/libclc/opencl/include/clc/opencl/workitem/get_global_id.h
index c60e9bc8b9b78..9172135fe6f8e 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_global_id.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_global_id(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h b/libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h
new file mode 100644
index 0000000000000..4ae60d730aa86
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/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_OPENCL_WORKITEM_GET_GLOBAL_LINEAR_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_LINEAR_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_global_linear_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_LINEAR_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h b/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h
index 7f06476048c52..4629dd16ef428 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_global_offset.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__
+#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_global_offset(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_OFFSET_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_size.h b/libclc/opencl/include/clc/opencl/workitem/get_global_size.h
index e235d990c79fb..b2e002ff1012f 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_size.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_global_size.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__
+#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_global_size(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_group_id.h b/libclc/opencl/include/clc/opencl/workitem/get_group_id.h
index 78b78e8e56922..5384ec4bce8cc 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_group_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_group_id.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_group_id(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_GROUP_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_id.h b/libclc/opencl/include/clc/opencl/workitem/get_local_id.h
index 82b569380d471..47acba78f996c 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_local_id.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_local_id.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_local_id(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h b/libclc/opencl/include/clc/opencl/workitem/get_local_linear_id.h
new file mode 100644
index 0000000000000..a73012affb3a2
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_local_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_OPENCL_WORKITEM_GET_LOCAL_LINEAR_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_LOCAL_LINEAR_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_local_linear_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_LINEAR_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_local_size.h b/libclc/opencl/include/clc/opencl/workitem/get_local_size.h
index 9458ba3923f7b..0407997482c7c 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_local_size.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_local_size.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__
+#define __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_local_size(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h b/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.h
new file mode 100644
index 0000000000000..207b5a4715954
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_max_sub_group_size.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_OPENCL_WORKITEM_GET_MAX_SUB_GROUP_SIZE_H__
+#define __CLC_OPENCL_WORKITEM_GET_MAX_SUB_GROUP_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_max_sub_group_size();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_MAX_SUB_GROUP_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h
index 3f0d3cb2c4fbd..9c511c30a14a3 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_num_groups.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim);
+#ifndef __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__
+#define __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_num_groups(uint dim);
+
+#endif // __CLC_OPENCL_WORKITEM_GET_NUM_GROUPS_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.h
new file mode 100644
index 0000000000000..b532ea100ab71
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_num_sub_groups.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_OPENCL_WORKITEM_GET_NUM_SUB_GROUPS_H__
+#define __CLC_OPENCL_WORKITEM_GET_NUM_SUB_GROUPS_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_num_sub_groups();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_NUM_SUB_GROUPS_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_id.h
new file mode 100644
index 0000000000000..eaf5303a3ff57
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_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_OPENCL_WORKITEM_GET_SUB_GROUP_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_sub_group_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_id.h
new file mode 100644
index 0000000000000..80fbff3955f76
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_local_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_OPENCL_WORKITEM_GET_SUB_GROUP_LOCAL_ID_H__
+#define __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_LOCAL_ID_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_sub_group_local_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_LOCAL_ID_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.h
new file mode 100644
index 0000000000000..613afb3f7b3e8
--- /dev/null
+++ b/libclc/opencl/include/clc/opencl/workitem/get_sub_group_size.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_OPENCL_WORKITEM_GET_SUB_GROUP_SIZE_H__
+#define __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_SIZE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_sub_group_size();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h b/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h
index dc6ae4e9f93bf..d54d5c4a8e3bd 100644
--- a/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h
+++ b/libclc/opencl/include/clc/opencl/workitem/get_work_dim.h
@@ -6,4 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-_CLC_DECL _CLC_OVERLOAD uint get_work_dim(void);
+#ifndef __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__
+#define __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_work_dim();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_WORK_DIM_H__
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 46ce6d6e36c24..6c4ac3cc95016 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -176,5 +176,20 @@ shared/max.cl
 shared/min.cl
 shared/vload.cl
 shared/vstore.cl
+workitem/get_enqueued_local_size.cl
+workitem/get_enqueued_num_sub_groups.cl
 workitem/get_global_id.cl
+workitem/get_global_linear_id.cl
+workitem/get_global_offset.cl
 workitem/get_global_size.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_groups.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
+workitem/get_work_dim.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl b/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
new file mode 100644
index 0000000000000..7c703770a3bd6
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_enqueued_local_size.h>
+#include <clc/workitem/clc_get_enqueued_local_size.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t get_enqueued_local_size(uint dim) {
+  return clc_get_enqueued_local_size(dim);
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl b/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
new file mode 100644
index 0000000000000..be2beca810b5d
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_enqueued_num_sub_groups.h>
+#include <clc/workitem/clc_get_enqueued_num_sub_groups.h>
+
+_CLC_OVERLOAD _CLC_DEF uint get_enqueued_num_sub_groups() {
+  return clc_get_enqueued_num_sub_groups();
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_global_id.cl b/libclc/opencl/lib/generic/workitem/get_global_id.cl
index 26c3bf528cd4d..d47c21c973539 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_id.cl
@@ -9,6 +9,8 @@
 #include <clc/opencl/clc.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
+  // FIXME call clc_get_global_id after amdgcn workitem functions are moved to
+  // clc.
   return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) +
          get_global_offset(dim);
 }
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..376d0c6957d7f
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_global_linear_id.h>
+#include <clc/workitem/clc_get_global_linear_id.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t get_global_linear_id() {
+  return clc_get_global_linear_id();
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_global_offset.cl b/libclc/opencl/lib/generic/workitem/get_global_offset.cl
new file mode 100644
index 0000000000000..a8057b7bf12d1
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_global_offset.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_global_offset.h>
+#include <clc/workitem/clc_get_global_offset.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t get_global_offset(uint dim) {
+  return clc_get_global_offset(dim);
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_global_size.cl b/libclc/opencl/lib/generic/workitem/get_global_size.cl
index 747115d524885..6dc59b1384e75 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_size.cl
@@ -6,8 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/clc.h>
+#include <clc/opencl/workitem/get_global_size.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);
+_CLC_OVERLOAD _CLC_DEF size_t get_global_size(uint dim) {
+  return clc_get_global_size(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_group_id.cl b/libclc/opencl/lib/generic/workitem/get_group_id.cl
new file mode 100644
index 0000000000000..8ceb47ffefbf6
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_group_id.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_group_id.h>
+#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/generic/workitem/get_local_id.cl b/libclc/opencl/lib/generic/workitem/get_local_id.cl
new file mode 100644
index 0000000000000..1309a4af0e5c2
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_local_id.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_local_id.h>
+#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/generic/workitem/get_local_linear_id.cl b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
new file mode 100644
index 0000000000000..425bd5218c943
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_local_linear_id.h>
+#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/generic/workitem/get_local_size.cl b/libclc/opencl/lib/generic/workitem/get_local_size.cl
new file mode 100644
index 0000000000000..f5dacabbaf983
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_local_size.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_local_size.h>
+#include <clc/workitem/clc_get_local_size.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t get_local_size(uint dim) {
+  return clc_get_local_size(dim);
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl b/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
new file mode 100644
index 0000000000000..23a33859cdcb1
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_max_sub_group_size.h>
+#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/generic/workitem/get_num_groups.cl b/libclc/opencl/lib/generic/workitem/get_num_groups.cl
new file mode 100644
index 0000000000000..eb53b0a292c25
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_num_groups.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_num_groups.h>
+#include <clc/workitem/clc_get_num_groups.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t get_num_groups(uint dim) {
+  return clc_get_num_groups(dim);
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl b/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
new file mode 100644
index 0000000000000..6ae834c9c16ca
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_num_sub_groups.h>
+#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/generic/workitem/get_sub_group_id.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
new file mode 100644
index 0000000000000..9ebe8b6a5bf62
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_sub_group_id.h>
+#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/generic/workitem/get_sub_group_local_id.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
new file mode 100644
index 0000000000000..050c3307758ab
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_sub_group_local_id.h>
+#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/generic/workitem/get_sub_group_size.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
new file mode 100644
index 0000000000000..d55f6d66acd8a
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_sub_group_size.h>
+#include <clc/workitem/clc_get_sub_group_size.h>
+
+_CLC_OVERLOAD _CLC_DEF uint get_sub_group_size() {
+  return clc_get_sub_group_size();
+}
diff --git a/libclc/opencl/lib/generic/workitem/get_work_dim.cl b/libclc/opencl/lib/generic/workitem/get_work_dim.cl
new file mode 100644
index 0000000000000..a94bb7caea91e
--- /dev/null
+++ b/libclc/opencl/lib/generic/workitem/get_work_dim.cl
@@ -0,0 +1,12 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/workitem/get_work_dim.h>
+#include <clc/workitem/clc_get_work_dim.h>
+
+_CLC_OVERLOAD _CLC_DEF uint clc_get_work_dim() { return clc_get_work_dim(); }
diff --git a/libclc/opencl/lib/ptx-nvidiacl/SOURCES b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
index c92c2a65d9aba..62e7ac8645ad7 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
@@ -1,7 +1,3 @@
 mem_fence/fence.cl
 synchronization/barrier.cl
 workitem/get_global_id.cl
-workitem/get_group_id.cl
-workitem/get_local_id.cl
-workitem/get_local_size.cl
-workitem/get_num_groups.cl

>From 3bab4ee8c7ff956fafac13975ca00a715b8ba460 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Tue, 17 Jun 2025 12:24:25 +0200
Subject: [PATCH 2/4] add __ prefix to clc functions

---
 .../workitem/clc_get_enqueued_local_size.h    |  2 +-
 .../clc_get_enqueued_num_sub_groups.h         |  2 +-
 .../include/clc/workitem/clc_get_global_id.h  |  2 +-
 .../clc/workitem/clc_get_global_linear_id.h   |  2 +-
 .../clc/workitem/clc_get_global_offset.h      |  2 +-
 .../clc/workitem/clc_get_global_size.h        |  2 +-
 .../include/clc/workitem/clc_get_group_id.h   |  2 +-
 .../include/clc/workitem/clc_get_local_id.h   |  2 +-
 .../clc/workitem/clc_get_local_linear_id.h    |  2 +-
 .../include/clc/workitem/clc_get_local_size.h |  2 +-
 .../clc/workitem/clc_get_max_sub_group_size.h |  2 +-
 .../include/clc/workitem/clc_get_num_groups.h |  2 +-
 .../clc/workitem/clc_get_num_sub_groups.h     |  2 +-
 .../clc/workitem/clc_get_sub_group_id.h       |  2 +-
 .../clc/workitem/clc_get_sub_group_local_id.h |  2 +-
 .../clc/workitem/clc_get_sub_group_size.h     |  2 +-
 .../include/clc/workitem/clc_get_work_dim.h   |  2 +-
 .../lib/generic/workitem/clc_get_global_id.cl |  6 ++---
 .../workitem/clc_get_global_linear_id.cl      | 22 +++++++++----------
 .../workitem/clc_get_local_linear_id.cl       |  8 ++++---
 .../workitem/clc_get_num_sub_groups.cl        |  8 +++----
 .../generic/workitem/clc_get_sub_group_id.cl  | 16 +++++++-------
 .../workitem/clc_get_sub_group_size.cl        | 16 +++++++-------
 .../workitem/clc_get_global_size.cl           |  4 ++--
 .../ptx-nvidiacl/workitem/clc_get_group_id.cl |  2 +-
 .../ptx-nvidiacl/workitem/clc_get_local_id.cl |  2 +-
 .../workitem/clc_get_local_size.cl            |  2 +-
 .../workitem/clc_get_max_sub_group_size.cl    |  2 +-
 .../workitem/clc_get_num_groups.cl            |  2 +-
 .../workitem/clc_get_sub_group_local_id.cl    |  2 +-
 .../workitem/get_enqueued_local_size.cl       |  2 +-
 .../workitem/get_enqueued_num_sub_groups.cl   |  2 +-
 .../lib/generic/workitem/get_global_id.cl     |  2 +-
 .../generic/workitem/get_global_linear_id.cl  |  2 +-
 .../lib/generic/workitem/get_global_offset.cl |  2 +-
 .../lib/generic/workitem/get_global_size.cl   |  2 +-
 .../lib/generic/workitem/get_group_id.cl      |  2 +-
 .../lib/generic/workitem/get_local_id.cl      |  2 +-
 .../generic/workitem/get_local_linear_id.cl   |  2 +-
 .../lib/generic/workitem/get_local_size.cl    |  2 +-
 .../workitem/get_max_sub_group_size.cl        |  2 +-
 .../lib/generic/workitem/get_num_groups.cl    |  2 +-
 .../generic/workitem/get_num_sub_groups.cl    |  2 +-
 .../lib/generic/workitem/get_sub_group_id.cl  |  2 +-
 .../workitem/get_sub_group_local_id.cl        |  2 +-
 .../generic/workitem/get_sub_group_size.cl    |  2 +-
 .../lib/generic/workitem/get_work_dim.cl      |  2 +-
 47 files changed, 81 insertions(+), 79 deletions(-)

diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
index 83aaed72c5036..cda0e6f3dbd93 100644
--- a/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_enqueued_local_size(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_enqueued_local_size(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_LOCAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
index 2a5af05f3f2d6..16b46075aa818 100644
--- a/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
+++ b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_DEF _CLC_OVERLOAD uint clc_get_enqueued_num_sub_groups();
+_CLC_DEF _CLC_OVERLOAD uint __clc_get_enqueued_num_sub_groups();
 
 #endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_id.h b/libclc/clc/include/clc/workitem/clc_get_global_id.h
index 697d7a629d794..da15ed8e53ca8 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_id(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_id(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_GLOBAL_ID_H__
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
index ac5a73ecd9e6e..0ab51ed8b064a 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_linear_id();
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_linear_id();
 
 #endif // __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_offset.h b/libclc/clc/include/clc/workitem/clc_get_global_offset.h
index ae51d88fea6a4..639be48d0b958 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_offset.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_offset.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_offset(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_offset(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_GLOBAL_OFFSET_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_global_size.h b/libclc/clc/include/clc/workitem/clc_get_global_size.h
index 67c315b1b6c38..fce11f417bf9d 100644
--- a/libclc/clc/include/clc/workitem/clc_get_global_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_global_size.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_global_size(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_global_size(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_GLOBAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_group_id.h b/libclc/clc/include/clc/workitem/clc_get_group_id.h
index a562fa175bf59..3fd43cb77f893 100644
--- a/libclc/clc/include/clc/workitem/clc_get_group_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_group_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_group_id(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_group_id(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_GROUP_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_id.h b/libclc/clc/include/clc/workitem/clc_get_local_id.h
index 482bb890d2159..c446127a5f90c 100644
--- a/libclc/clc/include/clc/workitem/clc_get_local_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_local_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_local_id(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_local_id(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_LOCAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h b/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
index ef27982cb176c..34c47f2afcb30 100644
--- a/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_local_linear_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_local_linear_id();
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_local_linear_id();
 
 #endif // __CLC_WORKITEM_CLC_GET_LOCAL_LINEAR_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_local_size.h b/libclc/clc/include/clc/workitem/clc_get_local_size.h
index 93fa68562f288..7fe3c193e0ec1 100644
--- a/libclc/clc/include/clc/workitem/clc_get_local_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_local_size.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_local_size(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_local_size(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_LOCAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
index 7593869650be4..8b521fe32862d 100644
--- a/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_max_sub_group_size.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_DEF _CLC_OVERLOAD uint clc_get_max_sub_group_size();
+_CLC_DEF _CLC_OVERLOAD uint __clc_get_max_sub_group_size();
 
 #endif // __CLC_WORKITEM_CLC_GET_MAX_SUB_GROUP_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_num_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_groups.h
index 03e0abb48ec33..f860944ad517f 100644
--- a/libclc/clc/include/clc/workitem/clc_get_num_groups.h
+++ b/libclc/clc/include/clc/workitem/clc_get_num_groups.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL size_t clc_get_num_groups(uint dim);
+_CLC_OVERLOAD _CLC_DECL size_t __clc_get_num_groups(uint dim);
 
 #endif // __CLC_WORKITEM_CLC_GET_NUM_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
index c17ebea1146c3..6965aef1bce30 100644
--- a/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
+++ b/libclc/clc/include/clc/workitem/clc_get_num_sub_groups.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_DEF _CLC_OVERLOAD uint clc_get_num_sub_groups();
+_CLC_DEF _CLC_OVERLOAD uint __clc_get_num_sub_groups();
 
 #endif // __CLC_WORKITEM_CLC_GET_NUM_SUB_GROUPS_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
index eadf08cb66f6b..ac3d7bd30e454 100644
--- a/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_DEF _CLC_OVERLOAD uint clc_get_sub_group_id();
+_CLC_DEF _CLC_OVERLOAD uint __clc_get_sub_group_id();
 
 #endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
index d631f37847a8a..06bb6f8b77a5c 100644
--- a/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_local_id.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_DEF _CLC_OVERLOAD uint clc_get_sub_group_local_id();
+_CLC_DEF _CLC_OVERLOAD uint __clc_get_sub_group_local_id();
 
 #endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_LOCAL_ID_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
index 56b736759a455..a6e8e49470e4f 100644
--- a/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
+++ b/libclc/clc/include/clc/workitem/clc_get_sub_group_size.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_DEF _CLC_OVERLOAD uint clc_get_sub_group_size();
+_CLC_DEF _CLC_OVERLOAD uint __clc_get_sub_group_size();
 
 #endif // __CLC_WORKITEM_CLC_GET_SUB_GROUP_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_work_dim.h b/libclc/clc/include/clc/workitem/clc_get_work_dim.h
index 1df6c4e819902..7e2da4570c4f5 100644
--- a/libclc/clc/include/clc/workitem/clc_get_work_dim.h
+++ b/libclc/clc/include/clc/workitem/clc_get_work_dim.h
@@ -11,6 +11,6 @@
 
 #include <clc/internal/clc.h>
 
-_CLC_OVERLOAD _CLC_DECL uint clc_get_work_dim();
+_CLC_OVERLOAD _CLC_DECL uint __clc_get_work_dim();
 
 #endif // __CLC_WORKITEM_CLC_GET_WORK_DIM_H__
diff --git a/libclc/clc/lib/generic/workitem/clc_get_global_id.cl b/libclc/clc/lib/generic/workitem/clc_get_global_id.cl
index bae19352a3e11..448109d9fe97b 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_global_id.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_global_id.cl
@@ -12,7 +12,7 @@
 #include <clc/workitem/clc_get_group_id.h>
 #include <clc/workitem/clc_get_local_id.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_global_id(uint dim) {
-  return clc_get_group_id(dim) * clc_get_enqueued_local_size(dim) +
-         clc_get_local_id(dim) + clc_get_global_offset(dim);
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_global_id(uint dim) {
+  return __clc_get_group_id(dim) * __clc_get_enqueued_local_size(dim) +
+         __clc_get_local_id(dim) + __clc_get_global_offset(dim);
 }
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
index 850c33f901d86..5b94afb419c77 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
@@ -12,21 +12,21 @@
 #include <clc/workitem/clc_get_global_size.h>
 #include <clc/workitem/clc_get_work_dim.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_global_linear_id() {
-  uint dim = clc_get_work_dim();
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_global_linear_id() {
+  uint dim = __clc_get_work_dim();
   switch (dim) {
   default:
   case 1:
-    return clc_get_global_id(0) - clc_get_global_offset(0);
+    return __clc_get_global_id(0) - __clc_get_global_offset(0);
   case 2:
-    return (clc_get_global_id(1) - clc_get_global_offset(1)) *
-               clc_get_global_size(0) +
-           (clc_get_global_id(0) - clc_get_global_offset(0));
+    return (__clc_get_global_id(1) - __clc_get_global_offset(1)) *
+               __clc_get_global_size(0) +
+           (__clc_get_global_id(0) - __clc_get_global_offset(0));
   case 3:
-    return ((clc_get_global_id(2) - clc_get_global_offset(2)) *
-            clc_get_global_size(1) * clc_get_global_size(0)) +
-           ((clc_get_global_id(1) - clc_get_global_offset(1)) *
-            clc_get_global_size(0)) +
-           (clc_get_global_id(0) - clc_get_global_offset(0));
+    return ((__clc_get_global_id(2) - __clc_get_global_offset(2)) *
+            __clc_get_global_size(1) * __clc_get_global_size(0)) +
+           ((__clc_get_global_id(1) - __clc_get_global_offset(1)) *
+            __clc_get_global_size(0)) +
+           (__clc_get_global_id(0) - __clc_get_global_offset(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 0ef7a238af606..fd7905568d595 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
@@ -10,7 +10,9 @@
 #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_get_local_size(0) +
-         clc_get_local_id(1) * clc_get_local_size(0) + clc_get_local_id(0);
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_linear_id() {
+  return __clc_get_local_id(2) * __clc_get_local_size(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/clc/lib/generic/workitem/clc_get_num_sub_groups.cl b/libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl
index 317619de91178..c61e838ebc693 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_num_sub_groups.cl
@@ -10,9 +10,9 @@
 #include <clc/workitem/clc_get_max_sub_group_size.h>
 #include <clc/workitem/clc_get_num_sub_groups.h>
 
-_CLC_OVERLOAD _CLC_DEF uint clc_get_num_sub_groups() {
-  size_t linear_size =
-      clc_get_local_size(0) * clc_get_local_size(1) * clc_get_local_size(2);
-  uint sg_size = clc_get_max_sub_group_size();
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_num_sub_groups() {
+  size_t linear_size = __clc_get_local_size(0) * __clc_get_local_size(1) *
+                       __clc_get_local_size(2);
+  uint sg_size = __clc_get_max_sub_group_size();
   return (uint)((linear_size + sg_size - 1) / sg_size);
 }
diff --git a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
index 59ce1a9c367bd..fa98e05550886 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_sub_group_id.cl
@@ -11,15 +11,15 @@
 #include <clc/workitem/clc_get_max_sub_group_size.h>
 #include <clc/workitem/clc_get_sub_group_id.h>
 
-_CLC_OVERLOAD _CLC_DEF uint clc_get_sub_group_id() {
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_sub_group_id() {
   // sreg.warpid is volatile and doesn't represent virtual warp index
   // see https://docs.nvidia.com/cuda/parallel-thread-execution/index.html
-  size_t id_x = clc_get_local_id(0);
-  size_t id_y = clc_get_local_id(1);
-  size_t id_z = clc_get_local_id(2);
-  size_t size_x = clc_get_local_size(0);
-  size_t size_y = clc_get_local_size(1);
-  size_t size_z = clc_get_local_size(2);
-  uint sg_size = clc_get_max_sub_group_size();
+  size_t id_x = __clc_get_local_id(0);
+  size_t id_y = __clc_get_local_id(1);
+  size_t id_z = __clc_get_local_id(2);
+  size_t size_x = __clc_get_local_size(0);
+  size_t size_y = __clc_get_local_size(1);
+  size_t size_z = __clc_get_local_size(2);
+  uint sg_size = __clc_get_max_sub_group_size();
   return (id_z * size_y * size_x + id_y * size_x + id_x) / sg_size;
 }
diff --git a/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl b/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl
index 1457bc67a4176..8ab4afe1ae05f 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl
+++ b/libclc/clc/lib/generic/workitem/clc_get_sub_group_size.cl
@@ -12,15 +12,15 @@
 #include <clc/workitem/clc_get_sub_group_id.h>
 #include <clc/workitem/clc_get_sub_group_size.h>
 
-_CLC_OVERLOAD _CLC_DEF uint clc_get_sub_group_size() {
-  if (clc_get_sub_group_id() != clc_get_num_sub_groups() - 1) {
-    return clc_get_max_sub_group_size();
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_sub_group_size() {
+  if (__clc_get_sub_group_id() != __clc_get_num_sub_groups() - 1) {
+    return __clc_get_max_sub_group_size();
   }
-  size_t size_x = clc_get_local_size(0);
-  size_t size_y = clc_get_local_size(1);
-  size_t size_z = clc_get_local_size(2);
+  size_t size_x = __clc_get_local_size(0);
+  size_t size_y = __clc_get_local_size(1);
+  size_t size_z = __clc_get_local_size(2);
   size_t linear_size = size_z * size_y * size_x;
-  size_t uniform_groups = clc_get_num_sub_groups() - 1;
-  size_t uniform_size = clc_get_max_sub_group_size() * uniform_groups;
+  size_t uniform_groups = __clc_get_num_sub_groups() - 1;
+  size_t uniform_size = __clc_get_max_sub_group_size() * uniform_groups;
   return linear_size - uniform_size;
 }
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
index 780341d3edb2a..3cba7463f3986 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
@@ -10,6 +10,6 @@
 #include <clc/workitem/clc_get_local_size.h>
 #include <clc/workitem/clc_get_num_groups.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_global_size(uint dim) {
-  return clc_get_num_groups(dim) * clc_get_local_size(dim);
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_global_size(uint dim) {
+  return __clc_get_num_groups(dim) * __clc_get_local_size(dim);
 }
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl
index 38e46d9a8998c..200b1bfd70992 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl
@@ -8,7 +8,7 @@
 
 #include <clc/workitem/clc_get_group_id.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_group_id(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_group_id(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_ctaid_x();
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl
index bcb408c7d8641..8c5b313b8135c 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl
@@ -8,7 +8,7 @@
 
 #include <clc/workitem/clc_get_local_id.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_local_id(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_id(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_tid_x();
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl
index e9342b4e0138b..4525c85f1e382 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl
@@ -8,7 +8,7 @@
 
 #include <clc/workitem/clc_get_local_size.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_local_size(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_ntid_x();
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl
index f776301db6d1d..6f0634d6f90e7 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_max_sub_group_size.cl
@@ -8,6 +8,6 @@
 
 #include <clc/workitem/clc_get_max_sub_group_size.h>
 
-_CLC_OVERLOAD _CLC_DEF uint clc_get_max_sub_group_size() {
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_max_sub_group_size() {
   return __nvvm_read_ptx_sreg_warpsize();
 }
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl
index ce6b9b99ac2fb..495864751ef68 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl
@@ -8,7 +8,7 @@
 
 #include <clc/workitem/clc_get_num_groups.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t clc_get_num_groups(uint dim) {
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_num_groups(uint dim) {
   switch (dim) {
   case 0:
     return __nvvm_read_ptx_sreg_nctaid_x();
diff --git a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl
index 481e030fa7af6..1448f8a8b97dc 100644
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_sub_group_local_id.cl
@@ -8,6 +8,6 @@
 
 #include <clc/workitem/clc_get_sub_group_local_id.h>
 
-_CLC_OVERLOAD _CLC_DEF uint clc_get_sub_group_local_id() {
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_sub_group_local_id() {
   return __nvvm_read_ptx_sreg_laneid();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl b/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
index 7c703770a3bd6..a2c040c6342d5 100644
--- a/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_enqueued_local_size.h>
 
 _CLC_OVERLOAD _CLC_DEF size_t get_enqueued_local_size(uint dim) {
-  return clc_get_enqueued_local_size(dim);
+  return __clc_get_enqueued_local_size(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl b/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
index be2beca810b5d..122b908554ebe 100644
--- a/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
+++ b/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_enqueued_num_sub_groups.h>
 
 _CLC_OVERLOAD _CLC_DEF uint get_enqueued_num_sub_groups() {
-  return clc_get_enqueued_num_sub_groups();
+  return __clc_get_enqueued_num_sub_groups();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_global_id.cl b/libclc/opencl/lib/generic/workitem/get_global_id.cl
index d47c21c973539..2ad3d9ff1051f 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_id.cl
@@ -9,7 +9,7 @@
 #include <clc/opencl/clc.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
-  // FIXME call clc_get_global_id after amdgcn workitem functions are moved to
+  // FIXME call __clc_get_global_id after amdgcn workitem functions are moved to
   // clc.
   return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) +
          get_global_offset(dim);
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 376d0c6957d7f..f3bf06553bd0e 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_global_linear_id.h>
 
 _CLC_OVERLOAD _CLC_DEF size_t get_global_linear_id() {
-  return clc_get_global_linear_id();
+  return __clc_get_global_linear_id();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_global_offset.cl b/libclc/opencl/lib/generic/workitem/get_global_offset.cl
index a8057b7bf12d1..c58696b50539f 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_offset.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_offset.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_global_offset.h>
 
 _CLC_OVERLOAD _CLC_DEF size_t get_global_offset(uint dim) {
-  return clc_get_global_offset(dim);
+  return __clc_get_global_offset(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_global_size.cl b/libclc/opencl/lib/generic/workitem/get_global_size.cl
index 6dc59b1384e75..4814208404dce 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_size.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_global_size.h>
 
 _CLC_OVERLOAD _CLC_DEF size_t get_global_size(uint dim) {
-  return clc_get_global_size(dim);
+  return __clc_get_global_size(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_group_id.cl b/libclc/opencl/lib/generic/workitem/get_group_id.cl
index 8ceb47ffefbf6..1bcd756bdd16e 100644
--- a/libclc/opencl/lib/generic/workitem/get_group_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_group_id.cl
@@ -10,5 +10,5 @@
 #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);
+  return __clc_get_group_id(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_local_id.cl b/libclc/opencl/lib/generic/workitem/get_local_id.cl
index 1309a4af0e5c2..f355b757b83b8 100644
--- a/libclc/opencl/lib/generic/workitem/get_local_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_local_id.cl
@@ -10,5 +10,5 @@
 #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);
+  return __clc_get_local_id(dim);
 }
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 425bd5218c943..8c7637b4fe31e 100644
--- a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
@@ -10,5 +10,5 @@
 #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();
+  return __clc_get_local_linear_id();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_local_size.cl b/libclc/opencl/lib/generic/workitem/get_local_size.cl
index f5dacabbaf983..b146c1ea1946d 100644
--- a/libclc/opencl/lib/generic/workitem/get_local_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_local_size.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_local_size.h>
 
 _CLC_OVERLOAD _CLC_DEF size_t get_local_size(uint dim) {
-  return clc_get_local_size(dim);
+  return __clc_get_local_size(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl b/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
index 23a33859cdcb1..fd76f7c5d0aa5 100644
--- a/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
@@ -10,5 +10,5 @@
 #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();
+  return __clc_get_max_sub_group_size();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_num_groups.cl b/libclc/opencl/lib/generic/workitem/get_num_groups.cl
index eb53b0a292c25..87e285d50205a 100644
--- a/libclc/opencl/lib/generic/workitem/get_num_groups.cl
+++ b/libclc/opencl/lib/generic/workitem/get_num_groups.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_num_groups.h>
 
 _CLC_OVERLOAD _CLC_DEF size_t get_num_groups(uint dim) {
-  return clc_get_num_groups(dim);
+  return __clc_get_num_groups(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl b/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
index 6ae834c9c16ca..1be77a325c0f3 100644
--- a/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
+++ b/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_num_sub_groups.h>
 
 _CLC_OVERLOAD _CLC_DEF uint get_num_sub_groups() {
-  return clc_get_num_sub_groups();
+  return __clc_get_num_sub_groups();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
index 9ebe8b6a5bf62..e538b86f12515 100644
--- a/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_sub_group_id.h>
 
 _CLC_OVERLOAD _CLC_DEF uint get_sub_group_id() {
-  return clc_get_sub_group_id();
+  return __clc_get_sub_group_id();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
index 050c3307758ab..2a8cfd57d85b2 100644
--- a/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
@@ -10,5 +10,5 @@
 #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();
+  return __clc_get_sub_group_local_id();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl b/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
index d55f6d66acd8a..f32d820d6418f 100644
--- a/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
@@ -10,5 +10,5 @@
 #include <clc/workitem/clc_get_sub_group_size.h>
 
 _CLC_OVERLOAD _CLC_DEF uint get_sub_group_size() {
-  return clc_get_sub_group_size();
+  return __clc_get_sub_group_size();
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_work_dim.cl b/libclc/opencl/lib/generic/workitem/get_work_dim.cl
index a94bb7caea91e..04ab7f599ac50 100644
--- a/libclc/opencl/lib/generic/workitem/get_work_dim.cl
+++ b/libclc/opencl/lib/generic/workitem/get_work_dim.cl
@@ -9,4 +9,4 @@
 #include <clc/opencl/workitem/get_work_dim.h>
 #include <clc/workitem/clc_get_work_dim.h>
 
-_CLC_OVERLOAD _CLC_DEF uint clc_get_work_dim() { return clc_get_work_dim(); }
+_CLC_OVERLOAD _CLC_DEF uint get_work_dim() { return __clc_get_work_dim(); }

>From 8b30c53d0959dfe1a9e49c1f8ee6c7f4695e351b Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Fri, 27 Jun 2025 11:31:23 +0200
Subject: [PATCH 3/4] move a few amdgcn functions into clc to add definition of
 __clc_get_work_dim/__clc_get_global_offset

---
 .../workitem/clc_get_enqueued_local_size.h    | 16 ----------
 .../clc_get_enqueued_num_sub_groups.h         | 16 ----------
 .../clc/workitem/clc_get_global_linear_id.h   | 16 ----------
 libclc/clc/lib/amdgcn/SOURCES                 |  5 +++
 .../amdgcn/workitem/clc_get_global_offset.cl} | 16 ++++++++--
 .../amdgcn/workitem/clc_get_global_size.cl}   | 16 +++++++---
 .../lib/amdgcn/workitem/clc_get_group_id.cl}  | 16 +++++++---
 .../lib/amdgcn/workitem/clc_get_local_id.cl}  | 16 +++++++---
 .../lib/amdgcn/workitem/clc_get_work_dim.cl}  | 14 ++++++--
 libclc/clc/lib/generic/SOURCES                |  4 +--
 .../workitem/clc_get_global_linear_id.cl      | 32 -------------------
 libclc/clc/lib/ptx-nvidiacl/SOURCES           | 11 ++++---
 .../workitem/clc_get_global_id.cl             |  7 ++--
 libclc/opencl/include/clc/opencl/clc.h        |  3 --
 .../opencl/workitem/get_enqueued_local_size.h | 16 ----------
 .../workitem/get_enqueued_num_sub_groups.h    | 16 ----------
 .../opencl/workitem/get_global_linear_id.h    | 16 ----------
 .../lib/amdgcn/workitem/get_global_offset.cl  | 14 ++------
 .../lib/amdgcn/workitem/get_global_size.cl    | 12 ++-----
 .../lib/amdgcn/workitem/get_group_id.cl       | 12 ++-----
 .../lib/amdgcn/workitem/get_local_id.cl       | 12 ++-----
 .../lib/amdgcn/workitem/get_work_dim.cl       |  6 ++--
 libclc/opencl/lib/generic/SOURCES             | 12 -------
 .../lib/generic/workitem/get_global_id.cl     |  3 --
 libclc/opencl/lib/ptx-nvidiacl/SOURCES        |  7 ++++
 .../ptx-nvidiacl/workitem/get_global_id.cl    |  5 +--
 .../workitem/get_local_linear_id.cl           |  0
 .../workitem/get_local_size.cl                |  0
 .../workitem/get_max_sub_group_size.cl        |  0
 .../workitem/get_num_sub_groups.cl            |  0
 .../workitem/get_sub_group_id.cl              |  0
 .../workitem/get_sub_group_local_id.cl        |  0
 .../workitem/get_sub_group_size.cl            |  0
 33 files changed, 96 insertions(+), 223 deletions(-)
 delete mode 100644 libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
 delete mode 100644 libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
 delete mode 100644 libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
 rename libclc/{opencl/lib/generic/workitem/get_global_offset.cl => clc/lib/amdgcn/workitem/clc_get_global_offset.cl} (51%)
 rename libclc/{opencl/lib/generic/workitem/get_enqueued_local_size.cl => clc/lib/amdgcn/workitem/clc_get_global_size.cl} (54%)
 rename libclc/{opencl/lib/generic/workitem/get_global_linear_id.cl => clc/lib/amdgcn/workitem/clc_get_group_id.cl} (53%)
 rename libclc/{opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl => clc/lib/amdgcn/workitem/clc_get_local_id.cl} (54%)
 rename libclc/{opencl/lib/generic/workitem/get_work_dim.cl => clc/lib/amdgcn/workitem/clc_get_work_dim.cl} (54%)
 delete mode 100644 libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
 rename libclc/clc/lib/{generic => ptx-nvidiacl}/workitem/clc_get_global_id.cl (70%)
 delete mode 100644 libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h
 delete mode 100644 libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h
 delete mode 100644 libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_local_linear_id.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_local_size.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_max_sub_group_size.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_num_sub_groups.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_sub_group_id.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_sub_group_local_id.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_sub_group_size.cl (100%)

diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
deleted file mode 100644
index cda0e6f3dbd93..0000000000000
--- a/libclc/clc/include/clc/workitem/clc_get_enqueued_local_size.h
+++ /dev/null
@@ -1,16 +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
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef __CLC_WORKITEM_CLC_GET_ENQUEUED_LOCAL_SIZE_H__
-#define __CLC_WORKITEM_CLC_GET_ENQUEUED_LOCAL_SIZE_H__
-
-#include <clc/internal/clc.h>
-
-_CLC_OVERLOAD _CLC_DECL size_t __clc_get_enqueued_local_size(uint dim);
-
-#endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_LOCAL_SIZE_H__
diff --git a/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h b/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
deleted file mode 100644
index 16b46075aa818..0000000000000
--- a/libclc/clc/include/clc/workitem/clc_get_enqueued_num_sub_groups.h
+++ /dev/null
@@ -1,16 +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
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__
-#define __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__
-
-#include <clc/internal/clc.h>
-
-_CLC_DEF _CLC_OVERLOAD uint __clc_get_enqueued_num_sub_groups();
-
-#endif // __CLC_WORKITEM_CLC_GET_ENQUEUED_NUM_SUB_GROUPS_H__
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
deleted file mode 100644
index 0ab51ed8b064a..0000000000000
--- a/libclc/clc/include/clc/workitem/clc_get_global_linear_id.h
+++ /dev/null
@@ -1,16 +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
-//
-//===----------------------------------------------------------------------===//
-
-#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_DECL size_t __clc_get_global_linear_id();
-
-#endif // __CLC_WORKITEM_CLC_GET_GLOBAL_LINEAR_ID_H__
diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
index 3a48049271aff..7bec1740f7636 100644
--- a/libclc/clc/lib/amdgcn/SOURCES
+++ b/libclc/clc/lib/amdgcn/SOURCES
@@ -1,3 +1,8 @@
 math/clc_fmax.cl
 math/clc_fmin.cl
 math/clc_ldexp_override.cl
+workitem/clc_get_global_offset.cl
+workitem/clc_get_global_size.cl
+workitem/clc_get_group_id.cl
+workitem/clc_get_local_id.cl
+workitem/clc_get_work_dim.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_global_offset.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl
similarity index 51%
rename from libclc/opencl/lib/generic/workitem/get_global_offset.cl
rename to libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl
index c58696b50539f..24b04fb1679d8 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_offset.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl
@@ -6,9 +6,19 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/workitem/get_global_offset.h>
 #include <clc/workitem/clc_get_global_offset.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t get_global_offset(uint dim) {
-  return __clc_get_global_offset(dim);
+#if __clang_major__ >= 8
+#define CONST_AS __constant
+#elif __clang_major__ >= 7
+#define CONST_AS __attribute__((address_space(4)))
+#else
+#define CONST_AS __attribute__((address_space(2)))
+#endif
+
+_CLC_DEF _CLC_OVERLOAD size_t __clc_get_global_offset(uint dim) {
+  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
+  if (dim < 3)
+    return ptr[dim + 1];
+  return 0;
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
similarity index 54%
rename from libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
rename to libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
index a2c040c6342d5..b1d8f27dc68c8 100644
--- a/libclc/opencl/lib/generic/workitem/get_enqueued_local_size.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl
@@ -6,9 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/workitem/get_enqueued_local_size.h>
-#include <clc/workitem/clc_get_enqueued_local_size.h>
+#include <clc/workitem/clc_get_global_size.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t get_enqueued_local_size(uint dim) {
-  return __clc_get_enqueued_local_size(dim);
+_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:
+    return 1;
+  }
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl
similarity index 53%
rename from libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
rename to libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl
index f3bf06553bd0e..aea927c3460b4 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_linear_id.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl
@@ -6,9 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/workitem/get_global_linear_id.h>
-#include <clc/workitem/clc_get_global_linear_id.h>
+#include <clc/workitem/clc_get_group_id.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t get_global_linear_id() {
-  return __clc_get_global_linear_id();
+_CLC_DEF _CLC_OVERLOAD size_t __clc_get_group_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __builtin_amdgcn_workgroup_id_x();
+  case 1:
+    return __builtin_amdgcn_workgroup_id_y();
+  case 2:
+    return __builtin_amdgcn_workgroup_id_z();
+  default:
+    return 1;
+  }
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl
similarity index 54%
rename from libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
rename to libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl
index 122b908554ebe..b7b7a43e735d3 100644
--- a/libclc/opencl/lib/generic/workitem/get_enqueued_num_sub_groups.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl
@@ -6,9 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/workitem/get_enqueued_num_sub_groups.h>
-#include <clc/workitem/clc_get_enqueued_num_sub_groups.h>
+#include <clc/workitem/clc_get_local_id.h>
 
-_CLC_OVERLOAD _CLC_DEF uint get_enqueued_num_sub_groups() {
-  return __clc_get_enqueued_num_sub_groups();
+_CLC_DEF _CLC_OVERLOAD size_t __clc_get_local_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __builtin_amdgcn_workitem_id_x();
+  case 1:
+    return __builtin_amdgcn_workitem_id_y();
+  case 2:
+    return __builtin_amdgcn_workitem_id_z();
+  default:
+    return 1;
+  }
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_work_dim.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl
similarity index 54%
rename from libclc/opencl/lib/generic/workitem/get_work_dim.cl
rename to libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl
index 04ab7f599ac50..93cab09b4109f 100644
--- a/libclc/opencl/lib/generic/workitem/get_work_dim.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl
@@ -6,7 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/workitem/get_work_dim.h>
 #include <clc/workitem/clc_get_work_dim.h>
 
-_CLC_OVERLOAD _CLC_DEF uint get_work_dim() { return __clc_get_work_dim(); }
+#if __clang_major__ >= 8
+#define CONST_AS __constant
+#elif __clang_major__ >= 7
+#define CONST_AS __attribute__((address_space(4)))
+#else
+#define CONST_AS __attribute__((address_space(2)))
+#endif
+
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_work_dim() {
+  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
+  return ptr[0];
+}
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index d840c54d2af20..bf8736a726315 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -152,9 +152,7 @@ shared/clc_max.cl
 shared/clc_min.cl
 shared/clc_vload.cl
 shared/clc_vstore.cl
-workitem/clc_get_sub_group_size.cl
-workitem/clc_get_global_id.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
+workitem/clc_get_sub_group_size.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
deleted file mode 100644
index 5b94afb419c77..0000000000000
--- a/libclc/clc/lib/generic/workitem/clc_get_global_linear_id.cl
+++ /dev/null
@@ -1,32 +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>
-#include <clc/workitem/clc_get_global_linear_id.h>
-#include <clc/workitem/clc_get_global_offset.h>
-#include <clc/workitem/clc_get_global_size.h>
-#include <clc/workitem/clc_get_work_dim.h>
-
-_CLC_OVERLOAD _CLC_DEF size_t __clc_get_global_linear_id() {
-  uint dim = __clc_get_work_dim();
-  switch (dim) {
-  default:
-  case 1:
-    return __clc_get_global_id(0) - __clc_get_global_offset(0);
-  case 2:
-    return (__clc_get_global_id(1) - __clc_get_global_offset(1)) *
-               __clc_get_global_size(0) +
-           (__clc_get_global_id(0) - __clc_get_global_offset(0));
-  case 3:
-    return ((__clc_get_global_id(2) - __clc_get_global_offset(2)) *
-            __clc_get_global_size(1) * __clc_get_global_size(0)) +
-           ((__clc_get_global_id(1) - __clc_get_global_offset(1)) *
-            __clc_get_global_size(0)) +
-           (__clc_get_global_id(0) - __clc_get_global_offset(0));
-  }
-}
diff --git a/libclc/clc/lib/ptx-nvidiacl/SOURCES b/libclc/clc/lib/ptx-nvidiacl/SOURCES
index 53b1c32de73c6..e058aa56a3d86 100644
--- a/libclc/clc/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/clc/lib/ptx-nvidiacl/SOURCES
@@ -1,7 +1,8 @@
-workitem/clc_get_num_groups.cl
-workitem/clc_get_max_sub_group_size.cl
-workitem/clc_get_local_size.cl
-workitem/clc_get_local_id.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
+workitem/clc_get_max_sub_group_size.cl
+workitem/clc_get_num_groups.cl
 workitem/clc_get_sub_group_local_id.cl
-workitem/clc_get_global_size.cl
diff --git a/libclc/clc/lib/generic/workitem/clc_get_global_id.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_id.cl
similarity index 70%
rename from libclc/clc/lib/generic/workitem/clc_get_global_id.cl
rename to libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_id.cl
index 448109d9fe97b..0782ef0b15813 100644
--- a/libclc/clc/lib/generic/workitem/clc_get_global_id.cl
+++ b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_id.cl
@@ -6,13 +6,12 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/workitem/clc_get_enqueued_local_size.h>
 #include <clc/workitem/clc_get_global_id.h>
-#include <clc/workitem/clc_get_global_offset.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>
 
 _CLC_OVERLOAD _CLC_DEF size_t __clc_get_global_id(uint dim) {
-  return __clc_get_group_id(dim) * __clc_get_enqueued_local_size(dim) +
-         __clc_get_local_id(dim) + __clc_get_global_offset(dim);
+  return __clc_get_group_id(dim) * __clc_get_local_size(dim) +
+         __clc_get_local_id(dim);
 }
diff --git a/libclc/opencl/include/clc/opencl/clc.h b/libclc/opencl/include/clc/opencl/clc.h
index 520ccf21df3d6..0df54e406ef87 100644
--- a/libclc/opencl/include/clc/opencl/clc.h
+++ b/libclc/opencl/include/clc/opencl/clc.h
@@ -36,10 +36,7 @@
 #include <clc/opencl/as_type.h>
 
 /* 6.11.1 Work-Item Functions */
-#include <clc/opencl/workitem/get_enqueued_local_size.h>
-#include <clc/opencl/workitem/get_enqueued_num_sub_groups.h>
 #include <clc/opencl/workitem/get_global_id.h>
-#include <clc/opencl/workitem/get_global_linear_id.h>
 #include <clc/opencl/workitem/get_global_offset.h>
 #include <clc/opencl/workitem/get_global_size.h>
 #include <clc/opencl/workitem/get_group_id.h>
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h b/libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h
deleted file mode 100644
index 58ac921de6203..0000000000000
--- a/libclc/opencl/include/clc/opencl/workitem/get_enqueued_local_size.h
+++ /dev/null
@@ -1,16 +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
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef __CLC_OPENCL_WORKITEM_GET_ENQUEUED_LOCAL_SIZE_H__
-#define __CLC_OPENCL_WORKITEM_GET_ENQUEUED_LOCAL_SIZE_H__
-
-#include <clc/internal/clc.h>
-
-_CLC_OVERLOAD _CLC_DECL size_t get_enqueued_local_size(uint dim);
-
-#endif // __CLC_OPENCL_WORKITEM_GET_ENQUEUED_LOCAL_SIZE_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h b/libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h
deleted file mode 100644
index 5c2e2737063bf..0000000000000
--- a/libclc/opencl/include/clc/opencl/workitem/get_enqueued_num_sub_groups.h
+++ /dev/null
@@ -1,16 +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
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef __CLC_OPENCL_WORKITEM_GET_ENQUEUED_NUM_SUB_GROUPS_H__
-#define __CLC_OPENCL_WORKITEM_GET_ENQUEUED_NUM_SUB_GROUPS_H__
-
-#include <clc/internal/clc.h>
-
-_CLC_OVERLOAD _CLC_DECL uint get_enqueued_num_sub_groups();
-
-#endif // __CLC_OPENCL_WORKITEM_GET_ENQUEUED_NUM_SUB_GROUPS_H__
diff --git a/libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h b/libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h
deleted file mode 100644
index 4ae60d730aa86..0000000000000
--- a/libclc/opencl/include/clc/opencl/workitem/get_global_linear_id.h
+++ /dev/null
@@ -1,16 +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
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef __CLC_OPENCL_WORKITEM_GET_GLOBAL_LINEAR_ID_H__
-#define __CLC_OPENCL_WORKITEM_GET_GLOBAL_LINEAR_ID_H__
-
-#include <clc/internal/clc.h>
-
-_CLC_OVERLOAD _CLC_DECL size_t get_global_linear_id();
-
-#endif // __CLC_OPENCL_WORKITEM_GET_GLOBAL_LINEAR_ID_H__
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl b/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl
index a1b3ce4192793..a7b3a5daf5b30 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl
@@ -7,18 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/clc.h>
-
-#if __clang_major__ >= 8
-#define CONST_AS __constant
-#elif __clang_major__ >= 7
-#define CONST_AS __attribute__((address_space(4)))
-#else
-#define CONST_AS __attribute__((address_space(2)))
-#endif
+#include <clc/workitem/clc_get_global_offset.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t get_global_offset(uint dim) {
-  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
-  if (dim < 3)
-    return ptr[dim + 1];
-  return 0;
+  return __clc_get_global_offset(dim);
 }
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl b/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl
index 8f1507765f934..cb8a60fd2f137 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_global_size.cl
@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_global_size.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t 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:
-    return 1;
-  }
+  return __clc_get_global_size(dim);
 }
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl b/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl
index 446cc63ab759d..7041cd751a31a 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_group_id.cl
@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_group_id.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
-  switch (dim) {
-  case 0:
-    return __builtin_amdgcn_workgroup_id_x();
-  case 1:
-    return __builtin_amdgcn_workgroup_id_y();
-  case 2:
-    return __builtin_amdgcn_workgroup_id_z();
-  default:
-    return 1;
-  }
+  return __clc_get_group_id(dim);
 }
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl b/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl
index cd07d8645cd53..415f1edd4aba2 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_local_id.cl
@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_local_id.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
-  switch (dim) {
-  case 0:
-    return __builtin_amdgcn_workitem_id_x();
-  case 1:
-    return __builtin_amdgcn_workitem_id_y();
-  case 2:
-    return __builtin_amdgcn_workitem_id_z();
-  default:
-    return 1;
-  }
+  return __clc_get_local_id(dim);
 }
diff --git a/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl b/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl
index 8ca8b0b61ce54..14b50d72ff3f8 100644
--- a/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl
+++ b/libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl
@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_work_dim.h>
 
 #if __clang_major__ >= 8
 #define CONST_AS __constant
@@ -16,7 +17,4 @@
 #define CONST_AS __attribute__((address_space(2)))
 #endif
 
-_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) {
-  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
-  return ptr[0];
-}
+_CLC_DEF _CLC_OVERLOAD uint get_work_dim() { return __clc_get_work_dim(); }
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 6c4ac3cc95016..369c521588a50 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -176,20 +176,8 @@ shared/max.cl
 shared/min.cl
 shared/vload.cl
 shared/vstore.cl
-workitem/get_enqueued_local_size.cl
-workitem/get_enqueued_num_sub_groups.cl
 workitem/get_global_id.cl
-workitem/get_global_linear_id.cl
-workitem/get_global_offset.cl
 workitem/get_global_size.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_groups.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
-workitem/get_work_dim.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_global_id.cl b/libclc/opencl/lib/generic/workitem/get_global_id.cl
index 2ad3d9ff1051f..78c0c2864724c 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_id.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_id.cl
@@ -5,12 +5,9 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-
 #include <clc/opencl/clc.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
-  // FIXME call __clc_get_global_id after amdgcn workitem functions are moved to
-  // clc.
   return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) +
          get_global_offset(dim);
 }
diff --git a/libclc/opencl/lib/ptx-nvidiacl/SOURCES b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
index 62e7ac8645ad7..69bb94082ed54 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
@@ -1,3 +1,10 @@
 mem_fence/fence.cl
 synchronization/barrier.cl
 workitem/get_global_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
+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
index 2f1f7cd4250fe..c4b0d8f805ec8 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl
+++ b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl
@@ -7,7 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/clc.h>
+#include <clc/workitem/clc_get_global_id.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
-  return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
+_CLC_OVERLOAD _CLC_DEF size_t get_global_id(uint dim) {
+  return __clc_get_global_id(dim);
 }
diff --git a/libclc/opencl/lib/generic/workitem/get_local_linear_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_linear_id.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_local_linear_id.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_linear_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_local_size.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_local_size.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_max_sub_group_size.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_max_sub_group_size.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_max_sub_group_size.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_sub_groups.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_num_sub_groups.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_sub_groups.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_id.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_sub_group_id.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_sub_group_local_id.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_sub_group_size.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_size.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_sub_group_size.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_sub_group_size.cl

>From b1397a4c0ba63a2fd2392aa064bfa351ce014a7d Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Thu, 3 Jul 2025 05:02:17 +0200
Subject: [PATCH 4/4] move get_group_id.cl, get_local_id.cl and
 get_num_groups.cl to ptx-nvidiacl, revert change to
 opencl/lib/generic/workitem/get_global_size.cl

---
 libclc/clc/lib/ptx-nvidiacl/SOURCES               |  1 -
 .../ptx-nvidiacl/workitem/clc_get_global_size.cl  | 15 ---------------
 libclc/opencl/lib/generic/SOURCES                 |  3 ---
 .../lib/generic/workitem/get_global_size.cl       |  7 +++----
 libclc/opencl/lib/ptx-nvidiacl/SOURCES            |  3 +++
 .../workitem/get_group_id.cl                      |  0
 .../workitem/get_local_id.cl                      |  0
 .../workitem/get_num_groups.cl                    |  0
 8 files changed, 6 insertions(+), 23 deletions(-)
 delete mode 100644 libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_group_id.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_local_id.cl (100%)
 rename libclc/opencl/lib/{generic => ptx-nvidiacl}/workitem/get_num_groups.cl (100%)

diff --git a/libclc/clc/lib/ptx-nvidiacl/SOURCES b/libclc/clc/lib/ptx-nvidiacl/SOURCES
index e058aa56a3d86..05368c5e4d4e3 100644
--- a/libclc/clc/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/clc/lib/ptx-nvidiacl/SOURCES
@@ -1,5 +1,4 @@
 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/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl b/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
deleted file mode 100644
index 3cba7463f3986..0000000000000
--- a/libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_size.cl
+++ /dev/null
@@ -1,15 +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>
-#include <clc/workitem/clc_get_local_size.h>
-#include <clc/workitem/clc_get_num_groups.h>
-
-_CLC_OVERLOAD _CLC_DEF 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/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 369c521588a50..46ce6d6e36c24 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -178,6 +178,3 @@ shared/vload.cl
 shared/vstore.cl
 workitem/get_global_id.cl
 workitem/get_global_size.cl
-workitem/get_group_id.cl
-workitem/get_local_id.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 4814208404dce..747115d524885 100644
--- a/libclc/opencl/lib/generic/workitem/get_global_size.cl
+++ b/libclc/opencl/lib/generic/workitem/get_global_size.cl
@@ -6,9 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/workitem/get_global_size.h>
-#include <clc/workitem/clc_get_global_size.h>
+#include <clc/opencl/clc.h>
 
-_CLC_OVERLOAD _CLC_DEF size_t get_global_size(uint dim) {
-  return __clc_get_global_size(dim);
+_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
+  return get_num_groups(dim) * get_local_size(dim);
 }
diff --git a/libclc/opencl/lib/ptx-nvidiacl/SOURCES b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
index 69bb94082ed54..eb28570a617af 100644
--- a/libclc/opencl/lib/ptx-nvidiacl/SOURCES
+++ b/libclc/opencl/lib/ptx-nvidiacl/SOURCES
@@ -1,9 +1,12 @@
 mem_fence/fence.cl
 synchronization/barrier.cl
 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_groups.cl
 workitem/get_num_sub_groups.cl
 workitem/get_sub_group_id.cl
 workitem/get_sub_group_local_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_group_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_group_id.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_local_id.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_local_id.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl
diff --git a/libclc/opencl/lib/generic/workitem/get_num_groups.cl b/libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl
similarity index 100%
rename from libclc/opencl/lib/generic/workitem/get_num_groups.cl
rename to libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl



More information about the cfe-commits mailing list