https://github.com/wenju-he created
https://github.com/llvm/llvm-project/pull/144333
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.
>From 4ef7a1c91eb33e72cc6c91725d59d6e5b6791ab9 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju...@intel.com>
Date: Mon, 16 Jun 2025 13:22:29 +0200
Subject: [PATCH] [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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
