From 6ea098f3458e481d1a0c5254d4ebde1df088443f Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Fri, 24 Mar 2023 10:52:49 -0700 Subject: [PATCH 1/2] Adds atomic_memory_scope context query. Signed-off-by: Maronas, Marcos --- sycl/source/detail/context_impl.cpp | 31 +++++++---- .../AtomicMemoryOrderCapabilities.cpp | 52 ++++++++++++++++-- .../AtomicMemoryScopeCapabilities.cpp | 54 +++++++++++++++++-- 3 files changed, 117 insertions(+), 20 deletions(-) diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index 198b4f1cc2d30..1a840a3ccfee0 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -176,7 +176,7 @@ context_impl::get_info() return CapabilityList; for (const sycl::device &Device : MDevices) { - std::vector NewCapabilityList(CapabilityList.size()); + std::vector NewCapabilityList; std::vector DeviceCapabilities = Device.get_info(); std::set_intersection( @@ -193,17 +193,26 @@ template <> std::vector context_impl::get_info() const { + std::vector CapabilityList{ + sycl::memory_scope::work_item, sycl::memory_scope::sub_group, + sycl::memory_scope::work_group, sycl::memory_scope::device, + sycl::memory_scope::system}; if (is_host()) - return {sycl::memory_scope::work_item, sycl::memory_scope::sub_group, - sycl::memory_scope::work_group, sycl::memory_scope::device, - sycl::memory_scope::system}; - - pi_memory_scope_capabilities Result; - getPlugin().call( - MContext, - PiInfoCode::value, - sizeof(Result), &Result, nullptr); - return readMemoryScopeBitfield(Result); + return CapabilityList; + + for (const sycl::device &Device : MDevices) { + std::vector NewCapabilityList; + std::vector DeviceCapabilities = + Device.get_info(); + std::set_intersection( + CapabilityList.begin(), CapabilityList.end(), + DeviceCapabilities.begin(), DeviceCapabilities.end(), + std::inserter(NewCapabilityList, NewCapabilityList.begin())); + CapabilityList = NewCapabilityList; + } + CapabilityList.shrink_to_fit(); + + return CapabilityList; } RT::PiContext &context_impl::getHandleRef() { return MContext; } diff --git a/sycl/unittests/SYCL2020/AtomicMemoryOrderCapabilities.cpp b/sycl/unittests/SYCL2020/AtomicMemoryOrderCapabilities.cpp index 6ed291c0da3e8..c320730d50fab 100644 --- a/sycl/unittests/SYCL2020/AtomicMemoryOrderCapabilities.cpp +++ b/sycl/unittests/SYCL2020/AtomicMemoryOrderCapabilities.cpp @@ -15,8 +15,7 @@ using namespace sycl; namespace { -static constexpr size_t expectedCapabilityVecSize = 5; -static thread_local bool deviceGetInfoCalled = false; +thread_local bool deviceGetInfoCalled; static bool has_capability(const std::vector &deviceCapabilities, memory_order capabilityToFind) { @@ -24,6 +23,18 @@ static bool has_capability(const std::vector &deviceCapabilities, capabilityToFind) != deviceCapabilities.end(); } +pi_result redefinedDevicesGet(pi_platform platform, pi_device_type device_type, + pi_uint32 num_entries, pi_device *devices, + pi_uint32 *num_devices) { + if (num_devices) + *num_devices = 2; + if (devices && num_entries > 0) { + devices[0] = reinterpret_cast(1); + devices[1] = reinterpret_cast(2); + } + return PI_SUCCESS; +} + pi_result redefinedDeviceGetInfo(pi_device device, pi_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret) { @@ -32,9 +43,14 @@ pi_result redefinedDeviceGetInfo(pi_device device, pi_device_info param_name, if (param_value) { pi_memory_order_capabilities *Capabilities = reinterpret_cast(param_value); - *Capabilities = PI_MEMORY_ORDER_RELAXED | PI_MEMORY_ORDER_ACQUIRE | - PI_MEMORY_ORDER_RELEASE | PI_MEMORY_ORDER_ACQ_REL | - PI_MEMORY_ORDER_SEQ_CST; + if (device == reinterpret_cast(1)) { + *Capabilities = PI_MEMORY_ORDER_RELAXED | PI_MEMORY_ORDER_ACQUIRE | + PI_MEMORY_ORDER_RELEASE | PI_MEMORY_ORDER_ACQ_REL | + PI_MEMORY_ORDER_SEQ_CST; + } + if (device == reinterpret_cast(2)) { + *Capabilities = PI_MEMORY_ORDER_RELAXED | PI_MEMORY_ORDER_SEQ_CST; + } } } return PI_SUCCESS; @@ -50,9 +66,12 @@ TEST(AtomicMemoryOrderCapabilities, DeviceQueryReturnsCorrectCapabilities) { const device Dev = Plt.get_devices()[0]; context Ctx{Dev}; + deviceGetInfoCalled = false; + auto Capabilities = Dev.get_info(); EXPECT_TRUE(deviceGetInfoCalled); + constexpr size_t expectedCapabilityVecSize = 5; EXPECT_EQ(Capabilities.size(), expectedCapabilityVecSize); EXPECT_TRUE(has_capability(Capabilities, memory_order::relaxed)); @@ -62,4 +81,27 @@ TEST(AtomicMemoryOrderCapabilities, DeviceQueryReturnsCorrectCapabilities) { EXPECT_TRUE(has_capability(Capabilities, memory_order::seq_cst)); } +TEST(AtomicMemoryOrderCapabilities, ContextQueryReturnsCorrectCapabilities) { + unittest::PiMock Mock; + platform Plt = Mock.getPlatform(); + + Mock.redefineAfter( + redefinedDeviceGetInfo); + Mock.redefineAfter(redefinedDevicesGet); + + auto devices = Plt.get_devices(); + context Ctx{devices}; + + deviceGetInfoCalled = false; + + auto Capabilities = + Ctx.get_info(); + EXPECT_TRUE(deviceGetInfoCalled); + constexpr size_t expectedCapabilityVecSize = 2; + EXPECT_EQ(Capabilities.size(), expectedCapabilityVecSize); + + EXPECT_TRUE(has_capability(Capabilities, memory_order::relaxed)); + EXPECT_TRUE(has_capability(Capabilities, memory_order::seq_cst)); +} + } // namespace diff --git a/sycl/unittests/SYCL2020/AtomicMemoryScopeCapabilities.cpp b/sycl/unittests/SYCL2020/AtomicMemoryScopeCapabilities.cpp index 776991b9e53bf..d862add5c2466 100644 --- a/sycl/unittests/SYCL2020/AtomicMemoryScopeCapabilities.cpp +++ b/sycl/unittests/SYCL2020/AtomicMemoryScopeCapabilities.cpp @@ -16,6 +16,18 @@ namespace { thread_local bool deviceGetInfoCalled; +pi_result redefinedDevicesGet(pi_platform platform, pi_device_type device_type, + pi_uint32 num_entries, pi_device *devices, + pi_uint32 *num_devices) { + if (num_devices) + *num_devices = 2; + if (devices && num_entries > 0) { + devices[0] = reinterpret_cast(1); + devices[1] = reinterpret_cast(2); + } + return PI_SUCCESS; +} + pi_result redefinedDeviceGetInfoAfter(pi_device device, pi_device_info param_name, size_t param_value_size, @@ -26,15 +38,21 @@ pi_result redefinedDeviceGetInfoAfter(pi_device device, if (param_value) { auto *Result = reinterpret_cast(param_value); - *Result = PI_MEMORY_SCOPE_WORK_ITEM | PI_MEMORY_SCOPE_SUB_GROUP | - PI_MEMORY_SCOPE_WORK_GROUP | PI_MEMORY_SCOPE_DEVICE | - PI_MEMORY_SCOPE_SYSTEM; + if (device == reinterpret_cast(1)) { + *Result = PI_MEMORY_SCOPE_WORK_ITEM | PI_MEMORY_SCOPE_SUB_GROUP | + PI_MEMORY_SCOPE_WORK_GROUP | PI_MEMORY_SCOPE_DEVICE | + PI_MEMORY_SCOPE_SYSTEM; + } + if (device == reinterpret_cast(2)) { + *Result = PI_MEMORY_SCOPE_WORK_ITEM | PI_MEMORY_SCOPE_SYSTEM; + } } } return PI_SUCCESS; } -TEST(AtomicMemoryScopeCapabilitiesCheck, CheckAtomicMemoryScopeCapabilities) { +TEST(AtomicMemoryScopeCapabilitiesCheck, + CheckDeviceAtomicMemoryScopeCapabilities) { sycl::unittest::PiMock Mock; sycl::platform Plt = Mock.getPlatform(); device Dev = Plt.get_devices()[0]; @@ -65,4 +83,32 @@ TEST(AtomicMemoryScopeCapabilitiesCheck, CheckAtomicMemoryScopeCapabilities) { sycl::memory_scope::system); EXPECT_FALSE(res == scope_capabilities.end()); } + +TEST(AtomicMemoryScopeCapabilitiesCheck, + CheckContextAtomicMemoryScopeCapabilities) { + sycl::unittest::PiMock Mock; + sycl::platform Plt = Mock.getPlatform(); + + Mock.redefineAfter( + redefinedDeviceGetInfoAfter); + Mock.redefineAfter(redefinedDevicesGet); + + auto devices = Plt.get_devices(); + context Ctx{devices}; + + deviceGetInfoCalled = false; + + auto scope_capabilities = + Ctx.get_info(); + EXPECT_TRUE(deviceGetInfoCalled); + size_t expectedSize = 2; + EXPECT_EQ(scope_capabilities.size(), expectedSize); + + auto res = std::find(scope_capabilities.begin(), scope_capabilities.end(), + sycl::memory_scope::work_item); + EXPECT_FALSE(res == scope_capabilities.end()); + res = std::find(scope_capabilities.begin(), scope_capabilities.end(), + sycl::memory_scope::system); + EXPECT_FALSE(res == scope_capabilities.end()); +} } // anonymous namespace From 2bd78b10e8badfd99efdce68051e982c8e4bcbee Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Mon, 27 Mar 2023 02:50:48 -0700 Subject: [PATCH 2/2] Address code review comments. Signed-off-by: Maronas, Marcos --- sycl/source/detail/context_impl.cpp | 28 ++++++---------------------- sycl/source/detail/context_impl.hpp | 15 +++++++++++++++ 2 files changed, 21 insertions(+), 22 deletions(-) diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index 1a840a3ccfee0..db251a2c95e93 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -175,17 +175,9 @@ context_impl::get_info() if (is_host()) return CapabilityList; - for (const sycl::device &Device : MDevices) { - std::vector NewCapabilityList; - std::vector DeviceCapabilities = - Device.get_info(); - std::set_intersection( - CapabilityList.begin(), CapabilityList.end(), - DeviceCapabilities.begin(), DeviceCapabilities.end(), - std::inserter(NewCapabilityList, NewCapabilityList.begin())); - CapabilityList = NewCapabilityList; - } - CapabilityList.shrink_to_fit(); + GetCapabilitiesIntersectionSet< + sycl::memory_order, info::device::atomic_memory_order_capabilities>( + MDevices, CapabilityList); return CapabilityList; } @@ -200,17 +192,9 @@ context_impl::get_info() if (is_host()) return CapabilityList; - for (const sycl::device &Device : MDevices) { - std::vector NewCapabilityList; - std::vector DeviceCapabilities = - Device.get_info(); - std::set_intersection( - CapabilityList.begin(), CapabilityList.end(), - DeviceCapabilities.begin(), DeviceCapabilities.end(), - std::inserter(NewCapabilityList, NewCapabilityList.begin())); - CapabilityList = NewCapabilityList; - } - CapabilityList.shrink_to_fit(); + GetCapabilitiesIntersectionSet< + sycl::memory_scope, info::device::atomic_memory_scope_capabilities>( + MDevices, CapabilityList); return CapabilityList; } diff --git a/sycl/source/detail/context_impl.hpp b/sycl/source/detail/context_impl.hpp index 4aa86ecb84326..ba6401a7da54c 100644 --- a/sycl/source/detail/context_impl.hpp +++ b/sycl/source/detail/context_impl.hpp @@ -270,6 +270,21 @@ class context_impl { std::mutex MDeviceGlobalInitializersMutex; }; +template +void GetCapabilitiesIntersectionSet(const std::vector &Devices, + std::vector &CapabilityList) { + for (const sycl::device &Device : Devices) { + std::vector NewCapabilityList; + std::vector DeviceCapabilities = Device.get_info(); + std::set_intersection( + CapabilityList.begin(), CapabilityList.end(), + DeviceCapabilities.begin(), DeviceCapabilities.end(), + std::inserter(NewCapabilityList, NewCapabilityList.begin())); + CapabilityList = NewCapabilityList; + } + CapabilityList.shrink_to_fit(); +} + } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl