From de581f686683d2be77a792e30cb0a153b65d6d0b Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Thu, 6 Aug 2026 07:10:40 -0700 Subject: [PATCH 1/2] add sycl::info::context queries --- dpctl/_backend.pxd | 14 ++ dpctl/_sycl_context.pyx | 141 +++++++++++++++++ dpctl/tests/test_sycl_context.py | 68 +++++++++ .../dpctl_sycl_context_interface.h | 69 +++++++++ .../source/dpctl_sycl_context_interface.cpp | 86 +++++++++++ .../tests/test_sycl_context_interface.cpp | 144 ++++++++++++++++++ 6 files changed, 522 insertions(+) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index a85565b729..ea0665b312 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -512,6 +512,20 @@ cdef extern from "syclinterface/dpctl_sycl_context_interface.h": cdef size_t DPCTLContext_Hash(const DPCTLSyclContextRef CRef) cdef _backend_type DPCTLContext_GetBackend(const DPCTLSyclContextRef) cdef void DPCTLContext_Delete(DPCTLSyclContextRef CtxRef) + cdef DPCTLSyclPlatformRef DPCTLContext_GetPlatform( + const DPCTLSyclContextRef CRef) + cdef int *DPCTLContext_GetAtomicMemoryOrderCapabilities( + const DPCTLSyclContextRef CRef, + size_t *res_len) + cdef int *DPCTLContext_GetAtomicFenceOrderCapabilities( + const DPCTLSyclContextRef CRef, + size_t *res_len) + cdef int *DPCTLContext_GetAtomicMemoryScopeCapabilities( + const DPCTLSyclContextRef CRef, + size_t *res_len) + cdef int *DPCTLContext_GetAtomicFenceScopeCapabilities( + const DPCTLSyclContextRef CRef, + size_t *res_len) cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": diff --git a/dpctl/_sycl_context.pyx b/dpctl/_sycl_context.pyx index 0cafc57c99..f48b6d34c3 100644 --- a/dpctl/_sycl_context.pyx +++ b/dpctl/_sycl_context.pyx @@ -34,7 +34,12 @@ from ._backend cimport ( # noqa: E211 DPCTLContext_CreateFromDevices, DPCTLContext_Delete, DPCTLContext_DeviceCount, + DPCTLContext_GetAtomicFenceOrderCapabilities, + DPCTLContext_GetAtomicFenceScopeCapabilities, + DPCTLContext_GetAtomicMemoryOrderCapabilities, + DPCTLContext_GetAtomicMemoryScopeCapabilities, DPCTLContext_GetDevices, + DPCTLContext_GetPlatform, DPCTLContext_Hash, DPCTLDeviceMgr_GetCachedContext, DPCTLDeviceVector_CreateFromArray, @@ -42,12 +47,17 @@ from ._backend cimport ( # noqa: E211 DPCTLDeviceVector_GetAt, DPCTLDeviceVector_Size, DPCTLDeviceVectorRef, + DPCTLInt_Array_Delete, DPCTLSyclContextRef, DPCTLSyclDeviceRef, + DPCTLSyclPlatformRef, error_handler_callback, ) from ._sycl_device cimport SyclDevice from ._sycl_device import SyclDeviceCreationError +from ._sycl_platform cimport SyclPlatform + +from .enum_types import memory_order, memory_scope __all__ = [ "SyclContext", @@ -84,6 +94,33 @@ cdef void _init_helper(_SyclContext context, DPCTLSyclContextRef CRef): context._ctxt_ref = CRef +cdef tuple _to_enum_tuple( + int *arr, size_t arr_len, object enum_type, str descr +): + """ + Converts an array of DPCTL enum values into a tuple of ``enum_type``s + + The DPCTL enums reserve value 0 for an unrecognized value, so a DPCTL + value of ``n`` corresponds to the ``n``-th member of ``enum_type``, whose + members are numbered from 1 by ``enum.auto()``. + """ + cdef list res = [] + cdef size_t i + + if arr is NULL: + return () + try: + for i in range(arr_len): + try: + res.append(enum_type(arr[i])) + except ValueError: + raise RuntimeError(f"Unrecognized {descr} reported") + finally: + DPCTLInt_Array_Delete(arr) + + return tuple(res) + + cdef class _SyclContext: """ Data owner for SyclContext """ @@ -442,6 +479,110 @@ cdef class SyclContext(_SyclContext): "associated with this context" ) + @property + def sycl_platform(self): + """ Returns the platform associated with this context. + + Returns: + :class:`dpctl.SyclPlatform`: + The platform associated with this context. + + Raises: + RuntimeError: + If ``DPCTLContext_GetPlatform`` fails to return a platform. + """ + cdef DPCTLSyclPlatformRef PRef = ( + DPCTLContext_GetPlatform(self.get_context_ref()) + ) + if (PRef == NULL): + raise RuntimeError("Could not get platform for context.") + else: + return SyclPlatform._create(PRef) + + @property + def atomic_memory_order_capabilities(self): + """ Returns a tuple of :class:`dpctl.memory_order` describing atomic + memory order capabilities of the context. + + Returns: + Tuple[:class:`dpctl.memory_order`]: + Tuple of supported memory orders. + + Raises: + RuntimeError: + If an unrecognized memory order is given by runtime. + """ + cdef int *arr = NULL + cdef size_t arr_len = 0 + + arr = DPCTLContext_GetAtomicMemoryOrderCapabilities( + self.get_context_ref(), &arr_len + ) + return _to_enum_tuple(arr, arr_len, memory_order, "memory order") + + @property + def atomic_fence_order_capabilities(self): + """ Returns a tuple of :class:`dpctl.memory_order` describing atomic + fence order capabilities of the context. + + Returns: + Tuple[:class:`dpctl.memory_order`]: + Tuple of supported fence orders. + + Raises: + RuntimeError: + If an unrecognized memory order is given by runtime. + """ + cdef int *arr = NULL + cdef size_t arr_len = 0 + + arr = DPCTLContext_GetAtomicFenceOrderCapabilities( + self.get_context_ref(), &arr_len + ) + return _to_enum_tuple(arr, arr_len, memory_order, "memory order") + + @property + def atomic_memory_scope_capabilities(self): + """ Returns a tuple of :class:`dpctl.memory_scope` describing atomic + memory scope capabilities of the context. + + Returns: + Tuple[:class:`dpctl.memory_scope`]: + Tuple of supported memory scopes. + + Raises: + RuntimeError: + If an unrecognized memory scope is given by runtime. + """ + cdef int *arr = NULL + cdef size_t arr_len = 0 + + arr = DPCTLContext_GetAtomicMemoryScopeCapabilities( + self.get_context_ref(), &arr_len + ) + return _to_enum_tuple(arr, arr_len, memory_scope, "memory scope") + + @property + def atomic_fence_scope_capabilities(self): + """ Returns a tuple of :class:`dpctl.memory_scope` describing atomic + fence scope capabilities of the context. + + Returns: + Tuple[:class:`dpctl.memory_scope`]: + Tuple of supported fence scopes. + + Raises: + RuntimeError: + If an unrecognized memory scope is given by runtime. + """ + cdef int *arr = NULL + cdef size_t arr_len = 0 + + arr = DPCTLContext_GetAtomicFenceScopeCapabilities( + self.get_context_ref(), &arr_len + ) + return _to_enum_tuple(arr, arr_len, memory_scope, "memory scope") + @property def __name__(self): return "SyclContext" diff --git a/dpctl/tests/test_sycl_context.py b/dpctl/tests/test_sycl_context.py index fc3dcfb52d..891e9093c7 100644 --- a/dpctl/tests/test_sycl_context.py +++ b/dpctl/tests/test_sycl_context.py @@ -273,3 +273,71 @@ def test_multi_device_different_platforms(): dpctl.SyclContext(devs) else: pytest.skip("Insufficient amount of available devices for this test") + + +def test_context_sycl_platform(valid_filter): + """ + Test that :attr:`dpctl.SyclContext.sycl_platform` returns the + platform shared by the context's devices. + """ + try: + ctx = dpctl.SyclContext(valid_filter) + except dpctl.SyclContextCreationError: + pytest.skip() + plat = ctx.sycl_platform + assert isinstance(plat, dpctl.SyclPlatform) + for d in ctx.get_devices(): + assert d.sycl_platform == plat + + +def test_context_atomic_memory_order_capabilities(valid_filter): + try: + ctx = dpctl.SyclContext(valid_filter) + except dpctl.SyclContextCreationError: + pytest.skip() + caps = ctx.atomic_memory_order_capabilities + assert isinstance(caps, tuple) + assert all(isinstance(m, dpctl.memory_order) for m in caps) + # SYCL 2020 requires at least these capabilities + assert dpctl.memory_order.relaxed in caps + # capabilities of the context must be a subset of every device + for d in ctx.get_devices(): + assert set(caps).issubset(set(d.atomic_memory_order_capabilities)) + + +def test_context_atomic_fence_order_capabilities(valid_filter): + try: + ctx = dpctl.SyclContext(valid_filter) + except dpctl.SyclContextCreationError: + pytest.skip() + caps = ctx.atomic_fence_order_capabilities + assert isinstance(caps, tuple) + assert all(isinstance(m, dpctl.memory_order) for m in caps) + for d in ctx.get_devices(): + assert set(caps).issubset(set(d.atomic_fence_order_capabilities)) + + +def test_context_atomic_memory_scope_capabilities(valid_filter): + try: + ctx = dpctl.SyclContext(valid_filter) + except dpctl.SyclContextCreationError: + pytest.skip() + caps = ctx.atomic_memory_scope_capabilities + assert isinstance(caps, tuple) + assert all(isinstance(m, dpctl.memory_scope) for m in caps) + # SYCL 2020 requires at least these capabilities + assert dpctl.memory_scope.work_group in caps + for d in ctx.get_devices(): + assert set(caps).issubset(set(d.atomic_memory_scope_capabilities)) + + +def test_context_atomic_fence_scope_capabilities(valid_filter): + try: + ctx = dpctl.SyclContext(valid_filter) + except dpctl.SyclContextCreationError: + pytest.skip() + caps = ctx.atomic_fence_scope_capabilities + assert isinstance(caps, tuple) + assert all(isinstance(m, dpctl.memory_scope) for m in caps) + for d in ctx.get_devices(): + assert set(caps).issubset(set(d.atomic_fence_scope_capabilities)) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_context_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_context_interface.h index c7c6ece7ed..c19eb024bc 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_context_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_context_interface.h @@ -161,4 +161,73 @@ void DPCTLContext_Delete(__dpctl_take DPCTLSyclContextRef CtxRef); DPCTL_API size_t DPCTLContext_Hash(__dpctl_keep DPCTLSyclContextRef CtxRef); +/*! + * @brief Wrapper over + * context.get_info(). + * + * @param CtxRef Opaque pointer to a ``sycl::context``. + * @return Returns an opaque pointer to the ``sycl::platform`` associated with + * the context. + * @ingroup ContextInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclPlatformRef +DPCTLContext_GetPlatform(__dpctl_keep const DPCTLSyclContextRef CtxRef); + +/*! + * @brief Wrapper over + * context.get_info(). + * + * @param CtxRef Opaque pointer to a ``sycl::context``. + * @param res_len Populated with size of the returned array. + * @return Returns an array of DPCTLMemoryOrderType values. + * @ingroup ContextInterface + */ +DPCTL_API +__dpctl_give int *DPCTLContext_GetAtomicMemoryOrderCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len); + +/*! + * @brief Wrapper over + * context.get_info(). + * + * @param CtxRef Opaque pointer to a ``sycl::context``. + * @param res_len Populated with size of the returned array. + * @return Returns an array of DPCTLMemoryOrderType values. + * @ingroup ContextInterface + */ +DPCTL_API +__dpctl_give int *DPCTLContext_GetAtomicFenceOrderCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len); + +/*! + * @brief Wrapper over + * context.get_info(). + * + * @param CtxRef Opaque pointer to a ``sycl::context``. + * @param res_len Populated with size of the returned array. + * @return Returns an array of DPCTLMemoryScopeType values. + * @ingroup ContextInterface + */ +DPCTL_API +__dpctl_give int *DPCTLContext_GetAtomicMemoryScopeCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len); + +/*! + * @brief Wrapper over + * context.get_info(). + * + * @param CtxRef Opaque pointer to a ``sycl::context``. + * @param res_len Populated with size of the returned array. + * @return Returns an array of DPCTLMemoryScopeType values. + * @ingroup ContextInterface + */ +DPCTL_API +__dpctl_give int *DPCTLContext_GetAtomicFenceScopeCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len); + DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/source/dpctl_sycl_context_interface.cpp b/libsyclinterface/source/dpctl_sycl_context_interface.cpp index c27cd9a343..cffadad953 100644 --- a/libsyclinterface/source/dpctl_sycl_context_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_context_interface.cpp @@ -28,6 +28,7 @@ #include "Config/dpctl_config.h" #include "dpctl_error_handlers.h" #include "dpctl_sycl_type_casters.hpp" +#include "dpctl_utils_helper.h" #include #include #include @@ -212,3 +213,88 @@ size_t DPCTLContext_Hash(__dpctl_keep const DPCTLSyclContextRef CtxRef) return 0; } } + +__dpctl_give DPCTLSyclPlatformRef +DPCTLContext_GetPlatform(__dpctl_keep const DPCTLSyclContextRef CtxRef) +{ + DPCTLSyclPlatformRef PRef = nullptr; + auto C = unwrap(CtxRef); + if (C) { + try { + PRef = wrap( + new platform(C->get_info())); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } + } + return PRef; +} + +namespace +{ + +template +int *get_context_info_enum_array(__dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len, + ConvertFn convert) +{ + int *arr = nullptr; + *res_len = 0; + auto C = unwrap(CtxRef); + if (C) { + try { + auto values = C->get_info(); + *res_len = values.size(); + if (*res_len > 0) { + arr = new int[*res_len]; + for (size_t i = 0; i < *res_len; ++i) { + arr[i] = convert(values[i]); + } + } + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + delete[] arr; + arr = nullptr; + *res_len = 0; + } + } + return arr; +} + +} // end of anonymous namespace + +__dpctl_give int *DPCTLContext_GetAtomicMemoryOrderCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len) +{ + return get_context_info_enum_array< + info::context::atomic_memory_order_capabilities>( + CtxRef, res_len, DPCTL_SyclMemoryOrderToDPCTLType); +} + +__dpctl_give int *DPCTLContext_GetAtomicFenceOrderCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len) +{ + return get_context_info_enum_array< + info::context::atomic_fence_order_capabilities>( + CtxRef, res_len, DPCTL_SyclMemoryOrderToDPCTLType); +} + +__dpctl_give int *DPCTLContext_GetAtomicMemoryScopeCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len) +{ + return get_context_info_enum_array< + info::context::atomic_memory_scope_capabilities>( + CtxRef, res_len, DPCTL_SyclMemoryScopeToDPCTLType); +} + +__dpctl_give int *DPCTLContext_GetAtomicFenceScopeCapabilities( + __dpctl_keep const DPCTLSyclContextRef CtxRef, + size_t *res_len) +{ + return get_context_info_enum_array< + info::context::atomic_fence_scope_capabilities>( + CtxRef, res_len, DPCTL_SyclMemoryScopeToDPCTLType); +} diff --git a/libsyclinterface/tests/test_sycl_context_interface.cpp b/libsyclinterface/tests/test_sycl_context_interface.cpp index 69f01d800a..a395e31b62 100644 --- a/libsyclinterface/tests/test_sycl_context_interface.cpp +++ b/libsyclinterface/tests/test_sycl_context_interface.cpp @@ -28,7 +28,10 @@ #include "dpctl_sycl_context_interface.h" #include "dpctl_sycl_device_interface.h" #include "dpctl_sycl_device_selector_interface.h" +#include "dpctl_sycl_enum_types.h" +#include "dpctl_sycl_platform_interface.h" #include "dpctl_sycl_types.h" +#include "dpctl_utils.h" #include @@ -181,6 +184,100 @@ TEST_P(TestDPCTLContextInterface, ChkGetBackend) EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(CRef)); } +TEST_P(TestDPCTLContextInterface, ChkGetPlatform) +{ + DPCTLSyclContextRef CRef = nullptr; + DPCTLSyclPlatformRef PRef = nullptr; + + EXPECT_NO_FATAL_FAILURE(CRef = DPCTLContext_Create(DRef, nullptr, 0)); + ASSERT_TRUE(CRef); + EXPECT_NO_FATAL_FAILURE(PRef = DPCTLContext_GetPlatform(CRef)); + ASSERT_TRUE(PRef); + + EXPECT_NO_FATAL_FAILURE(DPCTLPlatform_Delete(PRef)); + EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(CRef)); +} + +TEST_P(TestDPCTLContextInterface, ChkGetAtomicMemoryOrderCapabilities) +{ + DPCTLSyclContextRef CRef = nullptr; + int *arr = nullptr; + size_t len = 0; + + EXPECT_NO_FATAL_FAILURE(CRef = DPCTLContext_Create(DRef, nullptr, 0)); + ASSERT_TRUE(CRef); + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicMemoryOrderCapabilities(CRef, &len)); + EXPECT_TRUE(len > 0); + EXPECT_TRUE(arr != nullptr); + for (size_t i = 0; i < len; ++i) { + EXPECT_TRUE(arr[i] >= DPCTL_MEMORY_ORDER_RELAXED && + arr[i] <= DPCTL_MEMORY_ORDER_SEQ_CST); + } + EXPECT_NO_FATAL_FAILURE(DPCTLInt_Array_Delete(arr)); + EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(CRef)); +} + +TEST_P(TestDPCTLContextInterface, ChkGetAtomicFenceOrderCapabilities) +{ + DPCTLSyclContextRef CRef = nullptr; + int *arr = nullptr; + size_t len = 0; + + EXPECT_NO_FATAL_FAILURE(CRef = DPCTLContext_Create(DRef, nullptr, 0)); + ASSERT_TRUE(CRef); + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicFenceOrderCapabilities(CRef, &len)); + EXPECT_TRUE(len > 0); + EXPECT_TRUE(arr != nullptr); + for (size_t i = 0; i < len; ++i) { + EXPECT_TRUE(arr[i] >= DPCTL_MEMORY_ORDER_RELAXED && + arr[i] <= DPCTL_MEMORY_ORDER_SEQ_CST); + } + EXPECT_NO_FATAL_FAILURE(DPCTLInt_Array_Delete(arr)); + EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(CRef)); +} + +TEST_P(TestDPCTLContextInterface, ChkGetAtomicMemoryScopeCapabilities) +{ + DPCTLSyclContextRef CRef = nullptr; + int *arr = nullptr; + size_t len = 0; + + EXPECT_NO_FATAL_FAILURE(CRef = DPCTLContext_Create(DRef, nullptr, 0)); + ASSERT_TRUE(CRef); + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicMemoryScopeCapabilities(CRef, &len)); + EXPECT_TRUE(len > 0); + EXPECT_TRUE(arr != nullptr); + for (size_t i = 0; i < len; ++i) { + EXPECT_TRUE(arr[i] >= DPCTL_MEMORY_SCOPE_WORK_ITEM && + arr[i] <= DPCTL_MEMORY_SCOPE_SYSTEM); + } + EXPECT_NO_FATAL_FAILURE(DPCTLInt_Array_Delete(arr)); + EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(CRef)); +} + +TEST_P(TestDPCTLContextInterface, ChkGetAtomicFenceScopeCapabilities) +{ + DPCTLSyclContextRef CRef = nullptr; + int *arr = nullptr; + size_t len = 0; + + EXPECT_NO_FATAL_FAILURE(CRef = DPCTLContext_Create(DRef, nullptr, 0)); + ASSERT_TRUE(CRef); + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicFenceScopeCapabilities(CRef, &len)); + EXPECT_TRUE(len > 0); + EXPECT_TRUE(arr != nullptr); + for (size_t i = 0; i < len; ++i) { + EXPECT_TRUE(arr[i] >= DPCTL_MEMORY_SCOPE_WORK_ITEM && + arr[i] <= DPCTL_MEMORY_SCOPE_SYSTEM); + } + EXPECT_NO_FATAL_FAILURE(DPCTLInt_Array_Delete(arr)); + EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(CRef)); +} + INSTANTIATE_TEST_SUITE_P(DPCTLContextTests, TestDPCTLContextInterface, ::testing::Values("opencl", @@ -268,3 +365,50 @@ TEST_F(TestDPCTLContextNullArgs, ChkDelete) { EXPECT_NO_FATAL_FAILURE(DPCTLContext_Delete(Null_CRef)); } + +TEST_F(TestDPCTLContextNullArgs, ChkGetPlatform) +{ + DPCTLSyclPlatformRef PRef = nullptr; + EXPECT_NO_FATAL_FAILURE(PRef = DPCTLContext_GetPlatform(Null_CRef)); + ASSERT_FALSE(bool(PRef)); +} + +TEST_F(TestDPCTLContextNullArgs, ChkGetAtomicMemoryOrderCapabilities) +{ + int *arr = nullptr; + size_t len = 42; + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicMemoryOrderCapabilities(Null_CRef, &len)); + ASSERT_FALSE(bool(arr)); + ASSERT_EQ(len, 0ul); +} + +TEST_F(TestDPCTLContextNullArgs, ChkGetAtomicFenceOrderCapabilities) +{ + int *arr = nullptr; + size_t len = 42; + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicFenceOrderCapabilities(Null_CRef, &len)); + ASSERT_FALSE(bool(arr)); + ASSERT_EQ(len, 0ul); +} + +TEST_F(TestDPCTLContextNullArgs, ChkGetAtomicMemoryScopeCapabilities) +{ + int *arr = nullptr; + size_t len = 42; + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicMemoryScopeCapabilities(Null_CRef, &len)); + ASSERT_FALSE(bool(arr)); + ASSERT_EQ(len, 0ul); +} + +TEST_F(TestDPCTLContextNullArgs, ChkGetAtomicFenceScopeCapabilities) +{ + int *arr = nullptr; + size_t len = 42; + EXPECT_NO_FATAL_FAILURE( + arr = DPCTLContext_GetAtomicFenceScopeCapabilities(Null_CRef, &len)); + ASSERT_FALSE(bool(arr)); + ASSERT_EQ(len, 0ul); +} From 5a849fd69829916bcfd2412834c7b2944285f320 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Thu, 6 Aug 2026 07:16:05 -0700 Subject: [PATCH 2/2] add gh-2354 to changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 686d06eed7..65aadd438d 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +9,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added * `dpctl.SyclQueue.copy` and `dpctl.SyclQueue.copy_async` methods [gh-2273](https://github.com/IntelPython/dpctl/pull/2273) * Added a number of `sycl::device` info queries to `dpctl.SyclDevice` [gh-2324](https://github.com/IntelPython/dpctl/pull/2324) +* Added `sycl::info::context` queries `sycl_platform`, `atomic_memory_order_capabilities`, `atomic_fence_order_capabilities`, `atomic_memory_scope_capabilities`, and `atomic_fence_scope_capabilities` to `dpctl.SyclContext` [gh-2354](https://github.com/IntelPython/dpctl/pull/2354) ### Changed * Bump minimum NumPy version to 1.26 [gh-2192](https://github.com/IntelPython/dpctl/pull/2192)