[SYCL] Deprecate fallback assertions

sycl/doc/PreprocessorMacros.md

@@ -39,7 +39,7 @@ 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**
+- **SYCL_FALLBACK_ASSERT (deprecated)**
 
   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
@@ -52,6 +52,7 @@ This file describes macros that have effect on SYCL compiler and run-time.
   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.
+  This macro is deprecated and will be removed in the future.
 
 - **SYCL2020_CONFORMANT_APIS (deprecated)**
   This macro is used to comply with the SYCL 2020 specification, as some of the current 

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/defines_elementary.hpp

@@ -94,12 +94,24 @@
 static_assert(__cplusplus >= 201703L,
               "DPCPP does not support C++ version earlier than C++17.");
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 // 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
 #endif
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
+
+// 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__))
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+#warning "SYCL_FALLBACK_ASSERT has been removed and no longer has any effect."
+#else
+#warning "SYCL_FALLBACK_ASSERT is deprecated."
+#endif
+#endif
 
 #if defined(_WIN32) && !defined(_DLL) && !defined(__SYCL_DEVICE_ONLY__)
 // SYCL library is designed such a way that STL objects cross DLL boundary,

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -99,16 +99,25 @@ 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) {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  Q.submit_without_event(Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
+#else
   Q.submit_without_event<__SYCL_USE_FALLBACK_ASSERT>(
       Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
+#endif
 }
 
 template <typename CommandGroupFunc, typename PropertiesT>
 event submit_with_event_impl(const queue &Q, PropertiesT Props,
                              CommandGroupFunc &&CGF,
                              const sycl::detail::code_location &CodeLoc) {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  return Q.submit_with_event(Props, detail::type_erased_cgfo_ty{CGF}, nullptr,
+                             CodeLoc);
+#else
   return Q.submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
       Props, detail::type_erased_cgfo_ty{CGF}, nullptr, CodeLoc);
+#endif
 }
 } // namespace detail
 

sycl/include/sycl/queue.hpp

@@ -14,7 +14,9 @@
 #include <sycl/async_handler.hpp>             // for async_handler
 #include <sycl/backend_types.hpp>             // for backend, backe...
 #include <sycl/buffer.hpp>                    // for buffer
-#include <sycl/detail/assert_happened.hpp>    // for AssertHappened
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+#include <sycl/detail/assert_happened.hpp>
+#endif
 #include <sycl/detail/cg_types.hpp>           // for check_fn_signa...
 #include <sycl/detail/common.hpp>             // for code_location
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEP...
@@ -66,6 +68,7 @@ auto get_native(const SyclObjectT &Obj)
 namespace detail {
 class queue_impl;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 inline event submitAssertCapture(const queue &, event &, queue *,
                                  const detail::code_location &);
 
@@ -76,6 +79,7 @@ inline event submitAssertCapture(const queue &, event &, queue *,
 //                         meaningful when IsKernel is true
 // event &Event - event after which post processing should be executed
 using SubmitPostProcessF = std::function<void(bool, bool, event &)>;
+#endif
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 struct SubmissionInfoImpl;
@@ -84,8 +88,10 @@ class __SYCL_EXPORT SubmissionInfo {
 public:
   SubmissionInfo();
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   sycl::detail::optional<SubmitPostProcessF> &PostProcessorFunc();
   const sycl::detail::optional<SubmitPostProcessF> &PostProcessorFunc() const;
+#endif
 
   std::shared_ptr<detail::queue_impl> &SecondaryQueue();
   const std::shared_ptr<detail::queue_impl> &SecondaryQueue() const;
@@ -121,14 +127,14 @@ class __SYCL_EXPORT SubmissionInfo {
   SubmissionInfo(const detail::SubmissionInfo &SI)
       : MPostProcessorFunc(SI.PostProcessorFunc()),
         MSecondaryQueue(SI.SecondaryQueue()), MEventMode(SI.EventMode()) {}
-#endif
 
   sycl::detail::optional<SubmitPostProcessF> &PostProcessorFunc() {
     return MPostProcessorFunc;
   }
   const sycl::detail::optional<SubmitPostProcessF> &PostProcessorFunc() const {
     return MPostProcessorFunc;
   }
+#endif
 
   std::shared_ptr<detail::queue_impl> &SecondaryQueue() {
     return MSecondaryQueue;
@@ -143,7 +149,9 @@ class __SYCL_EXPORT SubmissionInfo {
   }
 
 private:
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   optional<detail::SubmitPostProcessF> MPostProcessorFunc = std::nullopt;
+#endif
   std::shared_ptr<detail::queue_impl> MSecondaryQueue = nullptr;
   ext::oneapi::experimental::event_mode_enum MEventMode =
       ext::oneapi::experimental::event_mode_enum::none;
@@ -426,9 +434,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
       T CGF,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+    return submit_with_event(
+        sycl::ext::oneapi::experimental::empty_properties_t{},
+        detail::type_erased_cgfo_ty{CGF}, CodeLoc);
+#else
     return submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
         sycl::ext::oneapi::experimental::empty_properties_t{},
         detail::type_erased_cgfo_ty{CGF}, CodeLoc);
+#endif
   }
 
   /// Submits a command group function object to the queue, in order to be
@@ -446,9 +460,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
       T CGF, queue &SecondaryQueue,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+    return submit_with_event(
+        sycl::ext::oneapi::experimental::empty_properties_t{},
+        detail::type_erased_cgfo_ty{CGF}, &SecondaryQueue, CodeLoc);
+#else
     return submit_with_event<__SYCL_USE_FALLBACK_ASSERT>(
         sycl::ext::oneapi::experimental::empty_properties_t{},
         detail::type_erased_cgfo_ty{CGF}, &SecondaryQueue, CodeLoc);
+#endif
   }
 
   /// Prevents any commands submitted afterward to this queue from executing
@@ -3581,9 +3601,11 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   friend auto get_native(const SyclObjectT &Obj)
       -> backend_return_t<BackendName, SyclObjectT>;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 #if __SYCL_USE_FALLBACK_ASSERT
   friend event detail::submitAssertCapture(const queue &, event &, queue *,
                                            const detail::code_location &);
+#endif
 #endif
 
   template <typename CommandGroupFunc, typename PropertiesT>
@@ -3698,7 +3720,11 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   // UseFallBackAssert as template param vs `#if` in function body is necessary
   // to prevent ODR-violation between TUs built with different fallback assert
   // modes.
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  template <typename PropertiesT>
+#else
   template <bool UseFallbackAssert, typename PropertiesT>
+#endif
   event submit_with_event(PropertiesT Props,
                           const detail::type_erased_cgfo_ty &CGF,
                           queue *SecondaryQueuePtr,
@@ -3709,6 +3735,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     ProcessSubmitProperties(Props, SI);
     if (SecondaryQueuePtr)
       SI.SecondaryQueue() = detail::getSyclObjImpl(*SecondaryQueuePtr);
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     if constexpr (UseFallbackAssert)
       SI.PostProcessorFunc() =
           [this, &SecondaryQueuePtr,
@@ -3723,6 +3750,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                                   TlsCodeLocCapture.query());
             }
           };
+#endif
     return submit_with_event_impl(CGF, SI, TlsCodeLocCapture.query(),
                                   TlsCodeLocCapture.isToplevel());
   }
@@ -3738,14 +3766,19 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   // UseFallBackAssert as template param vs `#if` in function body is necessary
   // to prevent ODR-violation between TUs built with different fallback assert
   // modes.
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  template <typename PropertiesT>
+#else
   template <bool UseFallbackAssert, typename PropertiesT>
+#endif
   event submit_with_event(PropertiesT Props,
                           const detail::type_erased_cgfo_ty &CGF,
                           const detail::code_location &CodeLoc =
                               detail::code_location::current()) const {
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
     detail::v1::SubmissionInfo SI{};
     ProcessSubmitProperties(Props, SI);
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     if constexpr (UseFallbackAssert)
       SI.PostProcessorFunc() = [this, &TlsCodeLocCapture](bool IsKernel,
                                                           bool KernelUsesAssert,
@@ -3759,6 +3792,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
           submitAssertCapture(*this, E, nullptr, TlsCodeLocCapture.query());
         }
       };
+#endif
     return submit_with_event_impl(CGF, SI, TlsCodeLocCapture.query(),
                                   TlsCodeLocCapture.isToplevel());
   }
@@ -3773,15 +3807,22 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   // UseFallBackAssert as template param vs `#if` in function body is necessary
   // to prevent ODR-violation between TUs built with different fallback assert
   // modes.
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  template <typename PropertiesT>
+#else
   template <bool UseFallbackAssert, typename PropertiesT>
+#endif
   void submit_without_event(PropertiesT Props,
                             const detail::type_erased_cgfo_ty &CGF,
                             const detail::code_location &CodeLoc) const {
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     if constexpr (UseFallbackAssert) {
       // If post-processing is needed, fall back to the regular submit.
       // TODO: Revisit whether we can avoid this.
       submit_with_event<UseFallbackAssert>(Props, CGF, CodeLoc);
-    } else {
+    } else
+#endif
+    {
       detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
       detail::v1::SubmissionInfo SI{};
       ProcessSubmitProperties(Props, SI);
@@ -3934,6 +3975,7 @@ template <> struct __SYCL_EXPORT hash<sycl::queue> {
 };
 } // namespace std
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 #if __SYCL_USE_FALLBACK_ASSERT
 // Explicitly request format macros
 #ifndef __STDC_FORMAT_MACROS
@@ -4032,3 +4074,4 @@ event submitAssertCapture(const queue &Self, event &Event,
 } // namespace _V1
 } // namespace sycl
 #endif // __SYCL_USE_FALLBACK_ASSERT
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES

sycl/source/detail/program_manager/program_manager.cpp

@@ -1526,10 +1526,13 @@ ProgramManager::getDeviceImage(KernelNameStrRefT KernelName,
         KernelId != m_KernelName2KernelIDs.end()) {
       Img = getBinImageFromMultiMap(m_KernelIDs2BinImage, KernelId->second,
                                     ContextImpl, DeviceImpl);
-    } else {
+    }
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+    else {
       Img = getBinImageFromMultiMap(m_ServiceKernels, KernelName, ContextImpl,
                                     DeviceImpl);
     }
+#endif
   }
 
   // Decompress the image if it is compressed.
@@ -2015,6 +2018,7 @@ void ProgramManager::addImage(sycl_device_binary RawImg,
 
     auto name = EntriesIt->GetName();
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     // Skip creating unique kernel ID if it is a service kernel.
     // SYCL service kernels are identified by having
     // __sycl_service_kernel__ in the mangled name, primarily as part of
@@ -2023,6 +2027,7 @@ void ProgramManager::addImage(sycl_device_binary RawImg,
       m_ServiceKernels.insert(std::make_pair(name, Img.get()));
       continue;
     }
+#endif
 
     // Skip creating unique kernel ID if it is an exported device
     // function. Exported device functions appear in the offload entries
@@ -2141,11 +2146,13 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
          EntriesIt = EntriesIt->Increment()) {
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {
         m_ServiceKernels.erase(EntriesIt->GetName());
         continue;
       }
+#endif
 
       // Exported device functions won't have a kernel ID
       if (m_ExportedSymbolImages.find(EntriesIt->GetName()) !=

sycl/source/detail/program_manager/program_manager.hpp

@@ -459,6 +459,7 @@ class ProgramManager {
   /// \ref Sync::getGlobalLock() while holding this mutex.
   std::mutex m_KernelIDsMutex;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   /// Caches all found service kernels to expedite future checks. A SYCL service
   /// kernel is a kernel that has not been defined by the user but is instead
   /// generated by the SYCL runtime. Service kernel name types must be declared
@@ -467,6 +468,7 @@ class ProgramManager {
   /// Access must be guarded by the m_KernelIDsMutex mutex.
   std::unordered_multimap<KernelNameStrT, const RTDeviceBinaryImage *>
       m_ServiceKernels;
+#endif
 
   /// Caches all exported symbols to allow faster lookup when excluding these
   /// from kernel bundles.