Merge pull request #894 from IntelPython/gh-886-more-device-attributes

gh-886: Added 3 new device attributes and kernel's device-specific attributes

Author: oleksandr-pavlyk

dpctl/__init__.py

@@ -73,7 +73,12 @@
 from ._device_selection import select_device_with_aspects
 from ._sycl_timer import SyclTimer
 from ._version import get_versions
-from .enum_types import backend_type, device_type, event_status_type
+from .enum_types import (
+    backend_type,
+    device_type,
+    event_status_type,
+    global_mem_cache_type,
+)
 
 __all__ = [
     "SyclContext",
@@ -127,6 +132,7 @@
     "device_type",
     "backend_type",
     "event_status_type",
+    "global_mem_cache_type",
 ]
 __all__ += [
     "get_include",

dpctl/_backend.pxd

@@ -21,7 +21,7 @@
 types defined by dpctl's C API.
 """
 
-from libc.stdint cimport int64_t, uint32_t
+from libc.stdint cimport int64_t, uint32_t, uint64_t
 from libcpp cimport bool
 
 
@@ -112,6 +112,12 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h":
         _RUNNING            'DPCTL_RUNNING'
         _COMPLETE           'DPCTL_COMPLETE'
 
+    ctypedef enum _global_mem_cache_type 'DPCTLGlobalMemCacheType':
+        _MEM_CACHE_TYPE_INDETERMINATE    'DPCTL_MEM_CACHE_TYPE_INDETERMINATE'
+        _MEM_CACHE_TYPE_NONE             'DPCTL_MEM_CACHE_TYPE_NONE'
+        _MEM_CACHE_TYPE_READ_ONLY        'DPCTL_MEM_CACHE_TYPE_READ_ONLY'
+        _MEM_CACHE_TYPE_READ_WRITE       'DPCTL_MEM_CACHE_TYPE_READ_WRITE'
+
 
 cdef extern from "syclinterface/dpctl_sycl_types.h":
     cdef struct DPCTLOpaqueSyclContext
@@ -195,6 +201,10 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h":
         _partition_affinity_domain_type PartitionAffinityDomainTy)
     cdef DPCTLSyclDeviceRef DPCTLDevice_GetParentDevice(const DPCTLSyclDeviceRef DRef)
     cdef size_t DPCTLDevice_GetProfilingTimerResolution(const DPCTLSyclDeviceRef DRef)
+    cdef uint32_t DPCTLDevice_GetGlobalMemCacheLineSize(const DPCTLSyclDeviceRef DRef)
+    cdef uint64_t DPCTLDevice_GetGlobalMemCacheSize(const DPCTLSyclDeviceRef DRef)
+    cdef _global_mem_cache_type DPCTLDevice_GetGlobalMemCacheType(
+        const DPCTLSyclDeviceRef DRef)
 
 
 cdef extern from "syclinterface/dpctl_sycl_device_manager.h":
@@ -254,6 +264,14 @@ cdef extern from "syclinterface/dpctl_sycl_event_interface.h":
 cdef extern from "syclinterface/dpctl_sycl_kernel_interface.h":
     cdef size_t DPCTLKernel_GetNumArgs(const DPCTLSyclKernelRef KRef)
     cdef void DPCTLKernel_Delete(DPCTLSyclKernelRef KRef)
+    cdef size_t DPCTLKernel_GetWorkGroupSize(const DPCTLSyclKernelRef KRef)
+    cdef size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple(const DPCTLSyclKernelRef KRef)
+    cdef size_t DPCTLKernel_GetPrivateMemSize(const DPCTLSyclKernelRef KRef)
+    cdef uint32_t DPCTLKernel_GetMaxNumSubGroups(const DPCTLSyclKernelRef KRef)
+##  Next line is commented out due to issue in DPC++ runtime
+#   cdef uint32_t DPCTLKernel_GetMaxSubGroupSize(const DPCTLSyclKernelRef KRef)
+    cdef uint32_t DPCTLKernel_GetCompileNumSubGroups(const DPCTLSyclKernelRef KRef)
+    cdef uint32_t DPCTLKernel_GetCompileSubGroupSize(const DPCTLSyclKernelRef KRef)
 
 
 cdef extern from "syclinterface/dpctl_sycl_platform_manager.h":

dpctl/_sycl_device.pyx

@@ -34,6 +34,9 @@ from ._backend cimport (  # noqa: E211
     DPCTLDevice_GetBackend,
     DPCTLDevice_GetDeviceType,
     DPCTLDevice_GetDriverVersion,
+    DPCTLDevice_GetGlobalMemCacheLineSize,
+    DPCTLDevice_GetGlobalMemCacheSize,
+    DPCTLDevice_GetGlobalMemCacheType,
     DPCTLDevice_GetGlobalMemSize,
     DPCTLDevice_GetImage2dMaxHeight,
     DPCTLDevice_GetImage2dMaxWidth,
@@ -87,12 +90,13 @@ from ._backend cimport (  # noqa: E211
     _aspect_type,
     _backend_type,
     _device_type,
+    _global_mem_cache_type,
     _partition_affinity_domain_type,
 )
 
-from .enum_types import backend_type, device_type
+from .enum_types import backend_type, device_type, global_mem_cache_type
 
-from libc.stdint cimport int64_t, uint32_t
+from libc.stdint cimport int64_t, uint32_t, uint64_t
 from libc.stdlib cimport free, malloc
 
 from ._sycl_platform cimport SyclPlatform
@@ -1098,6 +1102,52 @@ cdef class SyclDevice(_SyclDevice):
             raise RuntimeError("Failed to get device timer resolution.")
         return timer_res
 
+    @property
+    def global_mem_cache_type(self):
+        """ Global device cache memory type.
+
+        Returns:
+            global_mem_cache_type: type of cache memory
+        Raises:
+            A RuntimeError is raised if an unrecognized memory type
+            is reported by runtime.
+        """
+        cdef _global_mem_cache_type gmcTy = (
+           DPCTLDevice_GetGlobalMemCacheType(self._device_ref)
+        )
+        if gmcTy == _global_mem_cache_type._MEM_CACHE_TYPE_READ_WRITE:
+            return global_mem_cache_type.read_write
+        elif gmcTy == _global_mem_cache_type._MEM_CACHE_TYPE_READ_ONLY:
+            return global_mem_cache_type.read_only
+        elif gmcTy == _global_mem_cache_type._MEM_CACHE_TYPE_NONE:
+            return global_mem_cache_type.none
+        elif gmcTy == _global_mem_cache_type._MEM_CACHE_TYPE_INDETERMINATE:
+            raise RuntimeError("Unrecognized global memory cache type reported")
+
+    @property
+    def global_mem_cache_size(self):
+        """ Global device memory cache size.
+
+        Returns:
+            int: Cache size in bytes
+        """
+        cdef uint64_t cache_sz = DPCTLDevice_GetGlobalMemCacheSize(
+            self._device_ref
+	)
+        return cache_sz
+
+    @property
+    def global_mem_cache_line_size(self):
+        """ Global device memory cache line size.
+
+        Returns:
+            int: Cache size in bytes
+        """
+        cdef uint64_t cache_line_sz = DPCTLDevice_GetGlobalMemCacheLineSize(
+            self._device_ref
+	)
+        return cache_line_sz
+
     cdef cpp_bool equals(self, SyclDevice other):
         """ Returns ``True`` if the :class:`dpctl.SyclDevice` argument has the
         same _device_ref as this SyclDevice.

dpctl/enum_types.py

@@ -96,3 +96,22 @@ class event_status_type(Enum):
     submitted = auto()
     running = auto()
     complete = auto()
+
+
+class global_mem_cache_type(Enum):
+    """
+    An enumeration of global memory cache types for a device.
+
+    :Example:
+        .. code-block:: python
+
+            import dpctl
+            dev = dpctl.SyclDevice()
+            print(dev.global_mem_cache_type)
+            # Possible output: <global_mem_cache_type.read_write: 4>
+    """
+
+    indeterminate = auto()
+    none = auto()
+    read_only = auto()
+    read_write = auto()

dpctl/program/_program.pyx

@@ -26,11 +26,18 @@ a OpenCL source string or a SPIR-V binary file.
 """
 
 cimport cython.array
+from libc.stdint cimport uint32_t
 
-from dpctl._backend cimport (  # noqa: E211, E402
+from dpctl._backend cimport (  # noqa: E211, E402;
     DPCTLCString_Delete,
     DPCTLKernel_Delete,
+    DPCTLKernel_GetCompileNumSubGroups,
+    DPCTLKernel_GetCompileSubGroupSize,
+    DPCTLKernel_GetMaxNumSubGroups,
     DPCTLKernel_GetNumArgs,
+    DPCTLKernel_GetPreferredWorkGroupSizeMultiple,
+    DPCTLKernel_GetPrivateMemSize,
+    DPCTLKernel_GetWorkGroupSize,
     DPCTLKernelBundle_CreateFromOCLSource,
     DPCTLKernelBundle_CreateFromSpirv,
     DPCTLKernelBundle_Delete,
@@ -95,6 +102,68 @@ cdef class SyclKernel:
         """
         return int(<size_t>self._kernel_ref)
 
+    @property
+    def num_args(self):
+        """ Property equivalent to method call `SyclKernel.get_num_args()`
+        """
+        return self.get_num_args()
+
+    @property
+    def work_group_size(self):
+        """ Returns the maximum number of work-items in a work-group that can
+        be used to execute the kernel on device it was built for.
+        """
+        cdef size_t v = DPCTLKernel_GetWorkGroupSize(self._kernel_ref)
+        return v
+
+    @property
+    def preferred_work_group_size_multiple(self):
+        """ Returns a value, of which work-group size is preferred to be
+        a multiple, for executing the kernel on the device it was built for.
+        """
+        cdef size_t v = DPCTLKernel_GetPreferredWorkGroupSizeMultiple(
+            self._kernel_ref
+	)
+        return v
+
+    @property
+    def private_mem_size(self):
+        """ Returns the minimum amount of private memory, in bytes, used by each
+        work-item in the kernel.
+        """
+        cdef size_t v = DPCTLKernel_GetPrivateMemSize(self._kernel_ref)
+        return v
+
+    @property
+    def max_num_sub_groups(self):
+        """ Returns the maximum number of sub-groups for this kernel.
+        """
+        cdef uint32_t n = DPCTLKernel_GetMaxNumSubGroups(self._kernel_ref)
+        return n
+
+    @property
+    def max_sub_group_size(self):
+        """ Returns the maximum sub-groups size for this kernel.
+        """
+        cdef uint32_t sz = 0
+        return NotImplemented
+
+    @property
+    def compile_num_sub_groups(self):
+        """ Returns the number of sub-groups specified by this kernel,
+        or 0 (if not specified).
+        """
+        cdef size_t n = DPCTLKernel_GetCompileNumSubGroups(self._kernel_ref)
+        return n
+
+    @property
+    def compile_sub_group_size(self):
+        """ Returns the required sub-group size specified by this kernel,
+        or 0 (if not specified).
+        """
+        cdef size_t n = DPCTLKernel_GetCompileSubGroupSize(self._kernel_ref)
+        return n
+
 
 cdef class SyclProgram:
     """ Wraps a ``sycl::kernel_bundle<sycl::bundle_state::executable>`` object