Merge from 'sycl' to 'sycl-web'

Author: iclsrc

sycl/doc/extensions/ExplicitSIMD/README.md

@@ -0,0 +1,93 @@
+# Intel "Explicit SIMD" SYCL extension
+
+OneAPI provides the "Explicit SIMD" SYCL extension (or simply "ESIMD") for
+lower-level Intel GPU programming. It provides APIs closely matching Intel GPU ISA
+yet allows to write explicitly vectorized device code. This helps programmer to
+have more control over the generated code and depend less on compiler
+optimizations. The [specification](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md),
+[documented ESIMD APIs headers](https://github.com/intel/llvm/tree/sycl/sycl/include/CL/sycl/INTEL/esimd) and
+[working code examples](https://github.com/intel/llvm/blob/sycl/sycl/test/esimd/on-device) are available on the Intel DPC++ project's github.
+
+**_NOTE:_** _This extension is under active development and lots of APIs are
+subject to change. There are currenly a number of restrictions specified
+below._
+
+ESIMD kernels and functions always require the subgroup size of one, which means
+compiler never does vectorization across workitems in a subgroup. Instead,
+vectorization is experessed explicitly in the code by the programmer. Here is a
+trivial example which adds elements of two arrays and writes the results to the
+third:
+
+```cpp
+  float *A = static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
+  float *B = static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
+  float *C = static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
+
+  for (unsigned i = 0; i < Size; ++i) {
+    A[i] = B[i] = i;
+  }
+
+  // We need that many workitems. Each processes VL elements of data.
+  cl::sycl::range<1> GlobalRange{Size / VL};
+  // Number of workitems in each workgroup.
+  cl::sycl::range<1> LocalRange{GroupSize};
+
+  cl::sycl::nd_range<1> Range(GlobalRange, LocalRange);
+
+  auto e = q.submit([&](handler &cgh) {
+    cgh.parallel_for<class Test>(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+      using namespace sycl::INTEL::gpu;
+
+      int i = ndi.get_global_id(0);
+      simd<float, VL> va = block_load<float, VL>(A + i * VL);
+      simd<float, VL> vb = block_load<float, VL>(B + i * VL);
+      simd<float, VL> vc = va + vb;
+      block_store<float, VL>(C + i * VL, vc);
+    });
+  });
+```
+
+In this example the lambda function passed to the `parallel_for` is marked with
+a special attribute - `SYCL_ESIMD_KERNEL`. This tells the compiler that this
+kernel is a ESIMD one and ESIMD APIs can be used inside it. Here the `simd`
+objects and `block_load`/`block_store` intrinsics are used which are avaiable
+only in the ESIMD extension.
+Full runnable code sample can be found on the
+[github repo](https://github.com/intel/llvm/blob/sycl/sycl/test/esimd/on-device/vadd_usm.cpp).
+
+#### Compiling and running ESIMD code.
+To compile a code which uses the ESIMD extension, a special compiler switch
+`-fsycl-explicit-simd` switch must be used:
+
+> `$ clang++ -fsycl -fsycl-explicit-simd vadd_usm.cpp`
+
+The resulting executable can only be run on Intel GPU hardware, such as
+Intel HD Graphics 600 or later. To run it, couple additional environment
+variables must be used - `SYCL_BE=PI_OPENCL` and
+`SYCL_PROGRAM_COMPILE_OPTIONS=-vc-codegen`:
+
+> `$ SYCL_BE=PI_OPENCL SYCL_PROGRAM_COMPILE_OPTIONS=-vc-codegen ./a.out`
+
+#### Restrictions
+
+Here is a list of main restrictions imposed on using ESIMD extension. Note that
+some of them are not enforced by the compiler, which may lead to undefined
+program behavior if violated.
+
+##### Features not supported with ESIMD extension:
+- Windows target
+- Ahead-of-time compilation
+- The [C and C++ Standard libraries support](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst)
+- The [Device library extensions](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst)
+- Host device (in some cases)
+
+##### Unsupported standard SYCL APIs:
+- Local accessors
+- Most of image APIs
+- Specialization constants
+- Memory access through a raw pointer returned by `sycl::accessor::get_pointer()`
+
+##### Other restrictions:
+- Only Intel GPU device is supported
+- Usual and ESIMD DPC++ kernels can not co-exist in the same application in most
+  cases

sycl/doc/extensions/SpecConstants/README.md

@@ -0,0 +1,72 @@
+# Specialization constants.
+
+Specialization constant is basically a variable in a SYCL program set by host
+code and used in device code which appears to be constant for the online (JIT)
+compiler of the device code. Things like optimal tile size in a tiled matrix
+multiplication kernel may depend on the hardware and can be expressed via a
+specialization constant for better code generation.
+
+This version of oneAPI provides experimental implementation of specialization
+constants based on the
+[proposal](https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md)
+from Codeplay.
+
+**NOTE:** _In future versions it may be superseded by [SYCL 2020
+specification](https://www.khronos.org/registry/SYCL/specs/sycl-2020-provisional.pdf)._
+
+A specialization constant is identified by a C++ type name, similarly to a
+kernel, its value is set via `program::set_spec_constant` class API and is
+"frozen" once the program is built. The following example shows how
+different values of a specialization constant can be used within the same
+kernel:
+
+```cpp
+  for (int i = 0; i < n_sc_sets; i++) {
+    cl::sycl::program program(q.get_context());
+    const int *sc_set = &sc_vals[i][0];
+    cl::sycl::ONEAPI::experimental::spec_constant<int32_t, SC0> sc0 =
+        program.set_spec_constant<SC0>(sc_set[0]);
+    cl::sycl::ONEAPI::experimental::spec_constant<int32_t, SC1> sc1 =
+        program.set_spec_constant<SC1>(sc_set[1]);
+
+    program.build_with_kernel_type<KernelAAA>();
+
+    try {
+      cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());
+
+      q.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
+        cgh.single_task<KernelAAA>(
+            program.get_kernel<KernelAAA>(),
+            [=]() {
+              acc[i] = sc0.get() + sc1.get();
+            });
+      });
+    } catch (cl::sycl::exception &e) {
+      std::cout << "*** Exception caught: " << e.what() << "\n";
+      return 1;
+    }
+    ...
+  }
+```
+Here the values of specialization constants `SC0` and `SC1` are changed on
+every loop iteration. All what's needed is re-creating a `program` class
+instance, setting new values and rebuilding it via
+`program::build_with_kernel_type`. JIT compiler will effectively replace
+`sc0.get()` and  `sc1.get()` within thhe device code with the corresponding
+constant values (`sc_vals[i][0]` and `sc_vals[i][1]`). Full runnable example
+can be found on
+[github](https://github.com/intel/llvm/blob/sycl/sycl/test/spec_const/spec_const_redefine.cpp).
+
+Specialization constants can be used in programs compiled Ahead-Of-Time, in this
+case a specialization constant takes default value for its type (as specified by
+[C++ standard](https://en.cppreference.com/w/cpp/language/value_initialization)).
+
+#### Limitations
+- The implementation does not support the `template <unsigned NID> struct spec_constant_id`
+  API design for interoperability with OpenCL - to set specializataion constants
+  in SYCL programs originating from external SPIRV modules and wrapped by OpenCL
+  program objects. In SPIRV/OpenCL specialization constants are identified by an
+  integer number, and the `spec_constant_id` class models that.
+- Only primitive numeric types are supported.
+

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

@@ -214,6 +214,19 @@ enum class kernel_sub_group : cl_kernel_sub_group_info {
   compile_sub_group_size = CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL
 };
 
+enum class kernel_device_specific : cl_kernel_work_group_info {
+  global_work_size = CL_KERNEL_GLOBAL_WORK_SIZE,
+  work_group_size = CL_KERNEL_WORK_GROUP_SIZE,
+  compile_work_group_size = CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
+  preferred_work_group_size_multiple =
+      CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
+  private_mem_size = CL_KERNEL_PRIVATE_MEM_SIZE,
+  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
+};
+
 // A.6 Program information desctiptors
 enum class program : cl_program_info {
   context = CL_PROGRAM_CONTEXT,
@@ -242,6 +255,8 @@ enum class event_profiling : cl_profiling_info {
 // Provide an alias to the return type for each of the info parameters
 template <typename T, T param> class param_traits {};
 
+template <typename T, T param> struct compatibility_param_traits {};
+
 #define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
   template <> class param_traits<param_type, param_type::param> {              \
   public:                                                                      \
@@ -263,6 +278,7 @@ template <typename T, T param> class param_traits {};
 
 #include <CL/sycl/info/event_profiling_traits.def>
 
+#include <CL/sycl/info/kernel_device_specific_traits.def>
 #include <CL/sycl/info/kernel_sub_group_traits.def>
 #include <CL/sycl/info/kernel_traits.def>
 #include <CL/sycl/info/kernel_work_group_traits.def>
@@ -276,6 +292,24 @@ template <typename T, T param> class param_traits {};
 #undef PARAM_TRAITS_SPEC
 #undef PARAM_TRAITS_SPEC_WITH_INPUT
 
+#define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
+  template <>                                                                  \
+  struct compatibility_param_traits<param_type, param_type::param> {           \
+    static constexpr auto value = kernel_device_specific::param;               \
+  };
+
+#define PARAM_TRAITS_SPEC_WITH_INPUT(param_type, param, ret_type, in_type)     \
+  template <>                                                                  \
+  struct compatibility_param_traits<param_type, param_type::param> {           \
+    static constexpr auto value = kernel_device_specific::param;               \
+  };
+
+#include <CL/sycl/info/kernel_sub_group_traits.def>
+#include <CL/sycl/info/kernel_work_group_traits.def>
+
+#undef PARAM_TRAITS_SPEC
+#undef PARAM_TRAITS_SPEC_WITH_INPUT
+
 } // namespace info
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/info/kernel_device_specific_traits.def

@@ -0,0 +1,12 @@
+PARAM_TRAITS_SPEC(kernel_device_specific, compile_work_group_size,
+                  cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_device_specific, global_work_size, cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_device_specific, 
+                  preferred_work_group_size_multiple, size_t)
+PARAM_TRAITS_SPEC(kernel_device_specific, private_mem_size, cl_ulong)
+PARAM_TRAITS_SPEC(kernel_device_specific, work_group_size, size_t)
+PARAM_TRAITS_SPEC_WITH_INPUT(kernel_device_specific, max_sub_group_size,
+                             uint32_t, cl::sycl::range<3>)
+PARAM_TRAITS_SPEC(kernel_device_specific, max_num_sub_groups, uint32_t)
+PARAM_TRAITS_SPEC(kernel_device_specific, compile_num_sub_groups, uint32_t)
+PARAM_TRAITS_SPEC(kernel_device_specific, compile_sub_group_size, uint32_t)

sycl/include/CL/sycl/kernel.hpp

@@ -92,6 +92,27 @@ class __SYCL_EXPORT kernel {
   typename info::param_traits<info::kernel, param>::return_type
   get_info() const;
 
+  /// Query device-specific information from the kernel object using the
+  /// info::kernel_device_specific descriptor.
+  ///
+  /// \param Device is a valid SYCL device to query info for.
+  /// \return depends on information being queried.
+  template <info::kernel_device_specific param>
+  typename info::param_traits<info::kernel_device_specific, param>::return_type
+  get_info(const device &Device) const;
+
+  /// Query device-specific information from a kernel using the
+  /// info::kernel_device_specific descriptor for a specific device and value.
+  ///
+  /// \param Device is a valid SYCL device.
+  /// \param Value depends on information being queried.
+  /// \return depends on information being queried.
+  template <info::kernel_device_specific param>
+  typename info::param_traits<info::kernel_device_specific, param>::return_type
+  get_info(const device &Device,
+           typename info::param_traits<info::kernel_device_specific,
+                                       param>::input_type Value) const;
+
   /// Query work-group information from a kernel using the
   /// info::kernel_work_group descriptor for a specific device.
   ///
@@ -107,8 +128,11 @@ class __SYCL_EXPORT kernel {
   /// \param Device is a valid SYCL device.
   /// \return depends on information being queried.
   template <info::kernel_sub_group param>
+  // clang-format off
   typename info::param_traits<info::kernel_sub_group, param>::return_type
+  __SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
   get_sub_group_info(const device &Device) const;
+  // clang-format on
 
   /// Query sub-group information from a kernel using the
   /// info::kernel_sub_group descriptor for a specific device and value.
@@ -117,11 +141,13 @@ class __SYCL_EXPORT kernel {
   /// \param Value depends on information being queried.
   /// \return depends on information being queried.
   template <info::kernel_sub_group param>
+  // clang-format off
   typename info::param_traits<info::kernel_sub_group, param>::return_type
-  get_sub_group_info(
-      const device &Device,
-      typename info::param_traits<info::kernel_sub_group, param>::input_type
-          Value) const;
+  __SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
+  get_sub_group_info(const device &Device,
+                     typename info::param_traits<info::kernel_sub_group,
+                     param>::input_type Value) const;
+  // clang-format on
 
 private:
   /// Constructs a SYCL kernel object from a valid kernel_impl instance.

sycl/plugins/cuda/pi_cuda.cpp

@@ -4133,7 +4133,7 @@ pi_result cuda_piextUSMHostAlloc(void **result_ptr, pi_context context,
   } catch (pi_result error) {
     result = error;
   }
-  assert(reinterpret_cast<std::uintptr_t>(*result_ptr) % alignment == 0);
+
   return result;
 }
 
@@ -4154,7 +4154,7 @@ pi_result cuda_piextUSMDeviceAlloc(void **result_ptr, pi_context context,
   } catch (pi_result error) {
     result = error;
   }
-  assert(reinterpret_cast<std::uintptr_t>(*result_ptr) % alignment == 0);
+
   return result;
 }
 
@@ -4176,7 +4176,7 @@ pi_result cuda_piextUSMSharedAlloc(void **result_ptr, pi_context context,
   } catch (pi_result error) {
     result = error;
   }
-  assert(reinterpret_cast<std::uintptr_t>(*result_ptr) % alignment == 0);
+
   return result;
 }
 

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -1441,7 +1441,16 @@ piextDeviceSelectBinary(pi_device Device, // TODO: does this need to be context?
                         pi_device_binary *Binaries, pi_uint32 NumBinaries,
                         pi_uint32 *SelectedBinaryInd) {
 
-  // TODO dummy implementation.
+  assert(Device);
+  assert(SelectedBinaryInd);
+  assert(NumBinaries == 0 || Binaries);
+
+  // TODO: this is a bare-bones implementation for choosing a device image
+  // that would be compatible with the targeted device. An AOT-compiled
+  // image is preferred over SPIR-V for known devices (i.e. Intel devices)
+  // The implementation makes no effort to differentiate between multiple images
+  // for the given device, and simply picks the first one compatible.
+  //
   // Real implementation will use the same mechanism OpenCL ICD dispatcher
   // uses. Something like:
   //   PI_VALIDATE_HANDLE_RETURN_HANDLE(ctx, PI_INVALID_CONTEXT);
@@ -1450,9 +1459,28 @@ piextDeviceSelectBinary(pi_device Device, // TODO: does this need to be context?
   // where context->dispatch is set to the dispatch table provided by PI
   // plugin for platform/device the ctx was created for.
 
+  // Look for GEN binary, which we known can only be handled by Level-Zero now.
+  const char *BinaryTarget = PI_DEVICE_BINARY_TARGET_SPIRV64_GEN;
+
+  // Find the appropriate device image, fallback to spirv if not found
   constexpr pi_uint32 InvalidInd = std::numeric_limits<pi_uint32>::max();
-  *SelectedBinaryInd = NumBinaries > 0 ? 0 : InvalidInd;
-  return PI_SUCCESS;
+  pi_uint32 Spirv = InvalidInd;
+
+  for (pi_uint32 i = 0; i < NumBinaries; ++i) {
+    if (strcmp(Binaries[i]->DeviceTargetSpec, BinaryTarget) == 0) {
+      *SelectedBinaryInd = i;
+      return PI_SUCCESS;
+    }
+    if (strcmp(Binaries[i]->DeviceTargetSpec,
+               PI_DEVICE_BINARY_TARGET_SPIRV64) == 0)
+      Spirv = i;
+  }
+  // Points to a spirv image, if such indeed was found
+  if ((*SelectedBinaryInd = Spirv) != InvalidInd)
+    return PI_SUCCESS;
+
+  // No image can be loaded for the given device
+  return PI_INVALID_BINARY;
 }
 
 pi_result piextDeviceGetNativeHandle(pi_device Device,
@@ -3100,7 +3128,20 @@ pi_result piEventGetProfilingInfo(pi_event Event, pi_profiling_info ParamName,
   case PI_PROFILING_INFO_COMMAND_END: {
     zeEventQueryKernelTimestamp(Event->ZeEvent, &tsResult);
 
+    uint64_t ContextStartTime = tsResult.context.kernelStart;
     uint64_t ContextEndTime = tsResult.context.kernelEnd;
+    //
+    // Handle a possible wrap-around (the underlying HW counter is < 64-bit).
+    // Note, it will not report correct time if there were multiple wrap
+    // arounds, and the longer term plan is to enlarge the capacity of the
+    // HW timestamps.
+    //
+    if (ContextEndTime <= ContextStartTime) {
+      pi_device Device = Event->Context->Devices[0];
+      const uint64_t TimestampMaxValue =
+          (1LL << Device->ZeDeviceProperties.kernelTimestampValidBits) - 1;
+      ContextEndTime += TimestampMaxValue - ContextStartTime;
+    }
     ContextEndTime *= ZeTimerResolution;
 
     return ReturnValue(uint64_t{ContextEndTime});