[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