[SYCL] Switch to use of SPIR-V encoding for indexing built-ins.

This patch changes usages of OpenCL indexing biltins with special values
recognized by SPIRV translator, for instance:
get_global_id(0) replaced with VarLocalInvocationId.x

Signed-off-by: Romanov, Vlad <[email protected]>
Signed-off-by: Vladimir Lazarev <[email protected]>

Author: romanovvlad

sycl/include/CL/__spirv/spirv_vars.hpp

@@ -0,0 +1,41 @@
+//==---------- spirv_vars.hpp --- SPIRV variables -------------------------==//
+//
+// 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
+//
+// ===-------------------------------------------------------------------=== //
+
+#pragma once
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+namespace cl {
+namespace __spirv {
+typedef size_t size_t_vec __attribute__((ext_vector_type(3)));
+
+extern const __constant size_t_vec VarGlobalSize;
+extern const __constant size_t_vec VarGlobalInvocationId;
+extern const __constant size_t_vec VarWorkgroupSize;
+extern const __constant size_t_vec VarLocalInvocationId;
+extern const __constant size_t_vec VarWorkgroupId;
+extern const __constant size_t_vec VarGlobalOffset;
+
+#define DEFINE_INT_ID_TO_XYZ_CONVERTER(POSTFIX)                                \
+  template <int ID> static size_t get##POSTFIX();                              \
+  template <> static size_t get##POSTFIX<0>() { return Var##POSTFIX.x; }       \
+  template <> static size_t get##POSTFIX<1>() { return Var##POSTFIX.y; }       \
+  template <> static size_t get##POSTFIX<2>() { return Var##POSTFIX.z; }
+
+DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalSize);
+DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalInvocationId)
+DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupSize)
+DEFINE_INT_ID_TO_XYZ_CONVERTER(LocalInvocationId)
+DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupId)
+DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalOffset)
+
+#undef DEFINE_INT_ID_TO_XYZ_CONVERTER
+
+} // namespace __spirv
+} // namespace cl
+#endif // __SYCL_DEVICE_ONLY__

sycl/include/CL/sycl/handler.hpp

@@ -9,9 +9,11 @@
 
 #pragma once
 
+#include <CL/__spirv/spirv_vars.hpp>
 #include <CL/sycl/access/access.hpp>
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/scheduler/scheduler.h>
 #include <CL/sycl/event.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/kernel.hpp>
@@ -20,21 +22,10 @@
 #include <CL/sycl/property_list.hpp>
 #include <CL/sycl/stl.hpp>
 
-#include <CL/sycl/detail/scheduler/scheduler.h>
-
 #include <functional>
 #include <memory>
 #include <type_traits>
 
-#ifdef __SYCL_DEVICE_ONLY__
-size_t get_global_size(uint dimindx);
-size_t get_local_size(uint dimindx);
-size_t get_global_id(uint dimindx);
-size_t get_local_id(uint dimindx);
-size_t get_global_offset(uint dimindx);
-size_t get_group_id(uint dimindx);
-#endif
-
 template <typename T_src, int dim_src, cl::sycl::access::mode mode_src,
           cl::sycl::access::target tgt_src, typename T_dest, int dim_dest,
           cl::sycl::access::mode mode_dest, cl::sycl::access::target tgt_dest,
@@ -62,6 +53,48 @@ template <typename DataT, int Dimensions, access::mode AccessMode,
 class accessor;
 template <typename T, int Dimensions, typename AllocatorT> class buffer;
 namespace detail {
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+#define DEFINE_INIT_SIZES(POSTFIX)                                             \
+                                                                               \
+  template <int Dim, class DstT> struct InitSizesST##POSTFIX;                  \
+                                                                               \
+  template <class DstT> struct InitSizesST##POSTFIX<1, DstT> {                 \
+    static INLINE_IF_DEVICE void initSize(DstT &Dst) {                         \
+      Dst[0] = cl::__spirv::get##POSTFIX<0>();                                 \
+    }                                                                          \
+  };                                                                           \
+                                                                               \
+  template <class DstT> struct InitSizesST##POSTFIX<2, DstT> {                 \
+    static INLINE_IF_DEVICE void initSize(DstT &Dst) {                         \
+      Dst[1] = cl::__spirv::get##POSTFIX<1>();                                 \
+      InitSizesST##POSTFIX<1, DstT>::initSize(Dst);                            \
+    }                                                                          \
+  };                                                                           \
+                                                                               \
+  template <class DstT> struct InitSizesST##POSTFIX<3, DstT> {                 \
+    static INLINE_IF_DEVICE void initSize(DstT &Dst) {                         \
+      Dst[2] = cl::__spirv::get##POSTFIX<2>();                                 \
+      InitSizesST##POSTFIX<2, DstT>::initSize(Dst);                            \
+    }                                                                          \
+  };                                                                           \
+                                                                               \
+  template <int Dims, class DstT> static void init##POSTFIX(DstT &Dst) {       \
+    InitSizesST##POSTFIX<Dims, DstT>::initSize(Dst);                           \
+  }
+
+DEFINE_INIT_SIZES(GlobalSize);
+DEFINE_INIT_SIZES(GlobalInvocationId)
+DEFINE_INIT_SIZES(WorkgroupSize)
+DEFINE_INIT_SIZES(LocalInvocationId)
+DEFINE_INIT_SIZES(WorkgroupId)
+DEFINE_INIT_SIZES(GlobalOffset)
+
+#undef DEFINE_INIT_SIZES
+
+#endif //__SYCL_DEVICE_ONLY__
+
 class queue_impl;
 template <typename dataT, int dimensions, access::mode accessMode,
           access::target accessTarget, access::placeholder isPlaceholder,
@@ -263,9 +296,9 @@ class handler {
                                   (dimensions > 0 && dimensions < 4),
                               KernelType>::type kernelFunc) {
     id<dimensions> global_id;
-    for (int i = 0; i < dimensions; ++i) {
-      global_id[i] = get_global_id(i);
-    }
+
+    detail::initGlobalInvocationId<dimensions>(global_id);
+
     kernelFunc(global_id);
   }
 
@@ -277,10 +310,10 @@ class handler {
                               KernelType>::type kernelFunc) {
     id<dimensions> global_id;
     range<dimensions> global_size;
-    for (int i = 0; i < dimensions; ++i) {
-      global_id[i] = get_global_id(i);
-      global_size[i] = get_global_size(i);
-    }
+
+    detail::initGlobalInvocationId<dimensions>(global_id);
+    detail::initGlobalSize<dimensions>(global_size);
+
     item<dimensions, false> Item =
         detail::Builder::createItem<dimensions, false>(global_size, global_id);
     kernelFunc(Item);
@@ -299,14 +332,12 @@ class handler {
     id<dimensions> local_id;
     id<dimensions> global_offset;
 
-    for (int i = 0; i < dimensions; ++i) {
-      global_size[i] = get_global_size(i);
-      local_size[i] = get_local_size(i);
-      group_id[i] = get_group_id(i);
-      global_id[i] = get_global_id(i);
-      local_id[i] = get_local_id(i);
-      global_offset[i] = get_global_offset(i);
-    }
+    detail::initGlobalSize<dimensions>(global_size);
+    detail::initWorkgroupSize<dimensions>(local_size);
+    detail::initWorkgroupId<dimensions>(group_id);
+    detail::initGlobalInvocationId<dimensions>(global_id);
+    detail::initLocalInvocationId<dimensions>(local_id);
+    detail::initGlobalOffset<dimensions>(global_offset);
 
     group<dimensions> Group = detail::Builder::createGroup<dimensions>(
         global_size, local_size, group_id);