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

Changes in this PR:
* Declare most of 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.
* Move a few amdgcn workitem built-in implementations into clc.
* Include only needed headers in OpenCL workitem functions.
* Implement get_local_linear_id, get_max_sub_group_size,
get_num_sub_groups,
get_sub_group_id, get_sub_group_local_id, get_sub_group_size for
ptx-nvidiacl.

llvm-diff shows this PR adds a few new symbols to nvptx64--nvidiacl.bc.
llvm-diff shows no change to amdgcn--amdhsa.bc, nvptx--.bc and
nvptx64--.bc.

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__

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__

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__

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__

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__

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__

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__

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__

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__

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__

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__

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__

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__

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__

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

libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl

@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_offset.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
+
+_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;
+}

libclc/clc/lib/amdgcn/workitem/clc_get_global_size.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+_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;
+  }
+}

libclc/clc/lib/amdgcn/workitem/clc_get_group_id.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_group_id.h>
+
+_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;
+  }
+}

libclc/clc/lib/amdgcn/workitem/clc_get_local_id.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+_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;
+  }
+}

libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_work_dim.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
+
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_work_dim() {
+  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
+  return ptr[0];
+}

libclc/clc/lib/generic/SOURCES

@@ -152,3 +152,7 @@ shared/clc_max.cl
 shared/clc_min.cl
 shared/clc_vload.cl
 shared/clc_vstore.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

libclc/clc/lib/generic/workitem/clc_get_local_linear_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_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);
+}

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);
+}

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;
+}

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;
+}

libclc/clc/lib/ptx-nvidiacl/SOURCES

@@ -0,0 +1,7 @@
+workitem/clc_get_global_id.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

libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_global_id.cl

@@ -0,0 +1,17 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_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_local_size(dim) +
+         __clc_get_local_id(dim);
+}

libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_group_id.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_group_id.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_group_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __nvvm_read_ptx_sreg_ctaid_x();
+  case 1:
+    return __nvvm_read_ptx_sreg_ctaid_y();
+  case 2:
+    return __nvvm_read_ptx_sreg_ctaid_z();
+  default:
+    return 0;
+  }
+}

libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_id.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __nvvm_read_ptx_sreg_tid_x();
+  case 1:
+    return __nvvm_read_ptx_sreg_tid_y();
+  case 2:
+    return __nvvm_read_ptx_sreg_tid_z();
+  default:
+    return 0;
+  }
+}

libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_local_size.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_local_size(uint dim) {
+  switch (dim) {
+  case 0:
+    return __nvvm_read_ptx_sreg_ntid_x();
+  case 1:
+    return __nvvm_read_ptx_sreg_ntid_y();
+  case 2:
+    return __nvvm_read_ptx_sreg_ntid_z();
+  default:
+    return 0;
+  }
+}

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();
+}

libclc/clc/lib/ptx-nvidiacl/workitem/clc_get_num_groups.cl

@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_num_groups.h>
+
+_CLC_OVERLOAD _CLC_DEF size_t __clc_get_num_groups(uint dim) {
+  switch (dim) {
+  case 0:
+    return __nvvm_read_ptx_sreg_nctaid_x();
+  case 1:
+    return __nvvm_read_ptx_sreg_nctaid_y();
+  case 2:
+    return __nvvm_read_ptx_sreg_nctaid_z();
+  default:
+    return 0;
+  }
+}

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();
+}

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/opencl/opencl-base.h>
+
+_CLC_OVERLOAD _CLC_DECL size_t get_local_linear_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_LOCAL_LINEAR_ID_H__

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/opencl/opencl-base.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_max_sub_group_size();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_MAX_SUB_GROUP_SIZE_H__

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/opencl/opencl-base.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_num_sub_groups();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_NUM_SUB_GROUPS_H__

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/opencl/opencl-base.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_sub_group_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_ID_H__

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/opencl/opencl-base.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_sub_group_local_id();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_LOCAL_ID_H__

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/opencl/opencl-base.h>
+
+_CLC_OVERLOAD _CLC_DECL uint get_sub_group_size();
+
+#endif // __CLC_OPENCL_WORKITEM_GET_SUB_GROUP_SIZE_H__

libclc/opencl/lib/amdgcn/workitem/get_global_offset.cl

@@ -7,18 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_global_offset.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);
 }

libclc/opencl/lib/amdgcn/workitem/get_global_size.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #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) {
-  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);
 }

libclc/opencl/lib/amdgcn/workitem/get_group_id.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_group_id.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);
 }

libclc/opencl/lib/amdgcn/workitem/get_local_id.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_local_id.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);
 }

libclc/opencl/lib/amdgcn/workitem/get_work_dim.cl

@@ -7,16 +7,6 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_work_dim.h>
+#include <clc/workitem/clc_get_work_dim.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
-
-_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(); }

libclc/opencl/lib/ptx-nvidiacl/SOURCES

@@ -3,5 +3,11 @@ 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
+workitem/get_sub_group_size.cl

libclc/opencl/lib/ptx-nvidiacl/workitem/get_global_id.cl

@@ -7,10 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_global_id.h>
-#include <clc/opencl/workitem/get_group_id.h>
-#include <clc/opencl/workitem/get_local_id.h>
-#include <clc/opencl/workitem/get_local_size.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);
 }

libclc/opencl/lib/ptx-nvidiacl/workitem/get_group_id.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_group_id.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 __nvvm_read_ptx_sreg_ctaid_x();
-  case 1:
-    return __nvvm_read_ptx_sreg_ctaid_y();
-  case 2:
-    return __nvvm_read_ptx_sreg_ctaid_z();
-  default:
-    return 0;
-  }
+_CLC_OVERLOAD _CLC_DEF size_t get_group_id(uint dim) {
+  return __clc_get_group_id(dim);
 }

libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_id.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_local_id.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 __nvvm_read_ptx_sreg_tid_x();
-  case 1:
-    return __nvvm_read_ptx_sreg_tid_y();
-  case 2:
-    return __nvvm_read_ptx_sreg_tid_z();
-  default:
-    return 0;
-  }
+_CLC_OVERLOAD _CLC_DEF size_t get_local_id(uint dim) {
+  return __clc_get_local_id(dim);
 }

libclc/opencl/lib/ptx-nvidiacl/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();
+}

libclc/opencl/lib/ptx-nvidiacl/workitem/get_local_size.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_local_size.h>
+#include <clc/workitem/clc_get_local_size.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
-  switch (dim) {
-  case 0:
-    return __nvvm_read_ptx_sreg_ntid_x();
-  case 1:
-    return __nvvm_read_ptx_sreg_ntid_y();
-  case 2:
-    return __nvvm_read_ptx_sreg_ntid_z();
-  default:
-    return 0;
-  }
+_CLC_OVERLOAD _CLC_DEF size_t get_local_size(uint dim) {
+  return __clc_get_local_size(dim);
 }

libclc/opencl/lib/ptx-nvidiacl/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();
+}

libclc/opencl/lib/ptx-nvidiacl/workitem/get_num_groups.cl

@@ -7,16 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/workitem/get_num_groups.h>
+#include <clc/workitem/clc_get_num_groups.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
-  switch (dim) {
-  case 0:
-    return __nvvm_read_ptx_sreg_nctaid_x();
-  case 1:
-    return __nvvm_read_ptx_sreg_nctaid_y();
-  case 2:
-    return __nvvm_read_ptx_sreg_nctaid_z();
-  default:
-    return 0;
-  }
+_CLC_OVERLOAD _CLC_DEF size_t get_num_groups(uint dim) {
+  return __clc_get_num_groups(dim);
 }

libclc/opencl/lib/ptx-nvidiacl/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();
+}

libclc/opencl/lib/ptx-nvidiacl/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();
+}

libclc/opencl/lib/ptx-nvidiacl/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();
+}

libclc/opencl/lib/ptx-nvidiacl/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();
+}