diff --git a/dpctl/__init__.py b/dpctl/__init__.py index 23bddf1222..2e8a70f470 100644 --- a/dpctl/__init__.py +++ b/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", diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index c07df5097d..dba49a36e0 100644 --- a/dpctl/_backend.pxd +++ b/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": diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index 2da48ef333..eff7b82bd8 100644 --- a/dpctl/_sycl_device.pyx +++ b/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 @@ -1097,6 +1101,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. diff --git a/dpctl/enum_types.py b/dpctl/enum_types.py index bdf95959c0..d83aac5f87 100644 --- a/dpctl/enum_types.py +++ b/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: + """ + + indeterminate = auto() + none = auto() + read_only = auto() + read_write = auto() diff --git a/dpctl/program/_program.pyx b/dpctl/program/_program.pyx index f638087793..7c2341a883 100644 --- a/dpctl/program/_program.pyx +++ b/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(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`` object diff --git a/dpctl/tests/_device_attributes_checks.py b/dpctl/tests/_device_attributes_checks.py index 59c458fedb..14c0b973a5 100644 --- a/dpctl/tests/_device_attributes_checks.py +++ b/dpctl/tests/_device_attributes_checks.py @@ -36,52 +36,52 @@ # Unit test cases that will be run for every device -def check_get_max_compute_units(device): +def check_max_compute_units(device): max_compute_units = device.max_compute_units assert max_compute_units > 0 -def check_get_global_mem_size(device): +def check_global_mem_size(device): global_mem_size = device.global_mem_size assert global_mem_size > 0 -def check_get_local_mem_size(device): +def check_local_mem_size(device): local_mem_size = device.local_mem_size assert local_mem_size > 0 -def check_get_max_work_item_dims(device): +def check_max_work_item_dims(device): max_work_item_dims = device.max_work_item_dims assert max_work_item_dims > 0 -def check_get_max_work_item_sizes1d(device): +def check_max_work_item_sizes1d(device): max_work_item_sizes = device.max_work_item_sizes1d for size in max_work_item_sizes: assert size is not None -def check_get_max_work_item_sizes2d(device): +def check_max_work_item_sizes2d(device): max_work_item_sizes = device.max_work_item_sizes2d for size in max_work_item_sizes: assert size is not None -def check_get_max_work_item_sizes3d(device): +def check_max_work_item_sizes3d(device): max_work_item_sizes = device.max_work_item_sizes3d for size in max_work_item_sizes: assert size is not None @pytest.mark.filterwarnings("DeprecationWarning:") -def check_get_max_work_item_sizes(device): +def check_max_work_item_sizes(device): max_work_item_sizes = device.max_work_item_sizes for size in max_work_item_sizes: assert size is not None -def check_get_max_work_group_size(device): +def check_max_work_group_size(device): max_work_group_size = device.max_work_group_size # Special case for FPGA simulator if device.is_accelerator: @@ -90,7 +90,7 @@ def check_get_max_work_group_size(device): assert max_work_group_size > 0 -def check_get_max_num_sub_groups(device): +def check_max_num_sub_groups(device): max_num_sub_groups = device.max_num_sub_groups # Special case for FPGA simulator if device.is_accelerator or device.is_host: @@ -267,105 +267,105 @@ def check_is_host(device): pytest.fail("is_hostcall failed") -def check_get_max_read_image_args(device): +def check_max_read_image_args(device): try: device.max_read_image_args except Exception: pytest.fail("max_read_image_args call failed") -def check_get_max_write_image_args(device): +def check_max_write_image_args(device): try: device.max_write_image_args except Exception: pytest.fail("max_write_image_args call failed") -def check_get_image_2d_max_width(device): +def check_image_2d_max_width(device): try: device.image_2d_max_width except Exception: pytest.fail("image_2d_max_width call failed") -def check_get_image_2d_max_height(device): +def check_image_2d_max_height(device): try: device.image_2d_max_height except Exception: pytest.fail("image_2d_max_height call failed") -def check_get_image_3d_max_width(device): +def check_image_3d_max_width(device): try: device.image_3d_max_width except Exception: pytest.fail("image_3d_max_width call failed") -def check_get_image_3d_max_height(device): +def check_image_3d_max_height(device): try: device.image_3d_max_height except Exception: pytest.fail("image_3d_max_height call failed") -def check_get_image_3d_max_depth(device): +def check_image_3d_max_depth(device): try: device.image_3d_max_depth except Exception: pytest.fail("image_3d_max_depth call failed") -def check_get_sub_group_independent_forward_progress(device): +def check_sub_group_independent_forward_progress(device): try: device.sub_group_independent_forward_progress except Exception: pytest.fail("sub_group_independent_forward_progress call failed") -def check_get_preferred_vector_width_char(device): +def check_preferred_vector_width_char(device): try: device.preferred_vector_width_char except Exception: pytest.fail("preferred_vector_width_char call failed") -def check_get_preferred_vector_width_short(device): +def check_preferred_vector_width_short(device): try: device.preferred_vector_width_short except Exception: pytest.fail("preferred_vector_width_short call failed") -def check_get_preferred_vector_width_int(device): +def check_preferred_vector_width_int(device): try: device.preferred_vector_width_int except Exception: pytest.fail("preferred_vector_width_int call failed") -def check_get_preferred_vector_width_long(device): +def check_preferred_vector_width_long(device): try: device.preferred_vector_width_long except Exception: pytest.fail("preferred_vector_width_long call failed") -def check_get_preferred_vector_width_float(device): +def check_preferred_vector_width_float(device): try: device.preferred_vector_width_float except Exception: pytest.fail("preferred_vector_width_float call failed") -def check_get_preferred_vector_width_double(device): +def check_preferred_vector_width_double(device): try: device.preferred_vector_width_double except Exception: pytest.fail("preferred_vector_width_double call failed") -def check_get_preferred_vector_width_half(device): +def check_preferred_vector_width_half(device): try: device.preferred_vector_width_half except Exception: @@ -514,27 +514,93 @@ def check_platform(device): assert isinstance(p, dpctl.SyclPlatform) +def check_parent_device(device): + pd = device.parent_device + assert pd is None or isinstance(pd, dpctl.SyclDevice) + + +def check_filter_string(device): + try: + fs = device.filter_string + assert type(fs) is str + dd = dpctl.SyclDevice(fs) + assert device == dd + except TypeError: + pass + + +def check_name(device): + dn = device.name + assert dn + assert type(dn) is str + + +def check_driver_version(device): + dv = device.driver_version + assert dv + assert type(dv) is str + + +def check_vendor(device): + ve = device.vendor + assert ve or device.is_host + assert type(ve) is str + + +def check_default_selector_score(device): + sc = device.default_selector_score + assert type(sc) is int + assert sc > 0 + + +def check_backend(device): + be = device.backend + assert type(be) is dpctl.backend_type + + +def check_device_type(device): + dt = device.device_type + assert type(dt) is dpctl.device_type + + +def check_global_mem_cache_type(device): + gmc_ty = device.global_mem_cache_type + assert type(gmc_ty) is dpctl.global_mem_cache_type + + +def check_global_mem_cache_size(device): + gmc_sz = device.global_mem_cache_size + assert type(gmc_sz) is int + assert gmc_sz + + +def check_global_mem_cache_line_size(device): + gmc_sz = device.global_mem_cache_line_size + assert type(gmc_sz) is int + assert gmc_sz + + list_of_checks = [ - check_get_max_compute_units, - check_get_max_work_item_dims, - check_get_max_work_item_sizes1d, - check_get_max_work_item_sizes2d, - check_get_max_work_item_sizes3d, - check_get_max_work_item_sizes, - check_get_max_work_group_size, - check_get_max_num_sub_groups, + check_max_compute_units, + check_max_work_item_dims, + check_max_work_item_sizes1d, + check_max_work_item_sizes2d, + check_max_work_item_sizes3d, + check_max_work_item_sizes, + check_max_work_group_size, + check_max_num_sub_groups, check_is_accelerator, check_is_cpu, check_is_gpu, check_is_host, - check_get_sub_group_independent_forward_progress, - check_get_preferred_vector_width_char, - check_get_preferred_vector_width_short, - check_get_preferred_vector_width_int, - check_get_preferred_vector_width_long, - check_get_preferred_vector_width_float, - check_get_preferred_vector_width_double, - check_get_preferred_vector_width_half, + check_sub_group_independent_forward_progress, + check_preferred_vector_width_char, + check_preferred_vector_width_short, + check_preferred_vector_width_int, + check_preferred_vector_width_long, + check_preferred_vector_width_float, + check_preferred_vector_width_double, + check_preferred_vector_width_half, check_has_aspect_host, check_has_aspect_cpu, check_has_aspect_gpu, @@ -555,13 +621,13 @@ def check_platform(device): check_has_aspect_usm_atomic_host_allocations, check_has_aspect_usm_atomic_shared_allocations, check_has_aspect_host_debuggable, - check_get_max_read_image_args, - check_get_max_write_image_args, - check_get_image_2d_max_width, - check_get_image_2d_max_height, - check_get_image_3d_max_width, - check_get_image_3d_max_height, - check_get_image_3d_max_depth, + check_max_read_image_args, + check_max_write_image_args, + check_image_2d_max_width, + check_image_2d_max_height, + check_image_3d_max_width, + check_image_3d_max_height, + check_image_3d_max_depth, check_create_sub_devices_equally, check_create_sub_devices_by_counts, check_create_sub_devices_by_affinity_not_applicable, @@ -573,10 +639,21 @@ def check_platform(device): check_create_sub_devices_by_affinity_next_partitionable, check_print_device_info, check_repr, - check_get_global_mem_size, - check_get_local_mem_size, + check_global_mem_size, + check_local_mem_size, check_profiling_timer_resolution, check_platform, + check_parent_device, + check_filter_string, + check_vendor, + check_driver_version, + check_name, + check_default_selector_score, + check_backend, + check_device_type, + check_global_mem_cache_type, + check_global_mem_cache_size, + check_global_mem_cache_line_size, ] diff --git a/dpctl/tests/test_sycl_device.py b/dpctl/tests/test_sycl_device.py index c938aec466..4a6e42f8a4 100644 --- a/dpctl/tests/test_sycl_device.py +++ b/dpctl/tests/test_sycl_device.py @@ -129,6 +129,13 @@ def test_hashing_of_device(): assert device_dict +def test_equal(): + d1 = dpctl.SyclDevice() + d2 = dpctl.SyclDevice() + assert d1 != Ellipsis + assert d1 == d2 + + list_of_supported_aspects = [ "cpu", "gpu", diff --git a/dpctl/tests/test_sycl_program.py b/dpctl/tests/test_sycl_program.py index bd55e2b4cf..da435704f4 100644 --- a/dpctl/tests/test_sycl_program.py +++ b/dpctl/tests/test_sycl_program.py @@ -48,6 +48,26 @@ def _check_multi_kernel_program(prog): assert type(addKernel.addressof_ref()) is int assert type(axpyKernel.addressof_ref()) is int + for krn in [addKernel, axpyKernel]: + na = krn.num_args + assert na == krn.get_num_args() + wgsz = krn.work_group_size + assert type(wgsz) is int + pwgszm = krn.preferred_work_group_size_multiple + assert type(pwgszm) is int + pmsz = krn.private_mem_size + assert type(pmsz) is int + vmnsg = krn.max_num_sub_groups + assert type(vmnsg) is int + v = krn.max_sub_group_size + assert ( + v == NotImplemented + ), "SyclKernel.max_sub_group_size acquired implementation, fix the test" + cmnsg = krn.compile_num_sub_groups + assert type(cmnsg) is int + cmsgsz = krn.compile_sub_group_size + assert type(cmsgsz) is int + def test_create_program_from_source_ocl(): oclSrc = " \ diff --git a/libsyclinterface/include/dpctl_sycl_device_interface.h b/libsyclinterface/include/dpctl_sycl_device_interface.h index 0762f595e9..20677d5f4e 100644 --- a/libsyclinterface/include/dpctl_sycl_device_interface.h +++ b/libsyclinterface/include/dpctl_sycl_device_interface.h @@ -492,9 +492,6 @@ DPCTL_API __dpctl_give DPCTLDeviceVectorRef DPCTLDevice_CreateSubDevicesByAffinity( __dpctl_keep const DPCTLSyclDeviceRef DRef, DPCTLPartitionAffinityDomainType PartAffDomTy); - -DPCTL_C_EXTERN_C_END - /*! * @brief Wrapper over * device.get_info. @@ -631,3 +628,38 @@ size_t DPCTLDevice_Hash(__dpctl_keep const DPCTLSyclDeviceRef DRef); DPCTL_API size_t DPCTLDevice_GetProfilingTimerResolution( __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the size of global memory cache line in bytes as uint32_t. + */ +DPCTL_API +uint32_t DPCTLDevice_GetGlobalMemCacheLineSize( + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the size of global memory cache in bytes as uint64_t. + */ +DPCTL_API +uint64_t +DPCTLDevice_GetGlobalMemCacheSize(__dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Wrapper over + * device.get_info + * + * @param DRef Opaque pointer to a sycl::device + * @return Returns the type of global memory cache supported. + */ +DPCTL_API +DPCTLGlobalMemCacheType +DPCTLDevice_GetGlobalMemCacheType(__dpctl_keep const DPCTLSyclDeviceRef DRef); + +DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/dpctl_sycl_enum_types.h b/libsyclinterface/include/dpctl_sycl_enum_types.h index 6265850fbd..1ac169ce2c 100644 --- a/libsyclinterface/include/dpctl_sycl_enum_types.h +++ b/libsyclinterface/include/dpctl_sycl_enum_types.h @@ -161,4 +161,12 @@ typedef enum DPCTL_COMPLETE } DPCTLSyclEventStatusType; +typedef enum +{ + DPCTL_MEM_CACHE_TYPE_INDETERMINATE, + DPCTL_MEM_CACHE_TYPE_NONE, + DPCTL_MEM_CACHE_TYPE_READ_ONLY, + DPCTL_MEM_CACHE_TYPE_READ_WRITE +} DPCTLGlobalMemCacheType; + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/dpctl_sycl_kernel_interface.h b/libsyclinterface/include/dpctl_sycl_kernel_interface.h index 3bc0470bd7..6bc0ae83c8 100644 --- a/libsyclinterface/include/dpctl_sycl_kernel_interface.h +++ b/libsyclinterface/include/dpctl_sycl_kernel_interface.h @@ -39,11 +39,12 @@ DPCTL_C_EXTERN_C_BEGIN */ /*! - * @brief Returns the number of arguments for the OpenCL kernel. + * @brief Returns the number of arguments for the sycl + * interoperability kernel. * - * @param KRef DPCTLSyclKernelRef pointer to an OpenCL + * @param KRef DPCTLSyclKernelRef pointer to a SYCL * interoperability kernel. - * @return Returns the number of arguments for the OpenCL interoperability + * @return Returns the number of arguments for the interoperability * kernel. * @ingroup KernelInterface */ @@ -51,13 +52,112 @@ DPCTL_API size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef); /*! - * @brief Deletes the DPCTLSyclKernelRef after casting it to a sycl::kernel. + * @brief Deletes the DPCTLSyclKernelRef after casting it to a + * ``sycl::kernel``. * - * @param KRef DPCTLSyclKernelRef pointer to an OpenCL + * @param KRef DPCTLSyclKernelRef pointer to a SYCL * interoperability kernel. * @ingroup KernelInterface */ DPCTL_API void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef); +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to a SYCL + * interoperability kernel. + * @return Returns the maximum number of work-items in a work-group + * that can be used to execute a kernel on the device it was + * built for. + * @ingroup KernelInterface + */ +DPCTL_API +size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef); + +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to a SYCL + * interoperability kernel. + * @return Returns a value, of which work-group size is preferred to be a + * multiple, for executing a kernel on the device it was built for. + * @ingroup KernelInterface + */ +DPCTL_API +size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple( + __dpctl_keep const DPCTLSyclKernelRef KRef); + +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to a SYCL + * interoperability kernel. + * @return Returns the minimum amount of private memory, in bytes, + * used by each work-item in the kernel. + * @ingroup KernelInterface + */ +DPCTL_API +size_t +DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef); + +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to an SYCL + * interoperability kernel. + * @return Returns the maximum number of sub-groups for this kernel. + * @ingroup KernelInterface + */ +DPCTL_API +uint32_t +DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef); + +#if 0 +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to an SYCL + * interoperability kernel. + * @return Returns the maximum sub-group size for this kernel. + * @ingroup KernelInterface + */ +DPCTL_API +uint32_t +DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef); +#endif + +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to an SYCL + * interoperability kernel. + * @return Returns the number of sub-groups specified by the kernel, + * or 0 (if not specified). + * @ingroup KernelInterface + */ +DPCTL_API +uint32_t +DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef); + +/*! + * !brief Wrapper around + * `kernel::get_info()`. + * + * @param KRef DPCTLSyclKernelRef pointer to an SYCL + * interoperability kernel. + * @return Returns the required sub-group size specified by this kernel, + * or 0 (if not specified). + * @ingroup KernelInterface + */ +DPCTL_API +uint32_t +DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/dpctl_sycl_types.h b/libsyclinterface/include/dpctl_sycl_types.h index b00cab4702..2945fcef1b 100644 --- a/libsyclinterface/include/dpctl_sycl_types.h +++ b/libsyclinterface/include/dpctl_sycl_types.h @@ -60,7 +60,8 @@ typedef struct DPCTLOpaqueSyclEvent *DPCTLSyclEventRef; typedef struct DPCTLOpaqueSyclKernel *DPCTLSyclKernelRef; /*! - * @brief Opaque pointer to a ``sycl::kernel_bundle`` + * @brief Opaque pointer to a + * ``sycl::kernel_bundle`` * */ typedef struct DPCTLOpaqueSyclKernelBundle *DPCTLSyclKernelBundleRef; diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index c65f9ac38b..7494367924 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -691,3 +691,54 @@ size_t DPCTLDevice_GetProfilingTimerResolution( return 0; } } + +uint32_t DPCTLDevice_GetGlobalMemCacheLineSize( + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + if (DRef) { + auto D = unwrap(DRef); + return D->get_info(); + } + else { + error_handler("Argument DRef is null", __FILE__, __func__, __LINE__); + return 0; + } +} + +uint64_t +DPCTLDevice_GetGlobalMemCacheSize(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + if (DRef) { + auto D = unwrap(DRef); + return D->get_info(); + } + else { + error_handler("Argument DRef is null", __FILE__, __func__, __LINE__); + return 0; + } +} + +DPCTLGlobalMemCacheType +DPCTLDevice_GetGlobalMemCacheType(__dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + if (DRef) { + auto D = unwrap(DRef); + auto mem_type = D->get_info(); + switch (mem_type) { + case info::global_mem_cache_type::none: + return DPCTL_MEM_CACHE_TYPE_NONE; + case info::global_mem_cache_type::read_only: + return DPCTL_MEM_CACHE_TYPE_READ_ONLY; + case info::global_mem_cache_type::read_write: + return DPCTL_MEM_CACHE_TYPE_READ_WRITE; + } + // If execution reaches here unrecognized mem_type was returned. Check + // values in the enumeration `info::global_mem_cache_type` in SYCL specs + assert(false); + return DPCTL_MEM_CACHE_TYPE_INDETERMINATE; + } + else { + error_handler("Argument DRef is null", __FILE__, __func__, __LINE__); + return DPCTL_MEM_CACHE_TYPE_INDETERMINATE; + } +} diff --git a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp index e2fc31bba1..236546fb8d 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_interface.cpp @@ -29,6 +29,7 @@ #include "dpctl_error_handlers.h" #include "dpctl_string_utils.hpp" #include /* Sycl headers */ +#include using namespace cl::sycl; @@ -39,21 +40,177 @@ DEFINE_SIMPLE_CONVERSION_FUNCTIONS(kernel, DPCTLSyclKernelRef) } /* end of anonymous namespace */ -size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef Kernel) +size_t DPCTLKernel_GetNumArgs(__dpctl_keep const DPCTLSyclKernelRef KRef) { - if (!Kernel) { + if (!KRef) { error_handler("Cannot get the number of arguments from " "DPCTLSyclKernelRef as input is a nullptr.", __FILE__, __func__, __LINE__); return -1; } - auto SyclKernel = unwrap(Kernel); - auto num_args = SyclKernel->get_info(); - return (size_t)num_args; + auto sycl_kernel = unwrap(KRef); + auto num_args = sycl_kernel->get_info(); + return static_cast(num_args); } -void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef Kernel) +void DPCTLKernel_Delete(__dpctl_take DPCTLSyclKernelRef KRef) { - delete unwrap(Kernel); + delete unwrap(KRef); +} + +size_t DPCTLKernel_GetWorkGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = sycl_kern->get_info( + devs[0]); + return static_cast(v); +} + +size_t DPCTLKernel_GetPreferredWorkGroupSizeMultiple( + __dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = sycl_kern->get_info< + info::kernel_device_specific::preferred_work_group_size_multiple>( + devs[0]); + return static_cast(v); +} + +size_t DPCTLKernel_GetPrivateMemSize(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = + sycl_kern->get_info( + devs[0]); + return static_cast(v); +} + +uint32_t +DPCTLKernel_GetMaxNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = + sycl_kern->get_info( + devs[0]); + return static_cast(v); +} + +#if 0 +// commented out due to bug in DPC++ runtime, get_info for max_sub_group_size +// exported by libsycl has different, not SPEC-compliant signature +uint32_t +DPCTLKernel_GetMaxSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = sycl_kern + ->get_info(devs[0]); + return v; +} +#endif + +uint32_t +DPCTLKernel_GetCompileNumSubGroups(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = + sycl_kern + ->get_info( + devs[0]); + return static_cast(v); +} + +uint32_t +DPCTLKernel_GetCompileSubGroupSize(__dpctl_keep const DPCTLSyclKernelRef KRef) +{ + if (!KRef) { + error_handler("Input DPCTKSyclKernelRef is nullptr.", __FILE__, + __func__, __LINE__); + return 0; + } + + auto sycl_kern = unwrap(KRef); + auto devs = sycl_kern->get_kernel_bundle().get_devices(); + if (devs.empty()) { + error_handler("Input DPCTKSyclKernelRef has no associated device.", + __FILE__, __func__, __LINE__); + return 0; + } + auto v = + sycl_kern + ->get_info( + devs[0]); + return static_cast(v); } diff --git a/libsyclinterface/tests/test_sycl_device_interface.cpp b/libsyclinterface/tests/test_sycl_device_interface.cpp index 71be76fe80..ba3f7fb245 100644 --- a/libsyclinterface/tests/test_sycl_device_interface.cpp +++ b/libsyclinterface/tests/test_sycl_device_interface.cpp @@ -407,6 +407,30 @@ TEST_P(TestDPCTLSyclDeviceInterface, ChkGetProfilingTimerResolution) EXPECT_TRUE(res != 0); } +TEST_P(TestDPCTLSyclDeviceInterface, ChkGetGlobalMemCacheSize) +{ + uint64_t res = 0; + EXPECT_NO_FATAL_FAILURE(res = DPCTLDevice_GetGlobalMemCacheSize(DRef)); + EXPECT_TRUE(res != 0); +} + +TEST_P(TestDPCTLSyclDeviceInterface, ChkGetGlobalMemCacheLineSize) +{ + uint32_t res = 0; + EXPECT_NO_FATAL_FAILURE(res = DPCTLDevice_GetGlobalMemCacheLineSize(DRef)); + EXPECT_TRUE(res != 0); +} + +TEST_P(TestDPCTLSyclDeviceInterface, ChkGetGlobalMemCacheType) +{ + DPCTLGlobalMemCacheType res = DPCTL_MEM_CACHE_TYPE_INDETERMINATE; + EXPECT_NO_FATAL_FAILURE(res = DPCTLDevice_GetGlobalMemCacheType(DRef)); + EXPECT_TRUE(res != DPCTL_MEM_CACHE_TYPE_INDETERMINATE); + EXPECT_TRUE((res == DPCTL_MEM_CACHE_TYPE_NONE || + res == DPCTL_MEM_CACHE_TYPE_READ_ONLY || + res == DPCTL_MEM_CACHE_TYPE_READ_WRITE)); +} + INSTANTIATE_TEST_SUITE_P(DPCTLDeviceFns, TestDPCTLSyclDeviceInterface, ::testing::Values("opencl", @@ -713,3 +737,25 @@ TEST_F(TestDPCTLSyclDeviceNullArgs, ChkGetProfilingTimerResolution) res = DPCTLDevice_GetProfilingTimerResolution(Null_DRef)); ASSERT_TRUE(res == 0); } + +TEST_F(TestDPCTLSyclDeviceNullArgs, ChkGetGlobalMemCacheSize) +{ + uint64_t res = 1; + EXPECT_NO_FATAL_FAILURE(res = DPCTLDevice_GetGlobalMemCacheSize(Null_DRef)); + ASSERT_TRUE(res == 0); +} + +TEST_F(TestDPCTLSyclDeviceNullArgs, ChkGetGlobalMemCacheLineSize) +{ + uint32_t res = 1; + EXPECT_NO_FATAL_FAILURE( + res = DPCTLDevice_GetGlobalMemCacheLineSize(Null_DRef)); + ASSERT_TRUE(res == 0); +} + +TEST_F(TestDPCTLSyclDeviceNullArgs, ChkGetGlobalMemCacheType) +{ + DPCTLGlobalMemCacheType res = DPCTL_MEM_CACHE_TYPE_NONE; + EXPECT_NO_FATAL_FAILURE(res = DPCTLDevice_GetGlobalMemCacheType(Null_DRef)); + ASSERT_TRUE(res == DPCTL_MEM_CACHE_TYPE_INDETERMINATE); +} diff --git a/libsyclinterface/tests/test_sycl_kernel_interface.cpp b/libsyclinterface/tests/test_sycl_kernel_interface.cpp index 89cc586aab..016a3ccd3e 100644 --- a/libsyclinterface/tests/test_sycl_kernel_interface.cpp +++ b/libsyclinterface/tests/test_sycl_kernel_interface.cpp @@ -57,17 +57,33 @@ struct TestDPCTLSyclKernelInterface const char *CompileOpts = "-cl-fast-relaxed-math"; DPCTLSyclDeviceSelectorRef DSRef = nullptr; DPCTLSyclDeviceRef DRef = nullptr; + DPCTLSyclQueueRef QRef = nullptr; + DPCTLSyclContextRef CtxRef = nullptr; + DPCTLSyclKernelBundleRef KBRef = nullptr; + DPCTLSyclKernelRef AddKRef = nullptr; + DPCTLSyclKernelRef AxpyKRef = nullptr; TestDPCTLSyclKernelInterface() { DSRef = DPCTLFilterSelector_Create(GetParam()); DRef = DPCTLDevice_CreateFromSelector(DSRef); + QRef = DPCTLQueue_CreateForDevice(DRef, nullptr, 0); + CtxRef = DPCTLQueue_GetContext(QRef); + KBRef = DPCTLKernelBundle_CreateFromOCLSource( + CtxRef, DRef, CLProgramStr, CompileOpts); + AddKRef = DPCTLKernelBundle_GetKernel(KBRef, "add"); + AxpyKRef = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); } ~TestDPCTLSyclKernelInterface() { DPCTLDeviceSelector_Delete(DSRef); DPCTLDevice_Delete(DRef); + DPCTLQueue_Delete(QRef); + DPCTLContext_Delete(CtxRef); + DPCTLKernelBundle_Delete(KBRef); + DPCTLKernel_Delete(AddKRef); + DPCTLKernel_Delete(AxpyKRef); } void SetUp() @@ -83,30 +99,170 @@ struct TestDPCTLSyclKernelInterface TEST_P(TestDPCTLSyclKernelInterface, CheckGetNumArgs) { - auto QueueRef = DPCTLQueue_CreateForDevice(DRef, nullptr, 0); - auto CtxRef = DPCTLQueue_GetContext(QueueRef); - auto KBRef = DPCTLKernelBundle_CreateFromOCLSource( - CtxRef, DRef, CLProgramStr, CompileOpts); - auto AddKernel = DPCTLKernelBundle_GetKernel(KBRef, "add"); - auto AxpyKernel = DPCTLKernelBundle_GetKernel(KBRef, "axpy"); - ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKernel), 3ul); - ASSERT_EQ(DPCTLKernel_GetNumArgs(AxpyKernel), 4ul); + ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKRef), 3ul); + ASSERT_EQ(DPCTLKernel_GetNumArgs(AxpyKRef), 4ul); +} + +TEST_P(TestDPCTLSyclKernelInterface, CheckGetWorkGroupSize) +{ + + size_t add_wgsz = 0, axpy_wgsz = 0; + EXPECT_NO_FATAL_FAILURE(add_wgsz = DPCTLKernel_GetWorkGroupSize(AddKRef)); + EXPECT_NO_FATAL_FAILURE(axpy_wgsz = DPCTLKernel_GetWorkGroupSize(AxpyKRef)); + + ASSERT_TRUE(add_wgsz != 0); + ASSERT_TRUE(axpy_wgsz != 0); +} + +TEST_P(TestDPCTLSyclKernelInterface, CheckGetPreferredWorkGroupSizeMultiple) +{ + + size_t add_wgsz_m = 0, axpy_wgsz_m = 0; + EXPECT_NO_FATAL_FAILURE( + add_wgsz_m = DPCTLKernel_GetPreferredWorkGroupSizeMultiple(AddKRef)); + EXPECT_NO_FATAL_FAILURE( + axpy_wgsz_m = DPCTLKernel_GetPreferredWorkGroupSizeMultiple(AxpyKRef)); - DPCTLQueue_Delete(QueueRef); - DPCTLContext_Delete(CtxRef); - DPCTLKernelBundle_Delete(KBRef); - DPCTLKernel_Delete(AddKernel); - DPCTLKernel_Delete(AxpyKernel); + ASSERT_TRUE(add_wgsz_m != 0); + ASSERT_TRUE(axpy_wgsz_m != 0); } -TEST_P(TestDPCTLSyclKernelInterface, CheckNullPtrArg) +TEST_P(TestDPCTLSyclKernelInterface, CheckGetPrivateMemSize) { - DPCTLSyclKernelRef AddKernel = nullptr; - ASSERT_EQ(DPCTLKernel_GetNumArgs(AddKernel), -1); + size_t add_private_mem_sz = 0, axpy_private_mem_sz = 0; + EXPECT_NO_FATAL_FAILURE(add_private_mem_sz = + DPCTLKernel_GetPrivateMemSize(AddKRef)); + EXPECT_NO_FATAL_FAILURE(axpy_private_mem_sz = + DPCTLKernel_GetPrivateMemSize(AxpyKRef)); + + if (DPCTLDevice_IsGPU(DRef)) { + ASSERT_TRUE(add_private_mem_sz != 0); + ASSERT_TRUE(axpy_private_mem_sz != 0); + } + else { + ASSERT_TRUE(add_private_mem_sz >= 0); + ASSERT_TRUE(axpy_private_mem_sz >= 0); + } +} + +TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxNumSubGroups) +{ + + uint32_t add_mnsg = 0, axpy_mnsg = 0; + EXPECT_NO_FATAL_FAILURE(add_mnsg = DPCTLKernel_GetMaxNumSubGroups(AddKRef)); + EXPECT_NO_FATAL_FAILURE(axpy_mnsg = + DPCTLKernel_GetMaxNumSubGroups(AxpyKRef)); + + ASSERT_TRUE(add_mnsg != 0); + ASSERT_TRUE(axpy_mnsg != 0); +} + +/* +TEST_P(TestDPCTLSyclKernelInterface, CheckGetMaxSubGroupSize) +{ + + uint32_t add_msg_sz = 0, axpy_msg_sz = 0; + EXPECT_NO_FATAL_FAILURE(add_msg_sz = + DPCTLKernel_GetMaxSubGroupSize(AddKRef)); + EXPECT_NO_FATAL_FAILURE(axpy_msg_sz = + DPCTLKernel_GetMaxSubGroupSize(AxpyKRef)); + + ASSERT_TRUE(add_msg_sz != 0); + ASSERT_TRUE(axpy_msg_sz != 0); +} +*/ + +TEST_P(TestDPCTLSyclKernelInterface, CheckGetCompileNumSubGroups) +{ + + uint32_t add_cnsg = 0, axpy_cnsg = 0; + EXPECT_NO_FATAL_FAILURE(add_cnsg = + DPCTLKernel_GetCompileNumSubGroups(AddKRef)); + EXPECT_NO_FATAL_FAILURE(axpy_cnsg = + DPCTLKernel_GetCompileNumSubGroups(AxpyKRef)); + + EXPECT_TRUE(add_cnsg >= 0); + EXPECT_TRUE(axpy_cnsg >= 0); +} + +TEST_P(TestDPCTLSyclKernelInterface, CheckGetCompileSubGroupSize) +{ + + uint32_t add_csg_sz = 0, axpy_csg_sz = 0; + EXPECT_NO_FATAL_FAILURE(add_csg_sz = + DPCTLKernel_GetCompileSubGroupSize(AddKRef)); + EXPECT_NO_FATAL_FAILURE(axpy_csg_sz = + DPCTLKernel_GetCompileSubGroupSize(AxpyKRef)); + EXPECT_TRUE(add_csg_sz >= 0); + EXPECT_TRUE(axpy_csg_sz >= 0); } INSTANTIATE_TEST_SUITE_P(TestKernelInterfaceFunctions, TestDPCTLSyclKernelInterface, ::testing::Values("opencl:gpu:0", "opencl:cpu:0")); + +struct TestDPCTLSyclKernelNullArgs : public ::testing::Test +{ + DPCTLSyclKernelRef Null_KRef; + TestDPCTLSyclKernelNullArgs() : Null_KRef(nullptr) {} + ~TestDPCTLSyclKernelNullArgs() {} +}; + +TEST_F(TestDPCTLSyclKernelNullArgs, CheckNumArgsNullKRef) +{ + ASSERT_EQ(DPCTLKernel_GetNumArgs(Null_KRef), -1); +} + +TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetWorkGroupsSizeNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetWorkGroupSize(NullKRef), 0); +} + +TEST_F(TestDPCTLSyclKernelNullArgs, + CheckGetPreferredWorkGroupsSizeMultipleNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetPreferredWorkGroupSizeMultiple(NullKRef), 0); +} + +TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetPrivateMemSizeNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetPrivateMemSize(NullKRef), 0); +} + +TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetMaxNumSubGroupsNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetMaxNumSubGroups(NullKRef), 0); +} + +/* +TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetMaxSubGroupSizeNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetMaxSubGroupSize(NullKRef), 0); +} +*/ + +TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetCompileNumSubGroupsNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetCompileNumSubGroups(NullKRef), 0); +} + +TEST_F(TestDPCTLSyclKernelNullArgs, CheckGetCompileSubGroupSizeNullKRef) +{ + DPCTLSyclKernelRef NullKRef = nullptr; + + ASSERT_EQ(DPCTLKernel_GetCompileSubGroupSize(NullKRef), 0); +}