[UR] Add device info query for native assert.

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<UrApiKind::urDeviceGetInfo>(
+      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<info::device::ext_oneapi_srgb>();
   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,

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{};

sycl/unittests/assert/CMakeLists.txt

@@ -1,4 +1,5 @@
 add_sycl_unittest(AssertTests OBJECT
   assert.cpp
+  support_native.cpp
 )
 

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 <sycl/sycl.hpp>
+
+#include <helpers/UrMock.hpp>
+
+#include <gtest/gtest.h>
+
+template <bool Support>
+static ur_result_t redefinedDeviceGetInfoAfter(void *pParams) {
+  auto &Params = *reinterpret_cast<ur_device_get_info_params_t *>(pParams);
+  if (*Params.ppropName == UR_DEVICE_INFO_USE_NATIVE_ASSERT) {
+    if (*Params.ppPropValue)
+      *reinterpret_cast<ur_bool_t *>(*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<true>);
+
+  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<false>);
+
+  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));
+}

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,

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) {

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"

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:

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:

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:

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);

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<cl_device_id>(hDevice), {"cl_intel_devicelib_assert"},
+        Supported));
+    return ReturnValue(Supported);
+  }
   case UR_DEVICE_INFO_EXTENSIONS: {
     CL_RETURN_ON_FAILURE(clGetDeviceInfo(
         cl_adapter::cast<cl_device_id>(hDevice), CL_DEVICE_EXTENSIONS, propSize,

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,

unified-runtime/tools/urinfo/urinfo.hpp

@@ -332,6 +332,8 @@ inline void printDeviceInfos(ur_device_handle_t hDevice,
   printDeviceInfo<ur_bool_t>(
       hDevice, UR_DEVICE_INFO_PROGRAM_SET_SPECIALIZATION_CONSTANTS);
   std::cout << prefix;
+  printDeviceInfo<ur_bool_t>(hDevice, UR_DEVICE_INFO_USE_NATIVE_ASSERT);
+  std::cout << prefix;
   printDeviceInfo<ur_bool_t>(hDevice,
                              UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP);
   std::cout << prefix;