From cedd60211c5f8018aaed3457de04db5fa6397e3b Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Wed, 30 Oct 2024 17:15:45 +0000 Subject: [PATCH 1/5] [UR] Add device info query for devicelib assert. --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 2 +- sycl/cmake/modules/UnifiedRuntimeTag.cmake | 8 +------- sycl/source/detail/device_impl.cpp | 11 +++++------ sycl/source/detail/device_impl.hpp | 9 +++++++-- 4 files changed, 14 insertions(+), 16 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 72841724fa01d..d0eff6cbe84f1 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/aarongreig/unified-runtime.git") include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index 6f3f57255c75b..5cbeaba58a60a 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -1,7 +1 @@ -# commit 3d58884b4939d9bd095c917f8dd823ac8486684c -# Merge: 6ade245e b0bd146a -# Author: aarongreig -# Date: Thu Oct 31 14:05:55 2024 +0000 -# Merge pull request #2228 from nrspruit/copy_engine_refactor -# [L0] Refactor Copy Engine Usage checks for Performance -set(UNIFIED_RUNTIME_TAG 3d58884b4939d9bd095c917f8dd823ac8486684c) +set(UNIFIED_RUNTIME_TAG aaron/addAssertDeviceInfoQuery) diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 178634322f47e..dec1832c2965b 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -73,8 +73,9 @@ device_impl::device_impl(ur_native_handle_t InteropDeviceHandle, } MPlatform = Platform; - MIsAssertFailSupported = - has_extension(UR_DEVICE_INFO_EXTENSION_DEVICELIB_ASSERT); + Adapter->call( + MDevice, UR_DEVICE_INFO_USE_NATIVE_ASSERT, sizeof(ur_bool_t), + &MUseNativeAssert, nullptr); } device_impl::~device_impl() { @@ -472,7 +473,7 @@ bool device_impl::has(aspect Aspect) const { case aspect::ext_oneapi_srgb: return get_info(); case aspect::ext_oneapi_native_assert: - return isAssertFailSupported(); + return useNativeAssert(); case aspect::ext_oneapi_cuda_async_barrier: { int async_barrier_supported; bool call_successful = @@ -782,9 +783,7 @@ bool device_impl::has(aspect Aspect) const { return false; // This device aspect has not been implemented yet. } -bool device_impl::isAssertFailSupported() const { - return MIsAssertFailSupported; -} +bool device_impl::useNativeAssert() const { return MUseNativeAssert; } std::string device_impl::getDeviceName() const { std::call_once(MDeviceNameFlag, diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index b38b7582f3f28..2f45a0b288469 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -209,7 +209,12 @@ class device_impl { /// \return true if the SYCL device has the given feature. bool has(aspect Aspect) const; - bool isAssertFailSupported() const; + /// Indicates the SYCL device prefers to use its native assert + /// implementation. + /// + /// If this is false we will use the fallback assert implementation, + /// as detailed in doc/design/Assert.md + bool useNativeAssert() const; bool isRootDevice() const { return MRootDevice == nullptr; } @@ -302,7 +307,7 @@ class device_impl { ur_device_type_t MType; ur_device_handle_t MRootDevice = nullptr; PlatformImplPtr MPlatform; - bool MIsAssertFailSupported = false; + bool MUseNativeAssert = false; mutable std::string MDeviceName; mutable std::once_flag MDeviceNameFlag; mutable ext::oneapi::experimental::architecture MDeviceArch{}; From f124b82971cf0762600e91d129ff5cb2387fd323 Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Fri, 8 Nov 2024 11:33:19 +0000 Subject: [PATCH 2/5] Add unit test to check the aspect is correctly reported. --- sycl/unittests/assert/CMakeLists.txt | 1 + sycl/unittests/assert/support_native.cpp | 51 ++++++++++++++++++++++++ 2 files changed, 52 insertions(+) create mode 100644 sycl/unittests/assert/support_native.cpp diff --git a/sycl/unittests/assert/CMakeLists.txt b/sycl/unittests/assert/CMakeLists.txt index ebed0c40378ea..9e114f36441f4 100644 --- a/sycl/unittests/assert/CMakeLists.txt +++ b/sycl/unittests/assert/CMakeLists.txt @@ -1,4 +1,5 @@ add_sycl_unittest(AssertTests OBJECT assert.cpp + support_native.cpp ) diff --git a/sycl/unittests/assert/support_native.cpp b/sycl/unittests/assert/support_native.cpp new file mode 100644 index 0000000000000..00660ffbd65cf --- /dev/null +++ b/sycl/unittests/assert/support_native.cpp @@ -0,0 +1,51 @@ +//==---------- support_native.cpp --- Check support is correctly reported --==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "ur_mock_helpers.hpp" + +#include + +#include + +#include + +template +static ur_result_t redefinedDeviceGetInfoAfter(void *pParams) { + auto &Params = *reinterpret_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_USE_NATIVE_ASSERT) { + if (*Params.ppPropValue) + *reinterpret_cast(*Params.ppPropValue) = Support; + if (*Params.ppPropSizeRet) + **Params.ppPropSizeRet = sizeof(ur_bool_t); + } + return UR_RESULT_SUCCESS; +} + +TEST(SupportNativeAssert, True) { + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfoAfter); + + sycl::unittest::UrMock<> Mock; + sycl::platform Plt = sycl::platform(); + + const sycl::device Dev = Plt.get_devices()[0]; + + ASSERT_TRUE(Dev.has(sycl::aspect::ext_oneapi_native_assert)); +} + +TEST(SupportNativeAssert, False) { + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfoAfter); + + sycl::unittest::UrMock<> Mock; + sycl::platform Plt = sycl::platform(); + + const sycl::device Dev = Plt.get_devices()[0]; + + ASSERT_FALSE(Dev.has(sycl::aspect::ext_oneapi_native_assert)); +} From efd29e7e6208e8b236d1d011623fadbeefa36616 Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Wed, 30 Oct 2024 16:51:21 +0000 Subject: [PATCH 3/5] Add device info query to report support for native asserts. This allows cuda and hip to stop reporting the relevant opencl extension string, see issue #1374 --- unified-runtime/include/ur_api.h | 3 +++ unified-runtime/include/ur_print.hpp | 16 ++++++++++++++++ unified-runtime/scripts/core/device.yml | 2 ++ unified-runtime/source/adapters/cuda/device.cpp | 4 +++- unified-runtime/source/adapters/hip/device.cpp | 7 ++----- .../source/adapters/level_zero/device.cpp | 2 ++ .../source/adapters/native_cpu/device.cpp | 2 ++ .../source/adapters/opencl/device.cpp | 7 +++++++ .../test/conformance/device/urDeviceGetInfo.cpp | 6 ++++-- unified-runtime/tools/urinfo/urinfo.hpp | 2 ++ 10 files changed, 43 insertions(+), 8 deletions(-) diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index c390ed4410d16..056fe5b49917d 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -2194,6 +2194,9 @@ typedef enum ur_device_info_t { /// [::ur_bool_t] support the ::urProgramSetSpecializationConstants entry /// point UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS = 121, + /// [::ur_bool_t] return true if the device has a native assert + /// implementation. + UR_DEVICE_INFO_USE_NATIVE_ASSERT = 122, /// [::ur_bool_t] Returns true if the device supports the use of /// command-buffers. UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP = 0x1000, diff --git a/unified-runtime/include/ur_print.hpp b/unified-runtime/include/ur_print.hpp index 5c5f573477929..71728dbd1996a 100644 --- a/unified-runtime/include/ur_print.hpp +++ b/unified-runtime/include/ur_print.hpp @@ -2876,6 +2876,9 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_device_info_t value) { case UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS: os << "UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS"; break; + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: + os << "UR_DEVICE_INFO_USE_NATIVE_ASSERT"; + break; case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: os << "UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP"; break; @@ -4533,6 +4536,19 @@ inline ur_result_t printTagged(std::ostream &os, const void *ptr, os << ")"; } break; + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: { + const ur_bool_t *tptr = (const ur_bool_t *)ptr; + if (sizeof(ur_bool_t) > size) { + os << "invalid size (is: " << size + << ", expected: >=" << sizeof(ur_bool_t) << ")"; + return UR_RESULT_ERROR_INVALID_SIZE; + } + os << (const void *)(tptr) << " ("; + + os << *tptr; + + os << ")"; + } break; case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: { const ur_bool_t *tptr = (const ur_bool_t *)ptr; if (sizeof(ur_bool_t) > size) { diff --git a/unified-runtime/scripts/core/device.yml b/unified-runtime/scripts/core/device.yml index 1f94b20346042..6548f5e1b9819 100644 --- a/unified-runtime/scripts/core/device.yml +++ b/unified-runtime/scripts/core/device.yml @@ -447,6 +447,8 @@ etors: desc: "[uint32_t] the number of compute units for specific backend." - name: PROGRAM_SET_SPECIALIZATION_CONSTANTS desc: "[$x_bool_t] support the $xProgramSetSpecializationConstants entry point" + - name: USE_NATIVE_ASSERT + desc: "[$x_bool_t] return true if the device has a native assert implementation." --- #-------------------------------------------------------------------------- type: function desc: "Retrieves various information about device" diff --git a/unified-runtime/source/adapters/cuda/device.cpp b/unified-runtime/source/adapters/cuda/device.cpp index 3e0ce05c27306..2b14e0446e007 100644 --- a/unified-runtime/source/adapters/cuda/device.cpp +++ b/unified-runtime/source/adapters/cuda/device.cpp @@ -617,7 +617,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_EXTENSIONS: { std::string SupportedExtensions = "cl_khr_fp64 cl_khr_subgroups "; - SupportedExtensions += "cl_intel_devicelib_assert "; // Return supported for the UR command-buffer experimental feature SupportedExtensions += "ur_exp_command_buffer "; SupportedExtensions += "ur_exp_usm_p2p "; @@ -1114,6 +1113,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_LOW_POWER_EVENTS_EXP: return ReturnValue(false); + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: + return ReturnValue(true); + default: break; } diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index 783f4899b9f23..66f318a489162 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -558,12 +558,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(""); } case UR_DEVICE_INFO_EXTENSIONS: { - // TODO: Remove comment when HIP support native asserts. - // DEVICELIB_ASSERT extension is set so fallback assert - // postprocessing is NOP. HIP 4.3 docs indicate support for - // native asserts are in progress std::string SupportedExtensions = ""; - SupportedExtensions += "cl_intel_devicelib_assert "; SupportedExtensions += "ur_exp_usm_p2p "; int RuntimeVersion = 0; @@ -1107,6 +1102,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_LOW_POWER_EVENTS_EXP: { return ReturnValue(false); } + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: + return ReturnValue(true); default: break; } diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index 0d5323f6c3b03..d65bfe402366a 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1210,6 +1210,8 @@ ur_result_t urDeviceGetInfo( return ReturnValue(false); case UR_DEVICE_INFO_HOST_PIPE_READ_WRITE_SUPPORTED: return ReturnValue(false); + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: + return ReturnValue(false); default: logger::error("Unsupported ParamName in urGetDeviceInfo"); logger::error("ParamNameParamName={}(0x{})", ParamName, diff --git a/unified-runtime/source/adapters/native_cpu/device.cpp b/unified-runtime/source/adapters/native_cpu/device.cpp index 6deca1ac37ac2..bf2e035191aae 100644 --- a/unified-runtime/source/adapters/native_cpu/device.cpp +++ b/unified-runtime/source/adapters/native_cpu/device.cpp @@ -426,6 +426,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_USM_POOL_SUPPORT: return ReturnValue(false); + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: + return ReturnValue(false); case UR_DEVICE_INFO_LOW_POWER_EVENTS_EXP: return ReturnValue(false); diff --git a/unified-runtime/source/adapters/opencl/device.cpp b/unified-runtime/source/adapters/opencl/device.cpp index 1267b2e94177a..bad1e9281892b 100644 --- a/unified-runtime/source/adapters/opencl/device.cpp +++ b/unified-runtime/source/adapters/opencl/device.cpp @@ -1227,6 +1227,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(false); case UR_DEVICE_INFO_LOW_POWER_EVENTS_EXP: return ReturnValue(false); + case UR_DEVICE_INFO_USE_NATIVE_ASSERT: { + bool Supported = false; + UR_RETURN_ON_FAILURE(cl_adapter::checkDeviceExtensions( + cl_adapter::cast(hDevice), {"cl_intel_devicelib_assert"}, + Supported)); + return ReturnValue(Supported); + } default: { return UR_RESULT_ERROR_INVALID_ENUMERATION; } diff --git a/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp b/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp index 99ff34ad2a6b6..7310abe108c41 100644 --- a/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp +++ b/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp @@ -129,7 +129,8 @@ static std::unordered_map device_info_size_map = { {UR_DEVICE_INFO_ESIMD_SUPPORT, sizeof(ur_bool_t)}, {UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t)}, {UR_DEVICE_INFO_VIRTUAL_MEMORY_SUPPORT, sizeof(ur_bool_t)}, - {UR_DEVICE_INFO_NUM_COMPUTE_UNITS, sizeof(uint32_t)}}; + {UR_DEVICE_INFO_NUM_COMPUTE_UNITS, sizeof(uint32_t)}, + {UR_DEVICE_INFO_USE_NATIVE_ASSERT, sizeof(ur_bool_t)}}; using urDeviceGetInfoTest = uur::urDeviceTestWithParam; @@ -255,7 +256,8 @@ UUR_DEVICE_TEST_SUITE_WITH_PARAM( UR_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF, // UR_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT, // UR_DEVICE_INFO_NUM_COMPUTE_UNITS, // - UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS // + UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS, // + UR_DEVICE_INFO_USE_NATIVE_ASSERT // ), uur::deviceTestWithParamPrinter); diff --git a/unified-runtime/tools/urinfo/urinfo.hpp b/unified-runtime/tools/urinfo/urinfo.hpp index d01245138f28e..048a0f7b2a289 100644 --- a/unified-runtime/tools/urinfo/urinfo.hpp +++ b/unified-runtime/tools/urinfo/urinfo.hpp @@ -332,6 +332,8 @@ inline void printDeviceInfos(ur_device_handle_t hDevice, printDeviceInfo( hDevice, UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS); std::cout << prefix; + printDeviceInfo(hDevice, UR_DEVICE_INFO_USE_NATIVE_ASSERT); + std::cout << prefix; printDeviceInfo(hDevice, UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP); std::cout << prefix; From 4005ecd6db395df201fdae372ad927a41abcda00 Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Fri, 7 Mar 2025 13:06:44 +0000 Subject: [PATCH 4/5] Fix dodgy conflict resolution. --- unified-runtime/source/adapters/hip/device.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index a0b95840758e5..dc62f6d9d0aac 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -1075,10 +1075,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_LOW_POWER_EVENTS_EXP: { return ReturnValue(false); } -<<<<<<< HEAD case UR_DEVICE_INFO_USE_NATIVE_ASSERT: return ReturnValue(true); -======= case UR_DEVICE_INFO_USM_P2P_SUPPORT_EXP: return ReturnValue(true); case UR_DEVICE_INFO_LAUNCH_PROPERTIES_SUPPORT_EXP: @@ -1087,7 +1085,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(true); case UR_DEVICE_INFO_MULTI_DEVICE_COMPILE_SUPPORT_EXP: return ReturnValue(false); ->>>>>>> sycl default: break; } From 14c61f7c89d7039c3bc7794e76fb006738c01677 Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Mon, 10 Mar 2025 11:30:24 +0000 Subject: [PATCH 5/5] Fix test --- unified-runtime/test/conformance/device/urDeviceGetInfo.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp b/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp index 76aafb792934b..038543d01714a 100644 --- a/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp +++ b/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp @@ -2568,7 +2568,7 @@ TEST_P(urDeviceGetInfoTest, SuccessUseNativeAssert) { ASSERT_SUCCESS_OR_OPTIONAL_QUERY( urDeviceGetInfo(device, property_name, 0, nullptr, &property_size), property_name); - ASSERT_EQ(property_size, sizeof(uint32_t)); + ASSERT_EQ(property_size, sizeof(ur_bool_t)); uint32_t property_value = 0; ASSERT_QUERY_RETURNS_VALUE(urDeviceGetInfo(device, property_name,