Merge from 'sycl' to 'sycl-web'

Author: iclsrc

sycl/CMakeLists.txt

@@ -11,7 +11,7 @@ option(SYCL_ADD_DEV_VERSION_POSTFIX "Adds -V postfix to version string" ON)
 set(SYCL_MAJOR_VERSION 1)
 set(SYCL_MINOR_VERSION 0)
 set(SYCL_PATCH_VERSION 0)
-set(SYCL_DEV_ABI_VERSION 0)
+set(SYCL_DEV_ABI_VERSION 1)
 if (SYCL_ADD_DEV_VERSION_POSTFIX)
   set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}")
 endif()

sycl/include/CL/sycl/detail/defines.hpp

@@ -8,6 +8,8 @@
 
 #pragma once
 
+#include <climits>
+
 #ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
 #define __SYCL_INLINE_NAMESPACE(X) inline namespace X
 #else
@@ -18,6 +20,10 @@
 #define __has_attribute(x) 0
 #endif
 
+#ifndef __has_builtin
+#define __has_builtin(x) 0
+#endif
+
 #if __has_attribute(always_inline)
 #define ALWAYS_INLINE __attribute__((always_inline))
 #else
@@ -31,3 +37,12 @@
 #ifndef SYCL_EXTERNAL
 #define SYCL_EXTERNAL
 #endif
+
+#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && __has_builtin(__builtin_assume)
+#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX)
+#else
+#define __SYCL_ASSUME_INT(x)
+#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && !__has_builtin(__builtin_assume)
+#warning "No assumptions will be emitted due to no __builtin_assume available"
+#endif
+#endif

sycl/include/CL/sycl/detail/pi.h

@@ -45,7 +45,7 @@
 // TODO: we need a mapping of PI to OpenCL somewhere, and this can be done
 // elsewhere, e.g. in the pi_opencl, but constants/enums mapping is now
 // done here, for efficiency and simplicity.
-
+#include <CL/cl_ext_intel.h>
 #include <CL/cl_usm_ext.h>
 #include <CL/sycl/detail/cl.h>
 #include <CL/sycl/detail/export.hpp>
@@ -133,16 +133,9 @@ typedef enum {
 // make the translation to OpenCL transparent.
 //
 typedef enum : pi_uint64 {
-  PI_DEVICE_TYPE_DEFAULT =
-      CL_DEVICE_TYPE_DEFAULT, ///< The default device available in the PI
-                              ///< plugin.
-  PI_DEVICE_TYPE_ALL =
-      CL_DEVICE_TYPE_ALL, ///< All devices available in the PI plugin.
-  PI_DEVICE_TYPE_CPU =
-      CL_DEVICE_TYPE_CPU, ///< A PI device that is the host processor.
-  PI_DEVICE_TYPE_GPU = CL_DEVICE_TYPE_GPU, ///< A PI device that is a GPU.
-  PI_DEVICE_TYPE_ACC = CL_DEVICE_TYPE_ACCELERATOR ///< A PI device that is a
-                                                  ///< dedicated accelerator.
+  PI_DEVICE_TYPE_CPU = CL_DEVICE_TYPE_CPU,
+  PI_DEVICE_TYPE_GPU = CL_DEVICE_TYPE_GPU,
+  PI_DEVICE_TYPE_ACC = CL_DEVICE_TYPE_ACCELERATOR
 } _pi_device_type;
 
 typedef enum {
@@ -316,6 +309,14 @@ typedef enum {
   PI_IMAGE_INFO_DEPTH = CL_IMAGE_DEPTH
 } _pi_image_info;
 
+typedef enum {
+  PI_KERNEL_MAX_SUB_GROUP_SIZE = CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
+  PI_KERNEL_MAX_NUM_SUB_GROUPS = CL_KERNEL_MAX_NUM_SUB_GROUPS,
+  PI_KERNEL_COMPILE_NUM_SUB_GROUPS = CL_KERNEL_COMPILE_NUM_SUB_GROUPS,
+  PI_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL =
+      CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL
+} _pi_kernel_sub_group_info;
+
 typedef enum {
   PI_EVENT_INFO_COMMAND_QUEUE = CL_EVENT_COMMAND_QUEUE,
   PI_EVENT_INFO_CONTEXT = CL_EVENT_CONTEXT,
@@ -510,6 +511,7 @@ using pi_queue_info = _pi_queue_info;
 using pi_image_info = _pi_image_info;
 using pi_kernel_info = _pi_kernel_info;
 using pi_kernel_group_info = _pi_kernel_group_info;
+using pi_kernel_sub_group_info = _pi_kernel_sub_group_info;
 using pi_event_info = _pi_event_info;
 using pi_command_type = _pi_command_type;
 using pi_mem_type = _pi_mem_type;
@@ -1056,9 +1058,26 @@ __SYCL_EXPORT pi_result piKernelGetGroupInfo(pi_kernel kernel, pi_device device,
                                              void *param_value,
                                              size_t *param_value_size_ret);
 
+/// API to query information from the sub-group from a kernel
+///
+/// \param kernel is the pi_kernel to query
+/// \param device is the device the kernel is executed on
+/// \param param_name is a pi_kernel_sub_group_info enum value that
+///        specifies the informtation queried for.
+/// \param input_value_size is the size of input value passed in
+///        ptr input_value param
+/// \param input_value is the ptr to the input value passed.
+/// \param param_value_size is the size of the value in bytes.
+/// \param param_value is a pointer to the value to set.
+/// \param param_value_size_ret is a pointer to return the size of data in
+///        param_value ptr.
+///
+/// All queries expect a return of 4 bytes in param_value_size,
+/// param_value_size_ret, and a uint32_t value should to be written in
+/// param_value ptr.
+/// Note: This behaviour differs from OpenCL. OpenCL returns size_t.
 __SYCL_EXPORT pi_result piKernelGetSubGroupInfo(
-    pi_kernel kernel, pi_device device,
-    cl_kernel_sub_group_info param_name, // TODO: untie from OpenCL
+    pi_kernel kernel, pi_device device, pi_kernel_sub_group_info param_name,
     size_t input_value_size, const void *input_value, size_t param_value_size,
     void *param_value, size_t *param_value_size_ret);
 

sycl/include/CL/sycl/id.hpp

@@ -96,8 +96,10 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
    * conversion:
    * int a = id<1>(value); */
 
-  operator EnableIfT<(dimensions == 1), size_t>() const {
-    return this->common_array[0];
+  ALWAYS_INLINE operator EnableIfT<(dimensions == 1), size_t>() const {
+    size_t Result = this->common_array[0];
+    __SYCL_ASSUME_INT(Result);
+    return Result;
   }
 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
 

sycl/include/CL/sycl/info/info_desc.hpp

@@ -206,9 +206,7 @@ enum class kernel_work_group : cl_kernel_work_group_info {
 };
 
 enum class kernel_sub_group : cl_kernel_sub_group_info {
-  max_sub_group_size_for_ndrange = CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
-  sub_group_count_for_ndrange = CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE,
-  local_size_for_sub_group_count = CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
+  max_sub_group_size = CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
   max_num_sub_groups = CL_KERNEL_MAX_NUM_SUB_GROUPS,
   compile_num_sub_groups = CL_KERNEL_COMPILE_NUM_SUB_GROUPS,
   compile_sub_group_size = CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL

sycl/include/CL/sycl/info/kernel_sub_group_traits.def

@@ -1,10 +1,5 @@
-PARAM_TRAITS_SPEC_WITH_INPUT(kernel_sub_group, max_sub_group_size_for_ndrange,
-                             size_t, cl::sycl::range<3>)
-PARAM_TRAITS_SPEC_WITH_INPUT(kernel_sub_group, sub_group_count_for_ndrange,
-                             size_t, cl::sycl::range<3>)
-PARAM_TRAITS_SPEC_WITH_INPUT(kernel_sub_group, local_size_for_sub_group_count,
-                             cl::sycl::range<3>, size_t)
-PARAM_TRAITS_SPEC(kernel_sub_group, max_num_sub_groups, size_t)
-PARAM_TRAITS_SPEC(kernel_sub_group, compile_num_sub_groups, size_t)
-PARAM_TRAITS_SPEC(kernel_sub_group, compile_sub_group_size, size_t)
-
+PARAM_TRAITS_SPEC_WITH_INPUT(kernel_sub_group, max_sub_group_size,
+                             uint32_t, cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_sub_group, max_num_sub_groups, uint32_t)
+PARAM_TRAITS_SPEC(kernel_sub_group, compile_num_sub_groups, uint32_t)
+PARAM_TRAITS_SPEC(kernel_sub_group, compile_sub_group_size, uint32_t)

sycl/include/CL/sycl/item.hpp

@@ -8,12 +8,15 @@
 
 #pragma once
 
+#include <CL/sycl/detail/defines.hpp>
 #include <CL/sycl/detail/helpers.hpp>
 #include <CL/sycl/detail/item_base.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/range.hpp>
 
+#include <cstddef>
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
@@ -28,22 +31,37 @@ template <int dimensions = 1, bool with_offset = true> class item {
 
   id<dimensions> get_id() const { return MImpl.MIndex; }
 
-  size_t get_id(int dimension) const { return MImpl.MIndex[dimension]; }
+  size_t ALWAYS_INLINE get_id(int dimension) const {
+    size_t Id = MImpl.MIndex[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
-  size_t operator[](int dimension) const { return MImpl.MIndex[dimension]; }
+  size_t ALWAYS_INLINE operator[](int dimension) const {
+    size_t Id = MImpl.MIndex[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   range<dimensions> get_range() const { return MImpl.MExtent; }
 
-  size_t get_range(int dimension) const { return MImpl.MExtent[dimension]; }
+  size_t ALWAYS_INLINE get_range(int dimension) const {
+    size_t Id = MImpl.MExtent[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   template <bool has_offset = with_offset>
   detail::enable_if_t<has_offset, id<dimensions>> get_offset() const {
     return MImpl.MOffset;
   }
 
   template <bool has_offset = with_offset>
-  detail::enable_if_t<has_offset, size_t> get_offset(int dimension) const {
-    return MImpl.MOffset[dimension];
+  detail::enable_if_t<has_offset, size_t>
+      ALWAYS_INLINE get_offset(int dimension) const {
+    size_t Id = MImpl.MOffset[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
   }
 
   template <bool has_offset = with_offset>
@@ -52,7 +70,11 @@ template <int dimensions = 1, bool with_offset = true> class item {
         MImpl.MExtent, MImpl.MIndex, /*Offset*/ {});
   }
 
-  size_t get_linear_id() const { return MImpl.get_linear_id(); }
+  size_t ALWAYS_INLINE get_linear_id() const {
+    size_t Id = MImpl.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   item(const item &rhs) = default;
 

sycl/include/CL/sycl/nd_item.hpp

@@ -19,6 +19,7 @@
 #include <CL/sycl/nd_range.hpp>
 #include <CL/sycl/range.hpp>
 
+#include <cstddef>
 #include <stdexcept>
 #include <type_traits>
 
@@ -33,34 +34,57 @@ template <int dimensions = 1> class nd_item {
 
   id<dimensions> get_global_id() const { return globalItem.get_id(); }
 
-  size_t get_global_id(int dimension) const {
-    return globalItem.get_id(dimension);
+  size_t ALWAYS_INLINE get_global_id(int dimension) const {
+    size_t Id = globalItem.get_id(dimension);
+    __SYCL_ASSUME_INT(Id);
+    return Id;
   }
 
-  size_t get_global_linear_id() const { return globalItem.get_linear_id(); }
+  size_t ALWAYS_INLINE get_global_linear_id() const {
+    size_t Id = globalItem.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   id<dimensions> get_local_id() const { return localItem.get_id(); }
 
-  size_t get_local_id(int dimension) const {
-    return localItem.get_id(dimension);
+  size_t ALWAYS_INLINE get_local_id(int dimension) const {
+    size_t Id = localItem.get_id(dimension);
+    __SYCL_ASSUME_INT(Id);
+    return Id;
   }
 
-  size_t get_local_linear_id() const { return localItem.get_linear_id(); }
+  size_t get_local_linear_id() const {
+    size_t Id = localItem.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   group<dimensions> get_group() const { return Group; }
 
   intel::sub_group get_sub_group() const { return intel::sub_group(); }
 
-  size_t get_group(int dimension) const { return Group[dimension]; }
+  size_t ALWAYS_INLINE get_group(int dimension) const {
+    size_t Size = Group[dimension];
+    __SYCL_ASSUME_INT(Size);
+    return Size;
+  }
 
-  size_t get_group_linear_id() const { return Group.get_linear_id(); }
+  size_t ALWAYS_INLINE get_group_linear_id() const {
+    size_t Id = Group.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   range<dimensions> get_group_range() const {
     return Group.get_global_range() / Group.get_local_range();
   }
 
-  size_t get_group_range(int dimension) const {
-    return Group.get_global_range(dimension) / Group.get_local_range(dimension);
+  size_t ALWAYS_INLINE get_group_range(int dimension) const {
+    size_t Range =
+        Group.get_global_range(dimension) / Group.get_local_range(dimension);
+    __SYCL_ASSUME_INT(Range);
+    return Range;
   }
 
   range<dimensions> get_global_range() const { return globalItem.get_range(); }
@@ -101,39 +125,36 @@ template <int dimensions = 1> class nd_item {
     Group.mem_fence();
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(local_ptr<dataT> dest,
                                      global_ptr<dataT> src,
                                      size_t numElements) const {
     return Group.async_work_group_copy(dest, src, numElements);
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(global_ptr<dataT> dest,
                                      local_ptr<dataT> src,
                                      size_t numElements) const {
     return Group.async_work_group_copy(dest, src, numElements);
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(local_ptr<dataT> dest,
-                                     global_ptr<dataT> src,
-                                     size_t numElements,
+                                     global_ptr<dataT> src, size_t numElements,
                                      size_t srcStride) const {
 
     return Group.async_work_group_copy(dest, src, numElements, srcStride);
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(global_ptr<dataT> dest,
-                                     local_ptr<dataT> src,
-                                     size_t numElements,
+                                     local_ptr<dataT> src, size_t numElements,
                                      size_t destStride) const {
     return Group.async_work_group_copy(dest, src, numElements, destStride);
   }
 
-  template<typename... eventTN>
-  void wait_for(eventTN... events) const {
+  template <typename... eventTN> void wait_for(eventTN... events) const {
     Group.wait_for(events...);
   }
 

sycl/plugins/cuda/pi_cuda.cpp

@@ -716,18 +716,15 @@ pi_result cuda_piDevicesGet(pi_platform platform, pi_device_type device_type,
                             pi_uint32 *num_devices) {
 
   pi_result err = PI_SUCCESS;
-  const bool askingForDefault = device_type == PI_DEVICE_TYPE_DEFAULT;
-  const bool askingForGPU = device_type & PI_DEVICE_TYPE_GPU;
-  const bool returnDevices = askingForDefault || askingForGPU;
-
-  size_t numDevices = returnDevices ? platform->devices_.size() : 0;
+  const bool askingForGPU = (device_type & PI_DEVICE_TYPE_GPU);
+  size_t numDevices = askingForGPU ? platform->devices_.size() : 0;
 
   try {
     if (num_devices) {
       *num_devices = numDevices;
     }
 
-    if (returnDevices && devices) {
+    if (askingForGPU && devices) {
       for (size_t i = 0; i < std::min(size_t(num_entries), numDevices); ++i) {
         devices[i] = platform->devices_[i].get();
       }