[SYCL][NativeCPU] Unify definitions of __spirv_* functions. (#18010)

Several functions were being defined both in libclc and in libdevice,
but had different definitions between the two. This commit changes it so
that they are only defined once, in libclc, with the definition that we
had in libdevice that we were relying on taking precedence.

Author: hvdijk

libclc/libspirv/lib/native_cpu/SOURCES

@@ -11,4 +11,8 @@ math/native_log2.cl
 math/native_sin.cl
 math/native_sqrt.cl
 math/round.cl
+workitem/get_global_id.cl
+workitem/get_global_size.cl
+workitem/get_num_sub_groups.cl
+workitem/get_sub_group_size.cl
 cl_khr_int64_extended_atomics/minmax_helpers.ll

libclc/libspirv/lib/native_cpu/workitem/get_global_id.cl

@@ -0,0 +1,23 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+ulong __mux_get_global_id(int);
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_x() {
+  return __mux_get_global_id(0);
+}
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_y() {
+  return __mux_get_global_id(1);
+}
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalInvocationId_z() {
+  return __mux_get_global_id(2);
+}

libclc/libspirv/lib/native_cpu/workitem/get_global_size.cl

@@ -0,0 +1,23 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+
+ulong __mux_get_global_size(int);
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalSize_x() {
+  return __mux_get_global_size(0);
+}
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalSize_y() {
+  return __mux_get_global_size(1);
+}
+
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalSize_z() {
+  return __mux_get_global_size(2);
+}

libclc/libspirv/lib/native_cpu/workitem/get_num_sub_groups.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 <libspirv/spirv.h>
+
+uint __mux_get_num_sub_groups();
+
+_CLC_DEF _CLC_OVERLOAD uint __spirv_NumSubgroups() {
+  return __mux_get_num_sub_groups();
+}

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

libdevice/nativecpu_utils.cpp

@@ -323,12 +323,6 @@ DefineShuffleVec2to16(int32_t, i32, int32_t);
 DefineShuffleVec2to16(uint32_t, i32, int32_t);
 DefineShuffleVec2to16(float, f32, float);
 
-#define Define2ArgForward(Type, Name, Callee)                                  \
-  DEVICE_EXTERNAL Type Name(Type a, Type b) noexcept { return Callee(a, b); }  \
-  static_assert(true)
-
-Define2ArgForward(uint64_t, __spirv_ocl_u_min, std::min);
-
 #define GET_PROPS __attribute__((pure))
 #define GEN_u32(bname, muxname)                                                \
   DEVICE_EXTERN_C GET_PROPS uint32_t muxname();                                \
@@ -338,8 +332,6 @@ Define2ArgForward(uint64_t, __spirv_ocl_u_min, std::min);
 GEN_u32(__spirv_SubgroupLocalInvocationId, __mux_get_sub_group_local_id);
 GEN_u32(__spirv_SubgroupMaxSize, __mux_get_max_sub_group_size);
 GEN_u32(__spirv_SubgroupId, __mux_get_sub_group_id);
-GEN_u32(__spirv_NumSubgroups, __mux_get_num_sub_groups);
-GEN_u32(__spirv_SubgroupSize, __mux_get_sub_group_size);
 
 // I64_I32
 #define GEN_p(bname, muxname, arg)                                             \
@@ -352,8 +344,6 @@ GEN_u32(__spirv_SubgroupSize, __mux_get_sub_group_size);
   GEN_p(bname##_y, ncpu_name, 1);                                              \
   GEN_p(bname##_z, ncpu_name, 2)
 
-GEN_xyz(__spirv_GlobalInvocationId, __mux_get_global_id);
-GEN_xyz(__spirv_GlobalSize, __mux_get_global_size);
 GEN_xyz(__spirv_GlobalOffset, __mux_get_global_offset);
 GEN_xyz(__spirv_LocalInvocationId, __mux_get_local_id);
 GEN_xyz(__spirv_NumWorkgroups, __mux_get_num_groups);