diff --git a/docs/doc_sources/urls.json b/docs/doc_sources/urls.json index 647bbaea99..34801db920 100644 --- a/docs/doc_sources/urls.json +++ b/docs/doc_sources/urls.json @@ -5,6 +5,7 @@ "oneapi_filter_selection": "https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_filter_selector.asciidoc", "oneapi_default_context": "https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_default_context.asciidoc", "oneapi_enqueue_barrier": "https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_enqueue_barrier.asciidoc", + "oneapi_peer_access": "https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_peer_access.asciidoc", "sycl_aspects": "https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#table.device.aspect", "sycl_context": "https://sycl.readthedocs.io/en/latest/iface/context.html", "sycl_device": "https://sycl.readthedocs.io/en/latest/iface/device.html", diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index ca9e9ccb9f..93d9b5ef97 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -112,6 +112,10 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h": _L1_cache "L1_cache", _next_partitionable "next_partitionable", + ctypedef enum _peer_access "DPCTLPeerAccessType": + _access_supported "access_supported", + _atomics_supported "atomics_supported", + ctypedef enum _event_status_type "DPCTLSyclEventStatusType": _UNKNOWN_STATUS "DPCTL_UNKNOWN_STATUS" _SUBMITTED "DPCTL_SUBMITTED" @@ -278,7 +282,14 @@ cdef extern from "syclinterface/dpctl_sycl_device_interface.h": cdef DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices( const DPCTLSyclDeviceRef DRef ) + cdef bool DPCTLDevice_CanAccessPeer(const DPCTLSyclDeviceRef DRef, + const DPCTLSyclDeviceRef PDRef, + _peer_access PT) + cdef void DPCTLDevice_EnablePeerAccess(const DPCTLSyclDeviceRef DRef, + const DPCTLSyclDeviceRef PDRef) + cdef void DPCTLDevice_DisablePeerAccess(const DPCTLSyclDeviceRef DRef, + const DPCTLSyclDeviceRef PDRef) cdef extern from "syclinterface/dpctl_sycl_device_manager.h": cdef DPCTLDeviceVectorRef DPCTLDeviceVector_CreateFromArray( diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index 5b43ffed1a..2f8a1668ad 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -25,12 +25,15 @@ from ._backend cimport ( # noqa: E211 DPCTLCString_Delete, DPCTLDefaultSelector_Create, DPCTLDevice_AreEq, + DPCTLDevice_CanAccessPeer, DPCTLDevice_Copy, DPCTLDevice_CreateFromSelector, DPCTLDevice_CreateSubDevicesByAffinity, DPCTLDevice_CreateSubDevicesByCounts, DPCTLDevice_CreateSubDevicesEqually, DPCTLDevice_Delete, + DPCTLDevice_DisablePeerAccess, + DPCTLDevice_EnablePeerAccess, DPCTLDevice_GetBackend, DPCTLDevice_GetComponentDevices, DPCTLDevice_GetCompositeDevice, @@ -103,6 +106,7 @@ from ._backend cimport ( # noqa: E211 _device_type, _global_mem_cache_type, _partition_affinity_domain_type, + _peer_access, ) from .enum_types import backend_type, device_type, global_mem_cache_type @@ -213,6 +217,65 @@ cdef void _init_helper(_SyclDevice device, DPCTLSyclDeviceRef DRef) except *: raise RuntimeError("Descriptor 'max_work_item_sizes3d' not available") +cdef inline bint _check_peer_access(SyclDevice dev, SyclDevice peer) except *: + """ + Check peer access ahead of time to avoid errors from unified runtime or + compiler implementation. + """ + cdef list _peer_access_backends = [ + _backend_type._CUDA, + _backend_type._HIP, + _backend_type._LEVEL_ZERO + ] + cdef _backend_type BTy1 = DPCTLDevice_GetBackend(dev._device_ref) + cdef _backend_type BTy2 = DPCTLDevice_GetBackend(peer.get_device_ref()) + if ( + BTy1 == BTy2 and + BTy1 in _peer_access_backends and + BTy2 in _peer_access_backends and + dev != peer + ): + return True + return False + + +cdef inline void _raise_invalid_peer_access( + SyclDevice dev, + SyclDevice peer, +) except *: + """ + Check peer access ahead of time and raise errors for invalid cases. + """ + cdef list _peer_access_backends = [ + _backend_type._CUDA, + _backend_type._HIP, + _backend_type._LEVEL_ZERO + ] + cdef _backend_type BTy1 = DPCTLDevice_GetBackend(dev._device_ref) + cdef _backend_type BTy2 = DPCTLDevice_GetBackend(peer.get_device_ref()) + if (BTy1 != BTy2): + raise ValueError( + f"Device with backend {_backend_type_to_filter_string_part(BTy1)} " + "cannot peer access device with backend " + f"{_backend_type_to_filter_string_part(BTy2)}" + ) + if (BTy1 not in _peer_access_backends): + raise ValueError( + "Peer access not supported for backend " + f"{_backend_type_to_filter_string_part(BTy1)}" + ) + if (BTy2 not in _peer_access_backends): + raise ValueError( + "Peer access not supported for backend " + f"{_backend_type_to_filter_string_part(BTy2)}" + ) + if (dev == peer): + raise ValueError( + "Peer access cannot be enabled between a device and itself" + ) + return + + @functools.lru_cache(maxsize=None) def _cached_filter_string(d : SyclDevice): """ @@ -220,7 +283,7 @@ def _cached_filter_string(d : SyclDevice): and cached with `functools.cache`. Args: - d (dpctl.SyclDevice): + d (:class:`dpctl.SyclDevice`): A device for which to compute the filter string. Returns: out(str): @@ -1792,6 +1855,150 @@ cdef class SyclDevice(_SyclDevice): raise ValueError("Internal error: NULL device vector encountered") return _get_devices(cDVRef) + def can_access_peer(self, peer, value="access_supported"): + """ Returns ``True`` if this device (``self``) can enable peer access + to USM device memory on ``peer``, ``False`` otherwise. + + If peer access is supported, it may be enabled by calling + :meth:`.enable_peer_access`. + + For details, see + :oneapi_peer_access:`DPC++ peer access SYCL extension <>`. + + Args: + peer (:class:`dpctl.SyclDevice`): + The :class:`dpctl.SyclDevice` instance to check for peer access + by this device. + value (str, optional): + Specifies the kind of peer access being queried. + + The supported values are + + - ``"access_supported"`` + Returns ``True`` if it is possible for this device to + enable peer access to USM device memory on ``peer``. + + - ``"atomics_supported"`` + Returns ``True`` if it is possible for this device to + concurrently access and atomically modify USM device + memory on ``peer`` when enabled. Atomics must have + ``memory_scope::system`` when modifying memory on a peer + device. + + If ``False`` is returned, these operations result in + undefined behavior. + + Default: ``"access_supported"`` + + Returns: + bool: + ``True`` if the kind of peer access specified by ``value`` is + supported between this device and ``peer``, otherwise ``False``. + + Raises: + TypeError: + If ``peer`` is not :class:`dpctl.SyclDevice`. + """ + cdef SyclDevice p_dev + + if not isinstance(value, str): + raise TypeError( + f"Expected `value` to be of type str, got {type(value)}" + ) + if value == "access_supported": + access_type = _peer_access._access_supported + elif value == "atomics_supported": + access_type = _peer_access._atomics_supported + else: + raise ValueError( + "`value` must be 'access_supported' or 'atomics_supported', " + f"got {value}" + ) + if not isinstance(peer, SyclDevice): + raise TypeError( + "peer device must be a `dpctl.SyclDevice`, got " + f"{type(peer)}" + ) + p_dev = peer + if _check_peer_access(self, p_dev): + return DPCTLDevice_CanAccessPeer( + self._device_ref, + p_dev.get_device_ref(), + access_type + ) + return False + + def enable_peer_access(self, peer): + """ Enables this device (``self``) to access USM device allocations + located on ``peer``. + + Peer access may be disabled by calling :meth:`.disable_peer_access`. + + For details, see + :oneapi_peer_access:`DPC++ peer access SYCL extension <>`. + + Args: + peer (:class:`dpctl.SyclDevice`): + The :class:`dpctl.SyclDevice` instance to enable peer access + to. + + Raises: + TypeError: + If ``peer`` is not :class:`dpctl.SyclDevice`. + ValueError: + If the backend associated with this device or ``peer`` does not + support peer access. + """ + cdef SyclDevice p_dev + + if not isinstance(peer, SyclDevice): + raise TypeError( + "peer device must be a `dpctl.SyclDevice`, got " + f"{type(peer)}" + ) + p_dev = peer + _raise_invalid_peer_access(self, p_dev) + DPCTLDevice_EnablePeerAccess( + self._device_ref, + p_dev.get_device_ref() + ) + return + + def disable_peer_access(self, peer): + """ Disables peer access to ``peer`` from this device (``self``). + + Peer access may be enabled by calling :meth:`.enable_peer_access`. + + For details, see + :oneapi_peer_access:`DPC++ peer access SYCL extension <>`. + + Args: + peer (:class:`dpctl.SyclDevice`): + The :class:`dpctl.SyclDevice` instance to + disable peer access to. + + Raises: + TypeError: + If ``peer`` is not :class:`dpctl.SyclDevice`. + ValueError: + If the backend associated with this device or ``peer`` does not + support peer access. + """ + cdef SyclDevice p_dev + + if not isinstance(peer, SyclDevice): + raise TypeError( + "peer device must be a `dpctl.SyclDevice`, got " + f"{type(peer)}" + ) + p_dev = peer + _raise_invalid_peer_access(self, p_dev) + DPCTLDevice_DisablePeerAccess( + self._device_ref, + p_dev.get_device_ref() + ) + return + @property def profiling_timer_resolution(self): """ Profiling timer resolution. @@ -1912,7 +2119,7 @@ cdef class SyclDevice(_SyclDevice): same _device_ref as this SyclDevice. Args: - other (dpctl.SyclDevice): + other (:class:`dpctl.SyclDevice`): A :class:`dpctl.SyclDevice` instance to compare against. diff --git a/dpctl/tests/test_sycl_device.py b/dpctl/tests/test_sycl_device.py index 0d8025c060..f3ad9ee478 100644 --- a/dpctl/tests/test_sycl_device.py +++ b/dpctl/tests/test_sycl_device.py @@ -341,3 +341,120 @@ def test_get_component_devices_from_composite(): assert d.has_aspect_is_component # component devices are root devices assert d in devices + + +@pytest.mark.parametrize("platform_name", ["level_zero", "cuda", "hip"]) +def test_can_access_peer(platform_name): + """ + Test checks for peer access. + """ + try: + platform = dpctl.SyclPlatform(platform_name) + except ValueError as e: + pytest.skip(f"{str(e)} {platform_name}") + devices = platform.get_devices() + if len(devices) < 2: + pytest.skip( + f"Platform {platform_name} does not have enough devices to " + "test peer access" + ) + dev0 = devices[0] + dev1 = devices[1] + assert isinstance(dev0.can_access_peer(dev1), bool) + assert isinstance( + dev0.can_access_peer(dev1, value="atomics_supported"), bool + ) + + +@pytest.mark.parametrize("platform_name", ["level_zero", "cuda", "hip"]) +def test_enable_disable_peer_access(platform_name): + """ + Test that peer access can be enabled and disabled. + """ + try: + platform = dpctl.SyclPlatform(platform_name) + except ValueError as e: + pytest.skip(f"{str(e)} {platform_name}") + devices = platform.get_devices() + if len(devices) < 2: + pytest.skip( + f"Platform {platform_name} does not have enough devices to " + "test peer access" + ) + dev0 = devices[0] + dev1 = devices[1] + if dev0.can_access_peer(dev1): + dev0.enable_peer_access(dev1) + dev0.disable_peer_access(dev1) + else: + pytest.skip( + f"Provided {platform_name} devices do not support peer access" + ) + + +@pytest.mark.parametrize( + "method", + [ + "can_access_peer", + "enable_peer_access", + "disable_peer_access", + ], +) +def test_peer_device_arg_validation(method): + """ + Test for validation of arguments to peer access related methods. + """ + try: + dev = dpctl.SyclDevice() + except dpctl.SyclDeviceCreationError: + pytest.skip("No default device available") + bad_dev = dict() + callable = getattr(dev, method) + with pytest.raises(TypeError): + callable(bad_dev) + + +@pytest.mark.parametrize("platform_name", ["level_zero", "cuda", "hip"]) +def test_peer_access_to_self(platform_name): + """ + Validate behavior of a device attempting to enable peer access to itself. + """ + try: + platform = dpctl.SyclPlatform(platform_name) + except ValueError as e: + pytest.skip(f"{str(e)} {platform_name}") + dev = platform.get_devices()[0] + with pytest.raises(ValueError): + dev.enable_peer_access(dev) + with pytest.raises(ValueError): + dev.disable_peer_access(dev) + + +def test_peer_access_value_keyword_validation(): + """ + Validate behavior of `can_access_peer` for invalid `value` keyword. + """ + # we pick an arbitrary platform that supports peer access + platforms = dpctl.get_platforms() + peer_access_backends = [ + dpctl.backend_type.cuda, + dpctl.backend_type.hip, + dpctl.backend_type.hip, + ] + devs = None + for p in platforms: + if p.backend in peer_access_backends: + p_devs = p.get_devices() + if len(p_devs) >= 2: + devs = p_devs + break + if devs is None: + pytest.skip("No platform available with enough devices") + dev0 = devs[0] + dev1 = devs[1] + bad_type = 2 + with pytest.raises(TypeError): + dev0.can_access_peer(dev1, value=bad_type) + bad_value = "wrong" + with pytest.raises(ValueError): + dev0.can_access_peer(dev1, value=bad_value) diff --git a/libsyclinterface/helper/include/dpctl_utils_helper.h b/libsyclinterface/helper/include/dpctl_utils_helper.h index e1cb186d0e..9d9401b807 100644 --- a/libsyclinterface/helper/include/dpctl_utils_helper.h +++ b/libsyclinterface/helper/include/dpctl_utils_helper.h @@ -179,6 +179,33 @@ DPCTL_API DPCTLPartitionAffinityDomainType DPCTL_SyclPartitionAffinityDomainToDPCTLType( sycl::info::partition_affinity_domain PartitionAffinityDomain); +/*! + * @brief Converts a DPCTLPeerAccessType enum value to its corresponding + * sycl::ext::oneapi::peer_access enum value. + * + * @param PeerAccessTy A DPCTLPeerAccessType enum value + * @return A sycl::ext::oneapi::peer_access enum value for the input + * DPCTLPeerAccessType enum value. + * @throws runtime_error + */ +DPCTL_API +sycl::ext::oneapi::peer_access +DPCTL_DPCTLPeerAccessTypeToSycl(DPCTLPeerAccessType PeerAccessTy); + +/*! + * @brief Converts a sycl::ext::oneapi::peer_access enum value to + * corresponding DPCTLPeerAccessType enum value. + * + * @param PeerAccess sycl::ext::oneapi::peer_access to be + * converted to DPCTLPeerAccessType enum. + * @return A DPCTLPeerAccessType enum value for the input + * sycl::ext::oneapi::peer_access enum value. + * @throws runtime_error + */ +DPCTL_API +DPCTLPeerAccessType +DPCTL_SyclPeerAccessToDPCTLType(sycl::ext::oneapi::peer_access PeerAccess); + /*! * @brief Gives the index of the given device with respective to all the other * devices of the same type in the device's platform. diff --git a/libsyclinterface/helper/source/dpctl_utils_helper.cpp b/libsyclinterface/helper/source/dpctl_utils_helper.cpp index 88c1385a26..866e7e828e 100644 --- a/libsyclinterface/helper/source/dpctl_utils_helper.cpp +++ b/libsyclinterface/helper/source/dpctl_utils_helper.cpp @@ -452,6 +452,32 @@ DPCTLPartitionAffinityDomainType DPCTL_SyclPartitionAffinityDomainToDPCTLType( } } +ext::oneapi::peer_access +DPCTL_DPCTLPeerAccessTypeToSycl(DPCTLPeerAccessType PeerAccessTy) +{ + switch (PeerAccessTy) { + case DPCTLPeerAccessType::access_supported: + return ext::oneapi::peer_access::access_supported; + case DPCTLPeerAccessType::atomics_supported: + return ext::oneapi::peer_access::atomics_supported; + default: + throw std::runtime_error("Unsupported peer_access type"); + } +} + +DPCTLPeerAccessType +DPCTL_SyclPeerAccessToDPCTLType(ext::oneapi::peer_access PeerAccess) +{ + switch (PeerAccess) { + case ext::oneapi::peer_access::access_supported: + return DPCTLPeerAccessType::access_supported; + case ext::oneapi::peer_access::atomics_supported: + return DPCTLPeerAccessType::atomics_supported; + default: + throw std::runtime_error("Unsupported peer_access type"); + } +} + int64_t DPCTL_GetRelativeDeviceId(const device &Device) { auto relid = -1; diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h index 6fddb2967f..72b0261e1f 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_interface.h @@ -792,4 +792,40 @@ DPCTL_API __dpctl_give DPCTLDeviceVectorRef DPCTLDevice_GetComponentDevices(__dpctl_keep const DPCTLSyclDeviceRef DRef); +/*! + * @brief Checks if device supports peer access to another device. + * + * @param DRef Opaque pointer to a ``sycl::device`` + * @param PDRef Opaque pointer to a ``sycl::device`` + * @param PT DPCTLPeerAccessType of ``ext::oneapi::peer_access``. + * @return True if sycl::device supports the kind of peer access, else false. + * @ingroup DeviceInterface + */ +DPCTL_API +bool DPCTLDevice_CanAccessPeer(__dpctl_keep const DPCTLSyclDeviceRef DRef, + __dpctl_keep const DPCTLSyclDeviceRef PDRef, + DPCTLPeerAccessType PT); + +/*! + * @brief Enables peer access to another device. + * + * @param DRef Opaque pointer to a ``sycl::device`` + * @param PDRef Opaque pointer to a ``sycl::device`` + * @ingroup DeviceInterface + */ +DPCTL_API +void DPCTLDevice_EnablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, + __dpctl_keep const DPCTLSyclDeviceRef PDRef); + +/*! + * @brief Disables peer access to another device. + * + * @param DRef Opaque pointer to a ``sycl::device`` + * @param PDRef Opaque pointer to a ``sycl::device`` + * @ingroup DeviceInterface + */ +DPCTL_API +void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, + __dpctl_keep const DPCTLSyclDeviceRef PDRef); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h index 799f9d1484..76311daa7b 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h @@ -151,6 +151,16 @@ typedef enum next_partitionable } DPCTLPartitionAffinityDomainType; +/*! + * @brief DPCTL analogue of ``sycl::ext::oneapi::peer_access`` enum. + * + */ +typedef enum +{ + access_supported, + atomics_supported +} DPCTLPeerAccessType; + /*! * @brief Enums to depict the properties that can be passed to a sycl::queue * constructor. diff --git a/libsyclinterface/source/dpctl_sycl_device_interface.cpp b/libsyclinterface/source/dpctl_sycl_device_interface.cpp index 7b1e900b58..1378f6f818 100644 --- a/libsyclinterface/source/dpctl_sycl_device_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_device_interface.cpp @@ -33,7 +33,6 @@ #include "dpctl_sycl_type_casters.hpp" #include "dpctl_utils_helper.h" #include -#include #include #include /* SYCL headers */ #include @@ -116,7 +115,7 @@ __dpctl_give DPCTLSyclDeviceRef DPCTLDevice_CreateFromSelector( { auto Selector = unwrap(DSRef); if (!Selector) { - error_handler("Cannot difine device selector for DPCTLSyclDeviceRef " + error_handler("Cannot define device selector for DPCTLSyclDeviceRef " "as input is a nullptr.", __FILE__, __func__, __LINE__); return nullptr; @@ -184,8 +183,7 @@ DPCTLDevice_GetBackend(__dpctl_keep const DPCTLSyclDeviceRef DRef) DPCTLSyclBackendType BTy = DPCTLSyclBackendType::DPCTL_UNKNOWN_BACKEND; auto D = unwrap(DRef); if (D) { - BTy = DPCTL_SyclBackendToDPCTLBackendType( - D->get_platform().get_backend()); + BTy = DPCTL_SyclBackendToDPCTLBackendType(D->get_backend()); } return BTy; } @@ -903,3 +901,84 @@ DPCTLDevice_GetCompositeDevice(__dpctl_keep const DPCTLSyclDeviceRef DRef) else return nullptr; } + +static inline bool _CallPeerAccess(device dev, device peer) +{ + auto BE1 = dev.get_backend(); + auto BE2 = peer.get_backend(); + + if ((BE1 == BE2) && + (BE1 == sycl::backend::ext_oneapi_level_zero || + BE1 == sycl::backend::ext_oneapi_cuda || + BE1 == sycl::backend::ext_oneapi_hip) && + (BE2 == sycl::backend::ext_oneapi_level_zero || + BE2 == sycl::backend::ext_oneapi_cuda || + BE2 == sycl::backend::ext_oneapi_hip) && + (dev != peer)) + { + return true; + } + return false; +} + +bool DPCTLDevice_CanAccessPeer(__dpctl_keep const DPCTLSyclDeviceRef DRef, + __dpctl_keep const DPCTLSyclDeviceRef PDRef, + DPCTLPeerAccessType PT) +{ + bool canAccess = false; + auto D = unwrap(DRef); + auto PD = unwrap(PDRef); + if (D && PD) { + if (_CallPeerAccess(*D, *PD)) { + try { + canAccess = D->ext_oneapi_can_access_peer( + *PD, DPCTL_DPCTLPeerAccessTypeToSycl(PT)); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + } + return canAccess; +} + +void DPCTLDevice_EnablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, + __dpctl_keep const DPCTLSyclDeviceRef PDRef) +{ + auto D = unwrap(DRef); + auto PD = unwrap(PDRef); + if (D && PD) { + if (_CallPeerAccess(*D, *PD)) { + try { + D->ext_oneapi_enable_peer_access(*PD); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + else { + error_handler("Devices do not support peer access", __FILE__, + __func__, __LINE__); + } + } + return; +} + +void DPCTLDevice_DisablePeerAccess(__dpctl_keep const DPCTLSyclDeviceRef DRef, + __dpctl_keep const DPCTLSyclDeviceRef PDRef) +{ + auto D = unwrap(DRef); + auto PD = unwrap(PDRef); + if (D && PD) { + if (_CallPeerAccess(*D, *PD)) { + try { + D->ext_oneapi_disable_peer_access(*PD); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + else { + error_handler("Devices do not support peer access", __FILE__, + __func__, __LINE__); + } + } + return; +} diff --git a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp index c702018687..70d89fd4fb 100644 --- a/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_kernel_bundle_interface.cpp @@ -620,9 +620,9 @@ DPCTLKernelBundle_CreateFromSpirv(__dpctl_keep const DPCTLSyclContextRef CtxRef, break; #endif default: - error_handler("Backend " + std::to_string(static_cast(BE)) + - " is not supported", - __FILE__, __func__, __LINE__); + std::ostringstream os; + os << "Backend " << BE << " is not supported"; + error_handler(os.str(), __FILE__, __func__, __LINE__); break; } return KBRef; @@ -700,9 +700,9 @@ DPCTLKernelBundle_GetKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, return _GetKernel_ze_impl(*SyclKB, KernelName); #endif default: - error_handler("Backend " + std::to_string(static_cast(be)) + - " is not supported.", - __FILE__, __func__, __LINE__); + std::ostringstream os; + os << "Backend " << be << " is not supported"; + error_handler(os.str(), __FILE__, __func__, __LINE__); return nullptr; } } @@ -730,9 +730,9 @@ bool DPCTLKernelBundle_HasKernel(__dpctl_keep DPCTLSyclKernelBundleRef KBRef, return _HasKernel_ze_impl(*SyclKB, KernelName); #endif default: - error_handler("Backend " + std::to_string(static_cast(be)) + - " is not supported.", - __FILE__, __func__, __LINE__); + std::ostringstream os; + os << "Backend " << be << " is not supported"; + error_handler(os.str(), __FILE__, __func__, __LINE__); return false; } } diff --git a/libsyclinterface/tests/CMakeLists.txt b/libsyclinterface/tests/CMakeLists.txt index 424340260e..dc4465dc2e 100644 --- a/libsyclinterface/tests/CMakeLists.txt +++ b/libsyclinterface/tests/CMakeLists.txt @@ -47,6 +47,7 @@ add_sycl_to_target( ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_selector_interface.cpp ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_device_aspects.cpp ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_event_interface.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_peer_access.cpp ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_platform_interface.cpp ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_kernel_interface.cpp ${CMAKE_CURRENT_SOURCE_DIR}/test_sycl_kernel_bundle_interface.cpp diff --git a/libsyclinterface/tests/test_helper.cpp b/libsyclinterface/tests/test_helper.cpp index 8743d82dfd..25b5743ba2 100644 --- a/libsyclinterface/tests/test_helper.cpp +++ b/libsyclinterface/tests/test_helper.cpp @@ -184,3 +184,32 @@ TEST_F(TestHelperFns, SyclDeviceTypeToDPCTLDeviceType) sycl::info::device_type::custom)); ASSERT_TRUE(DTy == DPCTLSyclDeviceType::DPCTL_CUSTOM); } + +TEST_F(TestHelperFns, ChkDPCTLPeerAccessTypeToSycl) +{ + sycl::ext::oneapi::peer_access peer_type = + sycl::ext::oneapi::peer_access::atomics_supported; + + EXPECT_NO_FATAL_FAILURE(peer_type = DPCTL_DPCTLPeerAccessTypeToSycl( + DPCTLPeerAccessType::access_supported)); + ASSERT_TRUE(peer_type == sycl::ext::oneapi::peer_access::access_supported); + + EXPECT_NO_FATAL_FAILURE(peer_type = DPCTL_DPCTLPeerAccessTypeToSycl( + DPCTLPeerAccessType::atomics_supported)); + ASSERT_TRUE(peer_type == sycl::ext::oneapi::peer_access::atomics_supported); +} + +TEST_F(TestHelperFns, ChkSyclPeerAccessToDPCTLType) +{ + DPCTLPeerAccessType PTy = DPCTLPeerAccessType::atomics_supported; + + EXPECT_NO_FATAL_FAILURE( + PTy = DPCTL_SyclPeerAccessToDPCTLType( + sycl::ext::oneapi::peer_access::access_supported)); + ASSERT_TRUE(PTy == DPCTLPeerAccessType::access_supported); + + EXPECT_NO_FATAL_FAILURE( + PTy = DPCTL_SyclPeerAccessToDPCTLType( + sycl::ext::oneapi::peer_access::atomics_supported)); + ASSERT_TRUE(PTy == DPCTLPeerAccessType::atomics_supported); +} diff --git a/libsyclinterface/tests/test_sycl_peer_access.cpp b/libsyclinterface/tests/test_sycl_peer_access.cpp new file mode 100644 index 0000000000..8e3179b021 --- /dev/null +++ b/libsyclinterface/tests/test_sycl_peer_access.cpp @@ -0,0 +1,149 @@ +//===--- test_sycl_peer_access.cpp - Test cases for device peer access ===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2025 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file has unit test cases for peer access functions defined in +/// dpctl_sycl_device_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dpctl_sycl_device_interface.h" +#include "dpctl_sycl_device_selector_interface.h" +#include "dpctl_sycl_platform_interface.h" +#include "dpctl_utils_helper.h" + +#include +#include + +struct TestDPCTLPeerAccess : public ::testing::TestWithParam +{ + DPCTLSyclPlatformRef P = nullptr; + DPCTLDeviceVectorRef DV = nullptr; + + TestDPCTLPeerAccess() + { + auto DS = DPCTLFilterSelector_Create(GetParam()); + if (DS) { + EXPECT_NO_FATAL_FAILURE(P = DPCTLPlatform_CreateFromSelector(DS)); + } + DPCTLDeviceSelector_Delete(DS); + if (P) { + DV = DPCTLPlatform_GetDevices(P, DPCTLSyclDeviceType::DPCTL_ALL); + } + } + + void SetUp() + { + if (!P || !DV) { + auto message = "Skipping as no devices of type " + + std::string(GetParam()) + "."; + GTEST_SKIP_(message.c_str()); + } + + if (DPCTLDeviceVector_Size(DV) < 2) { + GTEST_SKIP_("Peer access tests require more than one device."); + } + } + + ~TestDPCTLPeerAccess() + { + DPCTLDeviceVector_Delete(DV); + DPCTLPlatform_Delete(P); + } +}; + +TEST_P(TestDPCTLPeerAccess, ChkAccessSupported) +{ + auto D0 = DPCTLDeviceVector_GetAt(DV, 0); + auto D1 = DPCTLDeviceVector_GetAt(DV, 1); + ASSERT_TRUE(D0 != nullptr); + ASSERT_TRUE(D1 != nullptr); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_CanAccessPeer( + D0, D1, DPCTLPeerAccessType::access_supported)); +} + +TEST_P(TestDPCTLPeerAccess, ChkAtomicsSupported) +{ + auto D0 = DPCTLDeviceVector_GetAt(DV, 0); + auto D1 = DPCTLDeviceVector_GetAt(DV, 1); + ASSERT_TRUE(D0 != nullptr); + ASSERT_TRUE(D1 != nullptr); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_CanAccessPeer( + D0, D1, DPCTLPeerAccessType::atomics_supported)); +} + +TEST_P(TestDPCTLPeerAccess, ChkPeerAccess) +{ + auto D0 = DPCTLDeviceVector_GetAt(DV, 0); + auto D1 = DPCTLDeviceVector_GetAt(DV, 1); + ASSERT_TRUE(D0 != nullptr); + ASSERT_TRUE(D1 != nullptr); + bool canEnable = false; + EXPECT_NO_FATAL_FAILURE(canEnable = DPCTLDevice_CanAccessPeer( + D0, D1, DPCTLPeerAccessType::access_supported)); + if (canEnable) { + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_EnablePeerAccess(D0, D1)); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_DisablePeerAccess(D0, D1)); + } +} + +TEST_P(TestDPCTLPeerAccess, ChkPeerAccessToSelf) +{ + auto D0 = DPCTLDeviceVector_GetAt(DV, 0); + ASSERT_TRUE(D0 != nullptr); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_CanAccessPeer( + D0, D0, DPCTLPeerAccessType::access_supported)); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_EnablePeerAccess(D0, D0)); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_DisablePeerAccess(D0, D0)); +} + +INSTANTIATE_TEST_SUITE_P(DPCTLDeviceFns, + TestDPCTLPeerAccess, + ::testing::Values("level_zero", "cuda", "hip")); + +struct TestDPCTLPeerAccessNullArgs : public ::testing::Test +{ + DPCTLSyclDeviceRef Null_DR0 = nullptr; + DPCTLSyclDeviceRef Null_DR1 = nullptr; +}; + +TEST_F(TestDPCTLPeerAccessNullArgs, ChkAccessSupported) +{ + bool accessSupported = true; + EXPECT_NO_FATAL_FAILURE( + accessSupported = DPCTLDevice_CanAccessPeer( + Null_DR0, Null_DR1, DPCTLPeerAccessType::access_supported)); + ASSERT_FALSE(accessSupported); +} + +TEST_F(TestDPCTLPeerAccessNullArgs, ChkAtomicsSupported) +{ + bool accessSupported = true; + EXPECT_NO_FATAL_FAILURE( + accessSupported = DPCTLDevice_CanAccessPeer( + Null_DR0, Null_DR1, DPCTLPeerAccessType::atomics_supported)); + ASSERT_FALSE(accessSupported); +} + +TEST_F(TestDPCTLPeerAccessNullArgs, ChkPeerAccess) +{ + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_EnablePeerAccess(Null_DR0, Null_DR1)); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_DisablePeerAccess(Null_DR0, Null_DR1)); +}