[SYCL] Remove fallback assertions (#18310)

This commit removes the fallback assertion implementation for SYCL
kernels. The expected behavior after this is that backends that do not
support native asserts, as reported through the ext_oneapi_native_assert
aspect, will ignore assertions in kernel code.

This commit also makes the L0 adapter report support for native assert,
as it was incorrectly reporting false before.

---------

Signed-off-by: Larsen, Steffen <[email protected]>
Co-authored-by: Alexey Bader <[email protected]>

Author: steffenlarsen

devops/compat_ci_exclude.sycl-rel-6_2

@@ -28,6 +28,13 @@ KernelCompiler/opencl.cpp
 KernelCompiler/opencl_cache_eviction.cpp
 KernelCompiler/opencl_queries.cpp
 
+# https://github.com/intel/llvm/pull/18310 removed fallback assertions, but due
+# to a bug in the L0 drivers, one test causes the implementation to continue
+# after assertions are triggered. We allow this regression while the drivers bug
+# gets addressed.
+# See GSD-11097.
+Assert/assert_in_kernels.cpp
+
 # Likely OK, but need author to provide justification, get approval/confirmation
 # from someone:
 

devops/compat_ci_exclude.sycl-rel-6_3

@@ -2,6 +2,13 @@
 Matrix/joint_matrix_bf16_fill_k_cache_arg_dim.cpp
 Matrix/joint_matrix_bf16_fill_k_cache_runtime_dim.cpp
 
+# https://github.com/intel/llvm/pull/18310 removed fallback assertions, but due
+# to a bug in the L0 drivers, one test causes the implementation to continue
+# after assertions are triggered. We allow this regression while the drivers bug
+# gets addressed.
+# See GSD-11097.
+Assert/assert_in_kernels.cpp
+
 # Likely OK, but need author to provide justification, get approval/confirmation
 # from someone:
 

sycl/doc/PreprocessorMacros.md

@@ -39,20 +39,6 @@ This file describes macros that have effect on SYCL compiler and run-time.
   Disable warning diagnostic issued when including `<sycl/sycl.hpp>` without
   `-fsycl` compiler flag.
 
-- **SYCL_FALLBACK_ASSERT**
-
-  Defining as non-zero enables the fallback assert feature even on devices
-  without native support. Be aware that this will add some overhead that is
-  associated with submitting kernels that call `assert()`. When this macro is
-  defined as 0 or is not defined, the logic for detecting assertion failures in kernels is
-  disabled, so a failed assert will not cause a message to be printed and will
-  not cause the program to abort. Some devices have native support for
-  assertions. The logic for detecting assertion failures is always enabled on
-  these devices regardless of whether this macro is defined because that logic
-  does not add any extra overhead. One can check to see if a device has native
-  support for `assert()` via `aspect::ext_oneapi_native_assert`.
-  This macro is undefined by default.
-
 - **SYCL2020_CONFORMANT_APIS (deprecated)**
   This macro is used to comply with the SYCL 2020 specification, as some of the current 
   implementations may be widespread and not conform to it.

sycl/doc/design/Assert.md

@@ -126,6 +126,8 @@ The following sequence of events describes how user code gets notified:
 
 ## Fallback approach
 
+**The fallback approach has been deprecated and will be removed in the future.**
+
 If Device-side Runtime doesn't support `__devicelib_assert_fail` (as reported
 via "cl_intel_devicelib_assert" extension query) then a fallback approach comes
 in place. The approach doesn't require any support from Device-side Runtime and

sycl/doc/extensions/supported/sycl_ext_oneapi_assert.asciidoc

@@ -62,18 +62,11 @@ and `+__LINE__+`, and the value of the standard variable `+__func__+`. If the
 failing assert comes from an `nd_range` `parallel_for` it will also include the
 global ID and the local ID of the failing work item.
 
-Some devices implement `assert()` natively while others use a fallback
-implementation, and the two implementations provide different guarantees. The
-native implementation is most similar to the way `assert()` works on the host. If
-an assertion fails in the native implementation, the assertion message is
-immediately printed to stderr and the program terminates by calling
-`std::abort()`. If an assertion fails with the fallback implementation, the
-failing assert() returns back to its caller and the device code must continue
-executing (without deadlocking) until the kernel completes. The implementation
-prints the assertion message to stderr and terminates with `std::abort()` only
-after the kernel completes execution. An application can determine which of the
-two mechanisms a device uses by testing the device aspect
-`aspect::ext_oneapi_native_assert`.
+Only some devices support `assert()` natively, as determinable by querying the
+new `aspect::ext_oneapi_native_assert` aspect. If an assertion fails in devices
+that support these natively, the assertion message is immediately printed to
+stderr and the program terminates by calling `std::abort()`. Failures in calls
+to `assert()` on devices that do not natively support it are ignored.
 
 The `assert()` macro is defined in system include headers, not in SYCL headers.
 On most of systems it is `<cassert>` and/or `<assert.h>` header files.
@@ -137,9 +130,8 @@ enum class aspect {
 ----
 
 If device has the `ext_oneapi_native_assert` aspect, then its Device-Side
-Runtime is capable of native support of `assert`. That is, safe implementation
-is used. If device doesn't have the aspect, then fallback implementation is
-used.
+Runtime is capable of native support of `assert`.  If device doesn't have the
+aspect, then assertions on the device will be silently ignored.
 
 == Version
 

sycl/include/sycl/detail/assert_happened.hpp

@@ -99,11 +99,10 @@
 static_assert(__cplusplus >= 201703L,
               "DPCPP does not support C++ version earlier than C++17.");
 
-// Helper macro to identify if fallback assert is needed
-#if defined(SYCL_FALLBACK_ASSERT)
-#define __SYCL_USE_FALLBACK_ASSERT SYCL_FALLBACK_ASSERT
-#else
-#define __SYCL_USE_FALLBACK_ASSERT 0
+// MSVC doesn't support #warning and we cannot use other methods to report a
+// warning from inside a system header (which SYCL is considered to be).
+#if defined(SYCL_FALLBACK_ASSERT) && (!defined(_MSC_VER) || defined(__clang__))
+#warning "SYCL_FALLBACK_ASSERT has been removed and no longer has any effect."
 #endif
 
 #if defined(_WIN32) && !defined(_DLL) && !defined(__SYCL_DEVICE_ONLY__)

sycl/include/sycl/detail/defines_elementary.hpp

@@ -100,16 +100,14 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
 template <typename CommandGroupFunc, typename PropertiesT>
 void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
-  Q.submit_without_event<__SYCL_USE_FALLBACK_ASSERT>(
-      Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
+  Q.submit_without_event(Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
 }
 
 template <typename CommandGroupFunc, typename PropertiesT>
 event submit_with_event_impl(const queue &Q, PropertiesT Props,
                              CommandGroupFunc &&CGF,
                              const sycl::detail::code_location &CodeLoc) {
-  return Q.submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
-      Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
+  return Q.submit_with_event(Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
 }
 } // namespace detail