[libspirv] Use clc workitem functions for libspirv workitem

libclc/libspirv/lib/amdgcn-amdhsa/SOURCES

@@ -58,11 +58,14 @@ math/sinh.cl
 math/tan.cl
 math/tanh.cl
 workitem/get_global_size.cl
+workitem/get_local_linear_id.cl
 workitem/get_local_size.cl
 workitem/get_num_groups.cl
+workitem/get_num_sub_groups.cl
 workitem/get_max_sub_group_size.cl
 workitem/get_sub_group_id.cl
 workitem/get_sub_group_local_id.cl
+workitem/get_sub_group_size.cl
 misc/sub_group_shuffle.cl
 async/wait_group_events.cl
 assert/__assert_fail.ll

libclc/libspirv/lib/amdgcn-amdhsa/workitem/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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInLocalInvocationIndex() {
+  return __spirv_BuiltInLocalInvocationId(2) * __spirv_BuiltInWorkgroupSize(1) *
+             __spirv_BuiltInWorkgroupSize(0) +
+         __spirv_BuiltInLocalInvocationId(1) * __spirv_BuiltInWorkgroupSize(0) +
+         __spirv_BuiltInLocalInvocationId(0);
+}

libclc/libspirv/lib/amdgcn-amdhsa/workitem/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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInNumSubgroups() {
+  size_t size_x = __spirv_BuiltInWorkgroupSize(0);
+  size_t size_y = __spirv_BuiltInWorkgroupSize(1);
+  size_t size_z = __spirv_BuiltInWorkgroupSize(2);
+  uint sg_size = __spirv_BuiltInSubgroupMaxSize();
+  size_t linear_size = size_z * size_y * size_x;
+  return (uint)((linear_size + sg_size - 1) / sg_size);
+}

libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_sub_group_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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInSubgroupSize() {
+  if (__spirv_BuiltInSubgroupId() != __spirv_BuiltInNumSubgroups() - 1) {
+    return __spirv_BuiltInSubgroupMaxSize();
+  }
+  size_t size_x = __spirv_BuiltInWorkgroupSize(0);
+  size_t size_y = __spirv_BuiltInWorkgroupSize(1);
+  size_t size_z = __spirv_BuiltInWorkgroupSize(2);
+  size_t linear_size = size_z * size_y * size_x;
+  size_t uniform_groups = __spirv_BuiltInNumSubgroups() - 1;
+  size_t uniform_size = __spirv_BuiltInSubgroupMaxSize() * uniform_groups;
+  return linear_size - uniform_size;
+}

libclc/libspirv/lib/generic/SOURCES

@@ -179,4 +179,5 @@ workitem/get_global_id.cl
 workitem/get_global_size.cl
 workitem/get_local_linear_id.cl
 workitem/get_num_sub_groups.cl
+workitem/get_sub_group_id.cl
 workitem/get_sub_group_size.cl

libclc/libspirv/lib/generic/workitem/get_local_linear_id.cl

@@ -6,11 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_local_linear_id.h>
 #include <libspirv/spirv.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInLocalInvocationIndex() {
-  return __spirv_BuiltInLocalInvocationId(2) * __spirv_BuiltInWorkgroupSize(1) *
-             __spirv_BuiltInWorkgroupSize(0) +
-         __spirv_BuiltInLocalInvocationId(1) * __spirv_BuiltInWorkgroupSize(0) +
-         __spirv_BuiltInLocalInvocationId(0);
+_CLC_OVERLOAD _CLC_DEF size_t __spirv_BuiltInLocalInvocationIndex() {
+  return __clc_get_local_linear_id();
 }

libclc/libspirv/lib/generic/workitem/get_num_sub_groups.cl

@@ -6,13 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_num_sub_groups.h>
 #include <libspirv/spirv.h>
 
-_CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInNumSubgroups() {
-  size_t size_x = __spirv_BuiltInWorkgroupSize(0);
-  size_t size_y = __spirv_BuiltInWorkgroupSize(1);
-  size_t size_z = __spirv_BuiltInWorkgroupSize(2);
-  uint sg_size = __spirv_BuiltInSubgroupMaxSize();
-  size_t linear_size = size_z * size_y * size_x;
-  return (uint)((linear_size + sg_size - 1) / sg_size);
+_CLC_OVERLOAD _CLC_DEF uint __spirv_BuiltInNumSubgroups() {
+  return __clc_get_num_sub_groups();
 }

libclc/libspirv/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/workitem/clc_get_sub_group_id.h>
+#include <libspirv/spirv.h>
+
+_CLC_OVERLOAD _CLC_DEF uint __spirv_BuiltInSubgroupId() {
+  return __clc_get_sub_group_id();
+}

libclc/libspirv/lib/native_cpu/SOURCES

@@ -15,6 +15,8 @@ atomic/atomic_max.ll
 atomic/atomic_min.ll
 workitem/get_global_id.cl
 workitem/get_global_size.cl
+workitem/get_local_linear_id.cl
 workitem/get_num_sub_groups.cl
+workitem/get_sub_group_id.cl
 workitem/get_sub_group_size.cl
 cl_khr_int64_extended_atomics/minmax_helpers.ll

libclc/libspirv/lib/native_cpu/workitem/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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInLocalInvocationIndex() {
+  return __spirv_BuiltInLocalInvocationId(2) * __spirv_BuiltInWorkgroupSize(1) *
+             __spirv_BuiltInWorkgroupSize(0) +
+         __spirv_BuiltInLocalInvocationId(1) * __spirv_BuiltInWorkgroupSize(0) +
+         __spirv_BuiltInLocalInvocationId(0);
+}

libclc/libspirv/lib/native_cpu/workitem/get_sub_group_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_sub_group_id.h>
+#include <libspirv/spirv.h>
+
+uint __mux_get_sub_group_id();
+
+_CLC_OVERLOAD _CLC_DEF uint __spirv_BuiltInSubgroupId() {
+  return __mux_get_sub_group_id();
+}

libclc/libspirv/lib/ptx-nvidiacl/SOURCES

@@ -82,7 +82,6 @@ workitem/get_local_id.cl
 workitem/get_local_size.cl
 workitem/get_max_sub_group_size.cl
 workitem/get_num_groups.cl
-workitem/get_sub_group_id.cl
 workitem/get_sub_group_local_id.cl
 images/image_helpers.ll
 images/image.cl

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

@@ -6,17 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_group_id.h>
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInWorkgroupId(int 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;
-  }
+  return __clc_get_group_id(dim);
 }

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

@@ -6,17 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_local_id.h>
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInLocalInvocationId(int 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;
-  }
+  return __clc_get_local_id(dim);
 }

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

@@ -6,17 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_local_size.h>
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInWorkgroupSize(int 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 1;
-  }
+  return __clc_get_local_size(dim);
 }

libclc/libspirv/lib/ptx-nvidiacl/workitem/get_max_sub_group_size.cl

@@ -6,10 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_max_sub_group_size.h>
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInSubgroupMaxSize() {
-  return 32;
-  // FIXME: warpsize is defined by NVVM IR but doesn't compile if used here
-  // return __nvvm_read_ptx_sreg_warpsize();
+  return __clc_get_max_sub_group_size();
 }

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

@@ -6,17 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_num_groups.h>
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInNumWorkgroups(int 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 1;
-  }
+  return __clc_get_num_groups(dim);
 }

libclc/libspirv/lib/ptx-nvidiacl/workitem/get_sub_group_local_id.cl

@@ -6,8 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/workitem/clc_get_sub_group_local_id.h>
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInSubgroupLocalInvocationId() {
-  return __nvvm_read_ptx_sreg_laneid();
+  return __clc_get_sub_group_local_id();
 }

libclc/libspirv/lib/r600/SOURCES

@@ -2,6 +2,10 @@ workitem/get_global_offset.cl
 workitem/get_group_id.cl
 workitem/get_global_size.cl
 workitem/get_local_id.cl
+workitem/get_local_linear_id.cl
 workitem/get_local_size.cl
 workitem/get_num_groups.cl
+workitem/get_num_sub_groups.cl
+workitem/get_sub_group_id.cl
+workitem/get_sub_group_size.cl
 workitem/get_work_dim.cl

libclc/libspirv/lib/r600/workitem/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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInLocalInvocationIndex() {
+  return __spirv_BuiltInLocalInvocationId(2) * __spirv_BuiltInWorkgroupSize(1) *
+             __spirv_BuiltInWorkgroupSize(0) +
+         __spirv_BuiltInLocalInvocationId(1) * __spirv_BuiltInWorkgroupSize(0) +
+         __spirv_BuiltInLocalInvocationId(0);
+}

libclc/libspirv/lib/r600/workitem/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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInNumSubgroups() {
+  size_t size_x = __spirv_BuiltInWorkgroupSize(0);
+  size_t size_y = __spirv_BuiltInWorkgroupSize(1);
+  size_t size_z = __spirv_BuiltInWorkgroupSize(2);
+  uint sg_size = __spirv_BuiltInSubgroupMaxSize();
+  size_t linear_size = size_z * size_y * size_x;
+  return (uint)((linear_size + sg_size - 1) / sg_size);
+}

libclc/libspirv/lib/ptx-nvidiacl/workitem/get_sub_group_id.cl → libclc/libspirv/lib/r600/workitem/get_sub_group_id.cl

@@ -9,14 +9,11 @@
 #include <libspirv/spirv.h>
 
 _CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInSubgroupId() {
-  // 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 = __spirv_BuiltInLocalInvocationId(0);
   size_t id_y = __spirv_BuiltInLocalInvocationId(1);
   size_t id_z = __spirv_BuiltInLocalInvocationId(2);
   size_t size_x = __spirv_BuiltInWorkgroupSize(0);
   size_t size_y = __spirv_BuiltInWorkgroupSize(1);
-  size_t size_z = __spirv_BuiltInWorkgroupSize(2);
   uint sg_size = __spirv_BuiltInSubgroupMaxSize();
   return (id_z * size_y * size_x + id_y * size_x + id_x) / sg_size;
 }

libclc/libspirv/lib/r600/workitem/get_sub_group_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 <libspirv/spirv.h>
+
+_CLC_DEF _CLC_OVERLOAD uint __spirv_BuiltInSubgroupSize() {
+  if (__spirv_BuiltInSubgroupId() != __spirv_BuiltInNumSubgroups() - 1) {
+    return __spirv_BuiltInSubgroupMaxSize();
+  }
+  size_t size_x = __spirv_BuiltInWorkgroupSize(0);
+  size_t size_y = __spirv_BuiltInWorkgroupSize(1);
+  size_t size_z = __spirv_BuiltInWorkgroupSize(2);
+  size_t linear_size = size_z * size_y * size_x;
+  size_t uniform_groups = __spirv_BuiltInNumSubgroups() - 1;
+  size_t uniform_size = __spirv_BuiltInSubgroupMaxSize() * uniform_groups;
+  return linear_size - uniform_size;
+}