diff --git a/CHANGELOG.md b/CHANGELOG.md index 6ecef415e2..83fb5ac1c4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,11 +11,14 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Added `add_event` and `add_cleanup_event` methods and `cleanup_events` and `num_cleanup_events` properties to the order manager, for tracking the events that gate the release of objects used by offloaded tasks [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * Added `SyclQueue_DeferRelease` C-API function, which defers a release callback until work submitted to a queue completes, and through which `dpctl::utils::keep_args_alive` keeps its arguments alive [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * Added `dpctl::utils::defer_release` to `dpctl4pybind11.hpp`, which calls a release callback once the given events have completed [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) +* Added `custom` device type to filter selector strings [gh-2420](https://github.com/IntelPython/dpctl/pull/2420) ### Changed * Object lifetime during offload is now managed by a list of releases deferred until the events that gate them have completed, rather than by a `host_task`. `dpctl` keeps one thread for looking over the list and running the releases that have come due, started by the first release deferred and again in the child of a `fork`, which takes the GIL to do so and lets a release call into the interpreter and free USM allocations [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * `import_dpctl()` now returns 0 on success and -1 with a Python exception set on failure, and `dpctl4pybind11.hpp` raises that exception rather than failing on the first C-API call [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * The event returned by `dpctl.SyclQueue.keep_args_alive`, `dpctl::utils::keep_args_alive` and `dpctl.SyclQueue._submit_keep_args_alive` is now for an empty kernel that gates the deferred release rather than for a `host_task` that performs it, so its completion means that the objects are no longer in use rather than that they have been released [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) +* Filter selector strings are implemented in `dpctl` instead of relying on the deprecated DPC++ `sycl_ext_oneapi_filter_selector` extension, and device numbers in a list of filters, such as `"0,1"`, now refer to the same devices as when used in separate filters [gh-2420](https://github.com/IntelPython/dpctl/pull/2420) +* `dpctl.SyclDevice.filter_string` and `dpctl.SyclDevice.get_filter_string` raise `ValueError` for backends and device types not supported in filter selector strings, instead of returning a string containing `"unknown"` [gh-2420](https://github.com/IntelPython/dpctl/pull/2420) ### Deprecated * Deprecated `dpctl.SyclQueue._submit_keep_args_alive` in favor of `dpctl.SyclQueue.keep_args_alive`, which it now calls and which behaves identically [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) @@ -28,6 +31,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Removed stale `select_host_device` and `has_host_device` names from `dpctl.__all__`, which made `from dpctl import *` fail with `AttributeError` [gh-2410](https://github.com/IntelPython/dpctl/pull/2410) * A USM allocation that fails now drops the references held for offloaded tasks that have completed, and is attempted once more if that released anything, so that memory only waiting to be given up is not reported as unavailable [gh-2359](https://github.com/IntelPython/dpctl/pull/2359) * `dpctl.SyclPlatform.get_devices` no longer prints Unified Runtime warnings for `device_type="custom"` on non-OpenCL platforms, where custom devices are not supported [gh-2306](https://github.com/IntelPython/dpctl/pull/2306) +* `DPCTLDeviceSelector_Score` returns correct scores for filter selectors with device numbers [gh-2420](https://github.com/IntelPython/dpctl/pull/2420) ### Maintenance * Removed the `-fno-sycl-id-queries-fit-in-int` compile option, which none of the `dpctl` extensions need since `dpctl.tensor` moved to `dpnp` [gh-2413](https://github.com/IntelPython/dpctl/pull/2413) diff --git a/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst b/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst index f96a8764df..2171e17bc4 100644 --- a/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst +++ b/docs/doc_sources/api_reference/dpctl/filter_selector_string.rst @@ -3,12 +3,10 @@ Filter Selector String ====================== -Filter selector string is a way to refer to unpartitioned SYCL devices -proposed in :oneapi_filter_selection:`sycl_ext_oneapi_filter_selector <>` -extension to SYCL standard. - -This document captures aspects of the referenced document relevant -to :mod:`dpctl`. +Filter selector string is a way to refer to unpartitioned SYCL devices. +:mod:`dpctl` implements the syntax of the DPC++ +:oneapi_filter_selection:`sycl_ext_oneapi_filter_selector <>` extension +without relying on the extension itself. A filter selector string defines one or more filters, which must be separated using ``","`` character. A filter is specified as a @@ -18,8 +16,8 @@ triple of the form: Backend:DeviceType:RelativeDeviceNumber -Every element of the triple is optional, but a filter must contain at -least one component. +Every element of the triple is optional and the elements may be given in +any order. ``Backend`` specifies the desired backend of targeted devices, while ``DeviceType`` specifies the type of targeted devices. @@ -27,6 +25,9 @@ least one component. any other given requirements, starting from 0 to marking the "first device that matches the requirements". +If several devices match, the one with the highest default selector +score is selected. + Attempting to use a non-conforming string in places where filter selector string is expected will raise an exception. @@ -45,6 +46,7 @@ Supported values for ``DeviceType`` are: * - accelerator - cpu + - custom - gpu Filter selector strings can be used as arguments to constructors of diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index 624575a5d8..85af18c276 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -239,6 +239,7 @@ cdef tuple _to_enum_tuple( cdef str _backend_type_to_filter_string_part(_backend_type BTy): + """ Returns the filter string part for the backend, or ``None``.""" if BTy == _backend_type._CUDA: return "cuda" elif BTy == _backend_type._HIP: @@ -248,20 +249,57 @@ cdef str _backend_type_to_filter_string_part(_backend_type BTy): elif BTy == _backend_type._OPENCL: return "opencl" else: - return "unknown" + return None cdef str _device_type_to_filter_string_part(_device_type DTy): + """ Returns the filter string part for the device type, or ``None``.""" if DTy == _device_type._ACCELERATOR: return "accelerator" - elif DTy == _device_type._AUTOMATIC: - return "automatic" elif DTy == _device_type._CPU: return "cpu" + elif DTy == _device_type._CUSTOM: + return "custom" elif DTy == _device_type._GPU: return "gpu" else: - return "unknown" + return None + + +cdef str _make_filter_string( + DPCTLSyclDeviceRef DRef, + bint include_backend, + bint include_device_type, + int64_t relId, +): + """ Joins the requested parts of the filter string of a root device. + + Raises ``ValueError`` if the backend or device type has no filter + string representation. + """ + cdef list parts = [] + if include_backend: + be_str = _backend_type_to_filter_string_part( + DPCTLDevice_GetBackend(DRef) + ) + if be_str is None: + raise ValueError( + "The backend of this SyclDevice is not supported in filter " + "strings" + ) + parts.append(be_str) + if include_device_type: + dt_str = _device_type_to_filter_string_part( + DPCTLDevice_GetDeviceType(DRef) + ) + if dt_str is None: + raise ValueError( + "The device type of this SyclDevice is not supported in " + "filter strings" + ) + parts.append(dt_str) + parts.append(str(relId)) + return ":".join(parts) cdef void _init_helper(_SyclDevice device, DPCTLSyclDeviceRef DRef) except *: @@ -317,22 +355,17 @@ cdef inline void _raise_invalid_peer_access( ] cdef _backend_type BTy1 = DPCTLDevice_GetBackend(dev._device_ref) cdef _backend_type BTy2 = DPCTLDevice_GetBackend(peer.get_device_ref()) + be1 = _backend_type_to_filter_string_part(BTy1) or "unknown" + be2 = _backend_type_to_filter_string_part(BTy2) or "unknown" 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)}" + f"Device with backend {be1} " + f"cannot peer access device with backend {be2}" ) if (BTy1 not in _peer_access_backends): - raise ValueError( - "Peer access not supported for backend " - f"{_backend_type_to_filter_string_part(BTy1)}" - ) + raise ValueError(f"Peer access not supported for backend {be1}") if (BTy2 not in _peer_access_backends): - raise ValueError( - "Peer access not supported for backend " - f"{_backend_type_to_filter_string_part(BTy2)}" - ) + raise ValueError(f"Peer access not supported for backend {be2}") if (dev == peer): raise ValueError( "Peer access cannot be enabled between a device and itself" @@ -355,20 +388,15 @@ def _cached_filter_string(d : SyclDevice): if the device is a root (unpartitioned) device. Raises: - ValueError: if the input device is a sub-device. + ValueError: if the input device is a sub-device, or its backend or + device type is not supported in filter strings. """ - cdef _backend_type BTy - cdef _device_type DTy cdef int64_t relId = -1 cdef SyclDevice cd = d relId = DPCTLDeviceMgr_GetRelativeId(cd._device_ref) if (relId == -1): raise ValueError("This SyclDevice is not a root device") - BTy = DPCTLDevice_GetBackend(cd._device_ref) - br_str = _backend_type_to_filter_string_part(BTy) - DTy = DPCTLDevice_GetDeviceType(cd._device_ref) - dt_str = _device_type_to_filter_string_part(DTy) - return ":".join((br_str, dt_str, str(relId))) + return _make_filter_string(cd._device_ref, True, True, relId) cdef class SyclDevice(_SyclDevice): @@ -379,7 +407,7 @@ cdef class SyclDevice(_SyclDevice): - by directly passing in a filter string to the class constructor. The filter string needs to conform to the - :oneapi_filter_selection:`DPC++ filter selector SYCL extension <>`. + :ref:`filter selector string ` syntax. :Example: @@ -2640,6 +2668,9 @@ cdef class SyclDevice(_SyclDevice): ValueError: If the device is a sub-device. + If the backend or device type of the device is not + supported in filter strings. + :Example: .. code-block:: python @@ -2746,6 +2777,9 @@ cdef class SyclDevice(_SyclDevice): If no match for the device was found in the vector returned by ``sycl::device::get_devices()`` + If the requested backend or device type of the device is + not supported in filter strings. + :Example: .. code-block:: python @@ -2760,8 +2794,6 @@ cdef class SyclDevice(_SyclDevice): """ cdef int relId = -1 cdef DPCTLSyclDeviceRef pDRef = NULL - cdef _device_type DTy - cdef _backend_type BTy if include_backend: if include_device_type: @@ -2783,22 +2815,9 @@ cdef class SyclDevice(_SyclDevice): DPCTLDevice_Delete(pDRef) raise ValueError("This SyclDevice is not a root device") else: - if include_backend: - BTy = DPCTLDevice_GetBackend(self._device_ref) - be_str = _backend_type_to_filter_string_part(BTy) - if include_device_type: - DTy = DPCTLDevice_GetDeviceType(self._device_ref) - dt_str = _device_type_to_filter_string_part(DTy) - return ":".join((be_str, dt_str, str(relId))) - else: - return ":".join((be_str, str(relId))) - else: - if include_device_type: - DTy = DPCTLDevice_GetDeviceType(self._device_ref) - dt_str = _device_type_to_filter_string_part(DTy) - return ":".join((dt_str, str(relId))) - else: - return str(relId) + return _make_filter_string( + self._device_ref, include_backend, include_device_type, relId + ) def get_unpartitioned_parent_device(self): """ get_unpartitioned_parent_device() diff --git a/dpctl/tests/_device_attributes_checks.py b/dpctl/tests/_device_attributes_checks.py index 7e00831fae..2c3785c915 100644 --- a/dpctl/tests/_device_attributes_checks.py +++ b/dpctl/tests/_device_attributes_checks.py @@ -46,6 +46,12 @@ "opencl:gpu:-1", "cuda:cpu:0", "abc", + "OPENCL", + "opencl:gpu:cpu", + "opencl:level_zero", + "0:1", + "opencl:gpu:0:0", + "", 1, ] diff --git a/libsyclinterface/helper/include/dpctl_utils_helper.h b/libsyclinterface/helper/include/dpctl_utils_helper.h index c441a40aec..70eeae3295 100644 --- a/libsyclinterface/helper/include/dpctl_utils_helper.h +++ b/libsyclinterface/helper/include/dpctl_utils_helper.h @@ -295,7 +295,7 @@ int64_t DPCTL_GetRelativeDeviceId(const sycl::device &Device); /*! * @brief Gives the filter string which would select given root device if - * used as argument to ``sycl::ext::oneapi::filter_selector``. Throws exception + * used as argument to ``DPCTLFilterSelector_Create``. Throws exception * if filter string can not be constructed. * * @param Device A ``sycl::device`` object whose filter selector diff --git a/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp b/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp index 7af8ef5e86..a5ed9e4016 100644 --- a/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp +++ b/libsyclinterface/include/syclinterface/dpctl_device_selection.hpp @@ -28,7 +28,10 @@ #pragma once #include "Support/DllExport.h" +#include +#include #include +#include namespace dpctl { @@ -71,14 +74,31 @@ class DPCTL_API dpctl_cpu_selector : public dpctl_device_selector int operator()(const sycl::device &d) const override; }; +/*! + * @brief Selects devices matching a filter selector string. + * + * The string is a comma-separated list of filters, each of the form + * ``backend:device_type:device_num`` with every component optional and in + * any order. ``device_num`` counts devices from ``sycl::device::get_devices()`` + * that match the filter's backend and device type. Matching devices are + * scored by ``sycl::default_selector_v``. + * + * Throws ``std::invalid_argument`` if the string is malformed. + */ class DPCTL_API dpctl_filter_selector : public dpctl_device_selector { public: - dpctl_filter_selector(const std::string &fs) : _impl(fs) {} + explicit dpctl_filter_selector(const std::string &fs); int operator()(const sycl::device &d) const override; private: - sycl::ext::oneapi::filter_selector _impl; + struct filter + { + std::optional backend; + std::optional device_type; + std::optional device; + }; + std::vector _filters; }; } // namespace syclinterface diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_device_selector_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_device_selector_interface.h index 74d04a256f..e4d7a83823 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_device_selector_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_device_selector_interface.h @@ -65,12 +65,12 @@ DPCTL_API __dpctl_give DPCTLSyclDeviceSelectorRef DPCTLCPUSelector_Create(void); /*! - * @brief Returns an opaque wrapper for sycl::ext::oneapi::filter_selector - * object based on the passed in filter string. + * @brief Returns a device selector for the passed in filter string. * * @param filter_str A C string providing a filter based on which to * create a device selector. - * @return An opaque pointer to a sycl::ext::oneapi::filter_selector object. + * @return An opaque pointer to a device selector, or nullptr if the + * filter string is malformed. * @ingroup DeviceSelectors */ DPCTL_API diff --git a/libsyclinterface/source/dpctl_device_selection.cpp b/libsyclinterface/source/dpctl_device_selection.cpp index 2e216162f4..5bf44e94df 100644 --- a/libsyclinterface/source/dpctl_device_selection.cpp +++ b/libsyclinterface/source/dpctl_device_selection.cpp @@ -27,14 +27,99 @@ #include "dpctl_device_selection.hpp" #include "Config/dpctl_config.h" +#include +#include +#include +#include +#include #include +#include +#include namespace { static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, "The compiler does not meet minimum version requirement"); + +/*! Splits s on delim, dropping empty tokens. */ +std::vector tokenize(const std::string &s, char delim) +{ + std::vector tokens; + std::size_t start = 0; + while (start <= s.size()) { + auto end = s.find(delim, start); + if (end == std::string::npos) + end = s.size(); + if (end > start) + tokens.emplace_back(s, start, end - start); + start = end + 1; + } + return tokens; +} + +std::optional backend_from_token(const std::string &tok) +{ + if (tok == "opencl") + return sycl::backend::opencl; + if (tok == "level_zero") + return sycl::backend::ext_oneapi_level_zero; + if (tok == "cuda") + return sycl::backend::ext_oneapi_cuda; + if (tok == "hip") + return sycl::backend::ext_oneapi_hip; + return std::nullopt; +} + +std::optional +device_type_from_token(const std::string &tok) +{ + if (tok == "cpu") + return sycl::info::device_type::cpu; + if (tok == "gpu") + return sycl::info::device_type::gpu; + if (tok == "accelerator") + return sycl::info::device_type::accelerator; + if (tok == "custom") + return sycl::info::device_type::custom; + return std::nullopt; +} + +/*! Parses a non-negative decimal int; nullopt on any other input. */ +std::optional device_num_from_token(const std::string &tok) +{ + for (char c : tok) { + if (c < '0' || c > '9') + return std::nullopt; + } + int num = 0; + const char *first = tok.data(); + const char *last = first + tok.size(); + auto [ptr, ec] = std::from_chars(first, last, num); + if (ec != std::errc{} || ptr != last) + return std::nullopt; + return num; } +bool matches_backend_and_type(const std::optional &be, + const std::optional &dt, + const sycl::device &d) +{ + if (be && d.get_backend() != *be) + return false; + if (dt && d.get_info() != *dt) + return false; + return true; +} + +[[noreturn]] void throw_invalid_filter(const std::string &fs) +{ + throw std::invalid_argument( + "Invalid filter string '" + fs + + "'. Filters are comma-separated and have the form " + "backend:device_type:device_num, where each component is optional."); +} +} // namespace + namespace dpctl { namespace syclinterface @@ -66,9 +151,63 @@ int dpctl_cpu_selector::operator()(const sycl::device &d) const return sycl::cpu_selector_v(d); } +dpctl_filter_selector::dpctl_filter_selector(const std::string &fs) +{ + std::vector>> parsed; + for (const auto &filter_str : tokenize(fs, ',')) { + auto tokens = tokenize(filter_str, ':'); + if (tokens.size() > 3) + throw_invalid_filter(fs); + + filter f; + std::optional device_num; + for (const auto &tok : tokens) { + if (auto be = backend_from_token(tok); be && !f.backend) + f.backend = be; + else if (auto dt = device_type_from_token(tok); + dt && !f.device_type) + f.device_type = dt; + else if (auto num = device_num_from_token(tok); num && !device_num) + device_num = num; + else + throw_invalid_filter(fs); + } + parsed.emplace_back(std::move(f), device_num); + } + + // Resolve device numbers; a filter whose device does not exist is dropped + std::vector devices; + bool have_devices = false; + for (auto &[f, device_num] : parsed) { + if (device_num) { + if (!have_devices) { + devices = sycl::device::get_devices(); + have_devices = true; + } + int idx = 0; + for (const auto &d : devices) { + if (!matches_backend_and_type(f.backend, f.device_type, d)) + continue; + if (idx++ == *device_num) { + f.device = d; + break; + } + } + if (!f.device) + continue; + } + _filters.push_back(std::move(f)); + } +} + int dpctl_filter_selector::operator()(const sycl::device &d) const { - return _impl(d); + for (const auto &f : _filters) { + if (matches_backend_and_type(f.backend, f.device_type, d) && + (!f.device || *f.device == d)) + return sycl::default_selector_v(d); + } + return REJECT_DEVICE; } } // namespace syclinterface diff --git a/libsyclinterface/tests/test_sycl_device_selector_interface.cpp b/libsyclinterface/tests/test_sycl_device_selector_interface.cpp index d6a9d3f4ee..37f5f17de2 100644 --- a/libsyclinterface/tests/test_sycl_device_selector_interface.cpp +++ b/libsyclinterface/tests/test_sycl_device_selector_interface.cpp @@ -28,7 +28,9 @@ #include "dpctl_sycl_device_manager.h" #include "dpctl_sycl_device_selector_interface.h" #include "dpctl_sycl_type_casters.hpp" +#include "dpctl_utils_helper.h" #include +#include #include using namespace sycl; @@ -205,6 +207,76 @@ TEST_F(TestDeviceSelectorInterface, ChkDPCTLGPUSelectorScore) ASSERT_TRUE(score < 0); } +TEST_F(TestDeviceSelectorInterface, ChkDPCTLFilterSelectorScore) +{ + using dpctl::syclinterface::unwrap; + using dpctl::syclinterface::wrap; + + const auto devices = sycl::device::get_devices(); + for (std::size_t i = 0; i < devices.size(); ++i) { + const std::string filters[] = {std::to_string(i), + DPCTL_GetDeviceFilterString(devices[i])}; + for (const auto &fs : filters) { + DPCTLSyclDeviceSelectorRef DSRef = nullptr; + EXPECT_NO_FATAL_FAILURE(DSRef = + DPCTLFilterSelector_Create(fs.c_str())); + ASSERT_TRUE(DSRef != nullptr); + // scores must not depend on earlier calls + for (int pass = 0; pass < 2; ++pass) { + for (std::size_t j = 0; j < devices.size(); ++j) { + auto D = devices[j]; + auto DRef = wrap(&D); + int score = DPCTLDeviceSelector_Score(DSRef, DRef); + if (i == j) + EXPECT_TRUE(score >= 0) << fs << " pass " << pass; + else + EXPECT_TRUE(score < 0) << fs << " pass " << pass; + } + } + EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef)); + } + } +} + +TEST_F(TestDeviceSelectorInterface, ChkDPCTLFilterSelectorList) +{ + const auto devices = sycl::device::get_devices(); + if (devices.size() < 2) + GTEST_SKIP_("Need at least two devices."); + + // "0,1" must accept both device 0 and device 1 + DPCTLSyclDeviceSelectorRef DSRef = nullptr; + EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLFilterSelector_Create("0,1")); + ASSERT_TRUE(DSRef != nullptr); + for (std::size_t j = 0; j < devices.size(); ++j) { + auto D = devices[j]; + auto DRef = dpctl::syclinterface::wrap(&D); + int score = DPCTLDeviceSelector_Score(DSRef, DRef); + if (j < 2) + EXPECT_TRUE(score >= 0); + else + EXPECT_TRUE(score < 0); + } + EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef)); +} + +TEST_F(TestDeviceSelectorInterface, ChkDPCTLFilterSelectorCustom) +{ + for (const char *fs : {"custom", "opencl:custom", "custom:0"}) { + DPCTLSyclDeviceSelectorRef DSRef = nullptr; + EXPECT_NO_FATAL_FAILURE(DSRef = DPCTLFilterSelector_Create(fs)); + ASSERT_TRUE(DSRef != nullptr) << fs; + for (auto D : sycl::device::get_devices()) { + if (D.get_info() == + sycl::info::device_type::custom) + continue; + auto DRef = dpctl::syclinterface::wrap(&D); + EXPECT_TRUE(DPCTLDeviceSelector_Score(DSRef, DRef) < 0) << fs; + } + EXPECT_NO_FATAL_FAILURE(DPCTLDeviceSelector_Delete(DSRef)); + } +} + INSTANTIATE_TEST_SUITE_P(FilterSelectorCreation, TestFilterSelector, ::testing::Values("opencl", @@ -220,8 +292,25 @@ INSTANTIATE_TEST_SUITE_P(FilterSelectorCreation, "gpu:0", "gpu:1", "1", - "0")); + "0", + "gpu:opencl", + "0:gpu", + "opencl::0", + ":", + "gpu,cpu", + "cuda,cpu")); INSTANTIATE_TEST_SUITE_P(NegativeFilterSelectorCreation, TestUnsupportedFilters, - ::testing::Values("abc", "-1", "cuda:cpu:0")); + ::testing::Values("abc", + "-1", + "cuda:cpu:0", + "", + ",", + "OPENCL", + " opencl", + "opencl:gpu:cpu", + "opencl:level_zero", + "0:1", + "opencl:gpu:0:0", + "99999999999"));