Merge remote-tracking branch 'origin/master' into gold/2021

Author: oleksandr-pavlyk

CHANGELOG.md

@@ -21,6 +21,8 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
 * Introduced `"syclinterface/dpctl_sycl_types_casters.hpp"` header file with declaration of conversion routines between SYCL type pointers and SyclInterface library opaque pointers [#960](https://github.com/IntelPython/dpctl/pull/960).
 * Added C-API to `dpctl.program.SyclKernel` and `dpctl.program.SyclProgram`. Added type casters for new types to "dpctl4pybind11" and added an example demonstrating its use [#970](https://github.com/IntelPython/dpctl/pull/970).
 * Introduced "dpctl/sycl.pxd" Cython declaration file to streamline use of SYCL functions from Cython, and added an example demonstrating its use [#981](https://github.com/IntelPython/dpctl/pull/981).
+* Added experimental support for sharing data allocated on sub-devices via dlpack [#984](https://github.com/IntelPython/dpctl/pull/984).
+* Added `dpctl.SyclDevice.sub_group_sizes` property to retrieve supported sizes of sub-group by the device [#985](https://github.com/IntelPython/dpctl/pull/985).
 
 ### Changed
 * Improved queue compatibility testing in `dpctl.tensor`'s implementation module [#900](https://github.com/IntelPython/dpctl/pull/900).

dpctl/_backend.pxd

@@ -204,6 +204,8 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h":
     cdef uint64_t DPCTLDevice_GetGlobalMemCacheSize(const DPCTLSyclDeviceRef DRef)
     cdef _global_mem_cache_type DPCTLDevice_GetGlobalMemCacheType(
         const DPCTLSyclDeviceRef DRef)
+    cdef size_t *DPCTLDevice_GetSubGroupSizes(const DPCTLSyclDeviceRef DRef,
+        size_t *res_len)
 
 
 cdef extern from "syclinterface/dpctl_sycl_device_manager.h":

dpctl/_sycl_device.pyx

@@ -65,6 +65,7 @@ from ._backend cimport (  # noqa: E211
     DPCTLDevice_GetPreferredVectorWidthShort,
     DPCTLDevice_GetProfilingTimerResolution,
     DPCTLDevice_GetSubGroupIndependentForwardProgress,
+    DPCTLDevice_GetSubGroupSizes,
     DPCTLDevice_GetVendor,
     DPCTLDevice_HasAspect,
     DPCTLDevice_Hash,
@@ -884,6 +885,28 @@ cdef class SyclDevice(_SyclDevice):
             self._device_ref
         )
 
+    @property
+    def sub_group_sizes(self):
+        """ Returns list of supported sub-group sizes for this device.
+
+        Returns:
+            List[int]: List of supported sub-group sizes.
+        """
+        cdef size_t *sg_sizes = NULL
+        cdef size_t sg_sizes_len = 0
+        cdef size_t i
+
+        sg_sizes = DPCTLDevice_GetSubGroupSizes(
+            self._device_ref, &sg_sizes_len)
+        if (sg_sizes is not NULL and sg_sizes_len > 0):
+            res = list()
+            for i in range(sg_sizes_len):
+                res.append(sg_sizes[i])
+            DPCTLSize_t_Array_Delete(sg_sizes)
+            return res
+        else:
+           return []
+
     @property
     def sycl_platform(self):
         """ Returns the platform associated with this device.

dpctl/_sycl_platform.pyx

@@ -272,7 +272,7 @@ cdef class SyclPlatform(_SyclPlatform):
         )
 
         if (CRef == NULL):
-            raise
+            raise RuntimeError("Getting default error ran into a problem")
         else:
             return SyclContext._create(CRef)
 

dpctl/tensor/_dlpack.pyx

@@ -33,12 +33,18 @@ from .._backend cimport (
 )
 from ._usmarray cimport usm_ndarray
 
+from platform import system as sys_platform
+
 import numpy as np
 
 import dpctl
 import dpctl.memory as dpmem
 
 
+cdef bint _IS_LINUX = sys_platform() == "Linux"
+
+del sys_platform
+
 cdef extern from 'dlpack/dlpack.h' nogil:
     cdef int DLPACK_VERSION
 
@@ -140,6 +146,7 @@ cpdef to_dlpack_capsule(usm_ndarray usm_ary) except+:
     cdef c_dpctl.SyclQueue ary_sycl_queue
     cdef c_dpctl.SyclDevice ary_sycl_device
     cdef DPCTLSyclDeviceRef pDRef = NULL
+    cdef DPCTLSyclDeviceRef tDRef = NULL
     cdef DLManagedTensor *dlm_tensor = NULL
     cdef DLTensor *dl_tensor = NULL
     cdef int nd = usm_ary.get_ndim()
@@ -157,19 +164,45 @@ cpdef to_dlpack_capsule(usm_ndarray usm_ary) except+:
     ary_sycl_queue = usm_ary.get_sycl_queue()
     ary_sycl_device = ary_sycl_queue.get_sycl_device()
 
-    # check that ary_sycl_device is a non-partitioned device
-    pDRef = DPCTLDevice_GetParentDevice(ary_sycl_device.get_device_ref())
-    if pDRef is not NULL:
-        DPCTLDevice_Delete(pDRef)
-        raise DLPackCreationError(
-            "to_dlpack_capsule: DLPack can only export arrays allocated on "
-            "non-partitioned SYCL devices."
-        )
-    default_context = dpctl.SyclQueue(ary_sycl_device).sycl_context
-    if not usm_ary.sycl_context == default_context:
+    try:
+        if _IS_LINUX:
+            default_context = ary_sycl_device.sycl_platform.default_context
+        else:
+            default_context = None
+    except RuntimeError:
+        # RT does not support default_context, e.g. Windows
+        default_context = None
+    if default_context is None:
+        # check that ary_sycl_device is a non-partitioned device
+        pDRef = DPCTLDevice_GetParentDevice(ary_sycl_device.get_device_ref())
+        if pDRef is not NULL:
+            DPCTLDevice_Delete(pDRef)
+            raise DLPackCreationError(
+                "to_dlpack_capsule: DLPack can only export arrays allocated "
+                "on non-partitioned SYCL devices on platforms where "
+                "default_context oneAPI extension is not supported."
+            )
+    else:
+        if not usm_ary.sycl_context == default_context:
+            raise DLPackCreationError(
+                "to_dlpack_capsule: DLPack can only export arrays based on USM "
+                "allocations bound to a default platform SYCL context"
+            )
+        # Find the unpartitioned parent of the allocation device
+        pDRef = DPCTLDevice_GetParentDevice(ary_sycl_device.get_device_ref())
+        if pDRef is not NULL:
+            tDRef = DPCTLDevice_GetParentDevice(pDRef)
+            while tDRef is not NULL:
+                DPCTLDevice_Delete(pDRef)
+                pDRef = tDRef
+                tDRef = DPCTLDevice_GetParentDevice(pDRef)
+            ary_sycl_device = c_dpctl.SyclDevice._create(pDRef)
+
+    # Find ordinal number of the parent device
+    device_id = ary_sycl_device.get_overall_ordinal()
+    if device_id < 0:
         raise DLPackCreationError(
-            "to_dlpack_capsule: DLPack can only export arrays based on USM "
-            "allocations bound to a default platform SYCL context"
+            "to_dlpack_capsule: failed to determine device_id"
         )
 
     dlm_tensor = <DLManagedTensor *> stdlib.malloc(
@@ -192,14 +225,6 @@ cpdef to_dlpack_capsule(usm_ndarray usm_ary) except+:
         for i in range(nd):
             shape_strides_ptr[nd + i] = strides_ptr[i]
 
-    device_id = ary_sycl_device.get_overall_ordinal()
-    if device_id < 0:
-        stdlib.free(shape_strides_ptr)
-        stdlib.free(dlm_tensor)
-        raise DLPackCreationError(
-            "to_dlpack_capsule: failed to determine device_id"
-        )
-
     ary_dt = usm_ary.dtype
     ary_dtk = ary_dt.kind
     element_offset = usm_ary.get_offset()
@@ -278,15 +303,16 @@ cpdef usm_ndarray from_dlpack_capsule(object py_caps) except +:
         success.
     Raises:
         TypeError: if argument is not a "dltensor" capsule.
-        ValueError: if argument is "used_dltensor" capsule,
-             if the USM pointer is not bound to the reconstructed
+        ValueError: if argument is "used_dltensor" capsule
+        BufferError:  if the USM pointer is not bound to the reconstructed
              sycl context, or the DLPack's device_type is not supported
              by dpctl.
     """
     cdef DLManagedTensor *dlm_tensor = NULL
     cdef bytes usm_type
     cdef size_t sz = 1
     cdef int i
+    cdef int device_id = -1
     cdef int element_bytesize = 0
     cdef Py_ssize_t offset_min = 0
     cdef Py_ssize_t offset_max = 0
@@ -308,26 +334,40 @@ cpdef usm_ndarray from_dlpack_capsule(object py_caps) except +:
             py_caps, "dltensor")
     # Verify that we can work with this device
     if dlm_tensor.dl_tensor.device.device_type == kDLOneAPI:
-        q = dpctl.SyclQueue(str(<int>dlm_tensor.dl_tensor.device.device_id))
+        device_id = dlm_tensor.dl_tensor.device.device_id
+        root_device = dpctl.SyclDevice(str(<int>device_id))
+        try:
+            if _IS_LINUX:
+                default_context = root_device.sycl_platform.default_context
+            else:
+                default_context = dpctl.SyclQueue(root_device).sycl_context
+        except RuntimeError:
+            default_context = dpctl.SyclQueue(root_device).sycl_context
         if dlm_tensor.dl_tensor.data is NULL:
             usm_type = b"device"
+            q = dpctl.SyclQueue(default_context, root_device)
         else:
             usm_type = c_dpmem._Memory.get_pointer_type(
                 <DPCTLSyclUSMRef> dlm_tensor.dl_tensor.data,
-                <c_dpctl.SyclContext>q.sycl_context)
-        if usm_type == b"unknown":
-            raise ValueError(
-                f"Data pointer in DLPack is not bound to default sycl "
-                "context of device '{device_id}', translated to "
-                "{q.sycl_device.filter_string}"
+                <c_dpctl.SyclContext>default_context)
+            if usm_type == b"unknown":
+                raise BufferError(
+                    "Data pointer in DLPack is not bound to default sycl "
+                    f"context of device '{device_id}', translated to "
+                    f"{root_device.filter_string}"
+                )
+            alloc_device = c_dpmem._Memory.get_pointer_device(
+                <DPCTLSyclUSMRef> dlm_tensor.dl_tensor.data,
+                <c_dpctl.SyclContext>default_context
             )
+            q = dpctl.SyclQueue(default_context, alloc_device)
         if dlm_tensor.dl_tensor.dtype.bits % 8:
-            raise ValueError(
+            raise BufferError(
                 "Can not import DLPack tensor whose element's "
                 "bitsize is not a multiple of 8"
             )
         if dlm_tensor.dl_tensor.dtype.lanes != 1:
-            raise ValueError(
+            raise BufferError(
                 "Can not import DLPack tensor with lanes != 1"
             )
         if dlm_tensor.dl_tensor.strides is NULL:

dpctl/tests/_device_attributes_checks.py

@@ -115,6 +115,11 @@ def check_max_num_sub_groups(device):
         assert max_num_sub_groups > 0
 
 
+def check_sub_group_sizes(device):
+    sg_sizes = device.sub_group_sizes
+    assert all(el > 0 for el in sg_sizes)
+
+
 def check_has_aspect_host(device):
     try:
         device.has_aspect_host
@@ -605,6 +610,7 @@ def check_global_mem_cache_line_size(device):
     check_max_work_item_sizes,
     check_max_work_group_size,
     check_max_num_sub_groups,
+    check_sub_group_sizes,
     check_is_accelerator,
     check_is_cpu,
     check_is_gpu,

examples/pybind11/use_dpctl_syclqueue/tests/test_queue_device.py

@@ -55,3 +55,11 @@ def test_offload_array_mod():
     Ynp = X % modulus_p
 
     assert np.array_equal(Y, Ynp)
+
+
+def test_get_sub_group_sizes():
+    d = dpctl.SyclDevice()
+    szs = uqd.get_sub_group_sizes(d)
+    assert type(szs) is list
+    assert all(type(el) is int for el in szs)
+    szs == d.sub_group_sizes

examples/pybind11/use_dpctl_syclqueue/use_queue_device/__init__.py

@@ -20,6 +20,7 @@
     get_device_global_mem_size,
     get_device_local_mem_size,
     get_max_compute_units,
+    get_sub_group_sizes,
     offloaded_array_mod,
 )
 
@@ -28,6 +29,7 @@
     "get_device_global_mem_size",
     "get_device_local_mem_size",
     "offloaded_array_mod",
+    "get_sub_group_sizes",
 ]
 
 __doc__ = """

examples/pybind11/use_dpctl_syclqueue/use_queue_device/_example.cpp

@@ -31,6 +31,7 @@
 #include <cstdint>
 #include <pybind11/numpy.h>
 #include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
 
 namespace py = pybind11;
 
@@ -84,6 +85,11 @@ offloaded_array_mod(sycl::queue q,
     return res;
 }
 
+std::vector<std::size_t> get_sub_group_sizes(const sycl::device &d)
+{
+    return d.get_info<sycl::info::device::sub_group_sizes>();
+}
+
 PYBIND11_MODULE(_use_queue_device, m)
 {
     m.def(
@@ -100,4 +106,6 @@ PYBIND11_MODULE(_use_queue_device, m)
           "Computes amount of local memory of the given dpctl.SyclDevice");
     m.def("offloaded_array_mod", &offloaded_array_mod,
           "Compute offloaded modular reduction of integer-valued NumPy array");
+    m.def("get_sub_group_sizes", &get_sub_group_sizes,
+          "Gets info::device::sub_group_sizes property of given device");
 }

libsyclinterface/include/dpctl_sycl_device_interface.h

@@ -651,4 +651,17 @@ DPCTL_API
 DPCTLGlobalMemCacheType
 DPCTLDevice_GetGlobalMemCacheType(__dpctl_keep const DPCTLSyclDeviceRef DRef);
 
+/*!
+ * @brief Wrapper for get_info<info::device::sub_group_sizes>().
+ *
+ * @param    DRef           Opaque pointer to a ``sycl::device``
+ * @param    res_len        Populated with size of the returned array
+ * @return   Returns the valid result if device exists else returns NULL.
+ * @ingroup DeviceInterface
+ */
+DPCTL_API
+__dpctl_keep size_t *
+DPCTLDevice_GetSubGroupSizes(__dpctl_keep const DPCTLSyclDeviceRef DRef,
+                             size_t *res_len);
+
 DPCTL_C_EXTERN_C_END