diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 9d630d20dbfeb..447fbfa4e2b54 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() { @@ -478,7 +479,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 = @@ -796,9 +797,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{}; 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)); +} diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index 6c2021069fa96..fa8b0af75f7ca 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -2222,6 +2222,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 04e1169d140cd..137614cd6f779 100644 --- a/unified-runtime/include/ur_print.hpp +++ b/unified-runtime/include/ur_print.hpp @@ -2930,6 +2930,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; @@ -4602,6 +4605,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 49777e3bd6905..4e613e0c923c7 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 c902df4bb3b0a..a4d33b989a2f8 100644 --- a/unified-runtime/source/adapters/cuda/device.cpp +++ b/unified-runtime/source/adapters/cuda/device.cpp @@ -615,9 +615,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(""); } case UR_DEVICE_INFO_EXTENSIONS: { - std::string SupportedExtensions = "cl_khr_fp64 "; - SupportedExtensions += "cl_intel_devicelib_assert "; int Major = 0; int Minor = 0; @@ -1113,6 +1111,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); case UR_DEVICE_INFO_USM_P2P_SUPPORT_EXP: return ReturnValue(true); case UR_DEVICE_INFO_LAUNCH_PROPERTIES_SUPPORT_EXP: diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index 3a9b087a4c1ac..dc62f6d9d0aac 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -546,12 +546,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 "; hipDeviceProp_t Props; detail::ur::assertion(hipGetDeviceProperties(&Props, hDevice->get()) == @@ -1080,6 +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); } + 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: diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index 04fa911b9ed73..e4e3a55a46384 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1203,6 +1203,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); case UR_DEVICE_INFO_USM_P2P_SUPPORT_EXP: return ReturnValue(true); case UR_DEVICE_INFO_LAUNCH_PROPERTIES_SUPPORT_EXP: diff --git a/unified-runtime/source/adapters/native_cpu/device.cpp b/unified-runtime/source/adapters/native_cpu/device.cpp index 177593b721558..397b30badbb20 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 47c648715651a..e06fc7ed06f26 100644 --- a/unified-runtime/source/adapters/opencl/device.cpp +++ b/unified-runtime/source/adapters/opencl/device.cpp @@ -1541,6 +1541,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue( ur::cl::getAdapter()->clSetProgramSpecializationConstant != nullptr); } + 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); + } case UR_DEVICE_INFO_EXTENSIONS: { CL_RETURN_ON_FAILURE(clGetDeviceInfo( cl_adapter::cast(hDevice), CL_DEVICE_EXTENSIONS, propSize, diff --git a/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp b/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp index 9834ab1bf459b..038543d01714a 100644 --- a/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp +++ b/unified-runtime/test/conformance/device/urDeviceGetInfo.cpp @@ -2561,6 +2561,22 @@ TEST_P(urDeviceGetInfoTest, Success2DBlockArrayCapabilities) { 0); } +TEST_P(urDeviceGetInfoTest, SuccessUseNativeAssert) { + size_t property_size = 0; + const ur_device_info_t property_name = UR_DEVICE_INFO_USE_NATIVE_ASSERT; + + ASSERT_SUCCESS_OR_OPTIONAL_QUERY( + urDeviceGetInfo(device, property_name, 0, nullptr, &property_size), + property_name); + ASSERT_EQ(property_size, sizeof(ur_bool_t)); + + uint32_t property_value = 0; + ASSERT_QUERY_RETURNS_VALUE(urDeviceGetInfo(device, property_name, + property_size, &property_value, + nullptr), + property_value); +} + TEST_P(urDeviceGetInfoTest, InvalidNullHandleDevice) { ur_device_type_t device_type; ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, diff --git a/unified-runtime/tools/urinfo/urinfo.hpp b/unified-runtime/tools/urinfo/urinfo.hpp index b00884fa9b9b8..90a1ea2350616 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;