[SYCL] Disable fallback assert by default (#4694)

Due excessive amount of issues submitted against fallback assert implementation the fallback assert is now disabled by default until fixed.

Author: sergei

sycl/doc/PreprocessorMacros.md

@@ -38,16 +38,19 @@ This file describes macros that have effect on SYCL compiler and run-time.
 
   Disables a message which warns about unsupported C++ version.
 
-- **SYCL_DISABLE_FALLBACK_ASSERT**
-
-  Defining this macro eliminates some overhead that is associated with
-  submitting kernels that call `assert()`. When this macro is 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. However, this macro only affects kernels that are submitted to devices
-  that do **not** have native support for `assert()` because devices with native
-  support do not impose any extra overhead. One can check to see if a device has
-  native support for `assert()` via `aspect::ext_oneapi_native_assert`.
+- **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**
   This macro is used to comply with the SYCL 2020 specification, as some of the current 

sycl/include/CL/sycl/queue.hpp

@@ -67,8 +67,8 @@
 
 // Helper macro to identify if fallback assert is needed
 // FIXME remove __NVPTX__ condition once devicelib supports CUDA
-#if !defined(SYCL_DISABLE_FALLBACK_ASSERT)
-#define __SYCL_USE_FALLBACK_ASSERT 1
+#if defined(SYCL_FALLBACK_ASSERT)
+#define __SYCL_USE_FALLBACK_ASSERT SYCL_FALLBACK_ASSERT
 #else
 #define __SYCL_USE_FALLBACK_ASSERT 0
 #endif

sycl/test/check_device_code/kernel_arguments_as.cpp

@@ -5,8 +5,8 @@
 //
 // Check the address space of the pointer in accessor class.
 //
-// CHECK: %struct.AccWrapper = type { %"class.cl::sycl::accessor.1" }
-// CHECK: %"class.cl::sycl::accessor.1" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
+// CHECK: %struct.AccWrapper = type { %"class.cl::sycl::accessor[[NUMBER_SUFFIX:\.?[0-9]*]]" }
+// CHECK: %"class.cl::sycl::accessor[[NUMBER_SUFFIX]]" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
 // CHECK-DISABLE: %[[UNION]] = type { i32 addrspace(1)* }
 // CHECK-ENABLE: %[[UNION]] = type { i32 addrspace(5)* }
 // CHECK: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::accessor.[[NUM:[0-9]+]]" }

sycl/unittests/assert/assert.cpp

@@ -18,6 +18,7 @@
  * pipe.
  */
 
+#define SYCL_FALLBACK_ASSERT 1
 // Enable use of interop kernel c-tor
 #define __SYCL_INTERNAL_API
 #include <CL/sycl.hpp>