[SYCL] Make handler-less path default for nd_range based functions (#20327)

Switch the default kernel submisson path from handler-based to
handler-less for:
- parallel_for queue shortcut function (nd_range based)
- nd_launch (nd_range overload only, excluding the launch_config
overload)
- launch_grouped (global and local range based)

This switch is limited to the cases, when:
- No kernel function properties are defined
- The kernel function does not use the kernel_handler parameter
- No reductions are used

Author: slawekptak

sycl/cmake/modules/AddSYCLUnitTest.cmake

@@ -1,6 +1,6 @@
 # Internal function to create SYCL unit tests with code reuse
-# add_sycl_unittest_internal(test_dirname SHARED|OBJECT is_preview is_no_cgh file1.cpp, file2.cpp ...)
-function(add_sycl_unittest_internal test_dirname link_variant is_preview is_no_cgh)
+# add_sycl_unittest_internal(test_dirname SHARED|OBJECT is_preview file1.cpp, file2.cpp ...)
+function(add_sycl_unittest_internal test_dirname link_variant is_preview)
   # Enable exception handling for these unit tests
   set(LLVM_REQUIRES_EH ON)
   set(LLVM_REQUIRES_RTTI ON)
@@ -37,10 +37,6 @@ function(add_sycl_unittest_internal test_dirname link_variant is_preview is_no_c
       set(CMAKE_CURRENT_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/Preview")
   endif()
 
-  if (${is_no_cgh})
-    set(CMAKE_CURRENT_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/NoCGH")
-  endif()
-
   if ("${link_variant}" MATCHES "SHARED")
     set(SYCL_LINK_LIBS ${sycl_so_target})
     add_unittest(SYCLUnitTests ${test_dirname} ${ARGN})
@@ -69,18 +65,6 @@ function(add_sycl_unittest_internal test_dirname link_variant is_preview is_no_c
     set(sycl_cache_suffix "_preview")
   endif()
 
-  if (${is_no_cgh})
-    set(sycl_cache_suffix "_no_cgh")
-  endif()
-
-  if (${is_no_cgh})
-    target_compile_definitions(
-      ${test_dirname}
-      PRIVATE
-        __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-    )
-  endif()
-
   if (SYCL_ENABLE_XPTI_TRACING)
     target_compile_definitions(${test_dirname}
       PRIVATE XPTI_ENABLE_INSTRUMENTATION XPTI_STATIC_LIBRARY)
@@ -175,7 +159,6 @@ endfunction()
 # the SYCL preview features enabled.
 # Produces two binaries, named `basename(test_name_prefix_non_preview)` and `basename(test_name_prefix_preview)`
 macro(add_sycl_unittest test_name_prefix link_variant)
-  add_sycl_unittest_internal(${test_name_prefix}_non_preview ${link_variant} FALSE FALSE ${ARGN})
-  add_sycl_unittest_internal(${test_name_prefix}_no_cgh ${link_variant} FALSE TRUE ${ARGN})
-  add_sycl_unittest_internal(${test_name_prefix}_preview ${link_variant} TRUE FALSE ${ARGN})
+  add_sycl_unittest_internal(${test_name_prefix}_non_preview ${link_variant} FALSE ${ARGN})
+  add_sycl_unittest_internal(${test_name_prefix}_preview ${link_variant} TRUE ${ARGN})
 endmacro()

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

@@ -259,7 +259,6 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
           typename KernelType, typename... ReductionsT>
 void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
                ReductionsT &&...Reductions) {
-#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   // TODO The handler-less path does not support reductions, kernel
   // function properties and kernel functions with the kernel_handler
   // type argument yet.
@@ -271,9 +270,7 @@ void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
                     KernelType, sycl::nd_item<Dimensions>>::value)) {
     detail::submit_kernel_direct<KernelName>(std::move(Q), empty_properties_t{},
                                              Range, KernelObj);
-  } else
-#endif
-  {
+  } else {
     submit(std::move(Q), [&](handler &CGH) {
       nd_launch<KernelName>(CGH, Range, KernelObj,
                             std::forward<ReductionsT>(Reductions)...);

sycl/include/sycl/khr/free_function_commands.hpp

@@ -157,7 +157,6 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
-#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   // TODO The handler-less path does not support kernel function properties
   // and kernel functions with the kernel_handler type argument yet.
   if constexpr (!(ext::oneapi::experimental::detail::
@@ -168,9 +167,7 @@ void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
     detail::submit_kernel_direct(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<1>(r, size), std::forward<KernelType>(k));
-  } else
-#endif
-  {
+  } else {
     submit(
         q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
         codeLoc);
@@ -181,7 +178,6 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
-#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   // TODO The handler-less path does not support kernel function properties
   // and kernel functions with the kernel_handler type argument yet.
   if constexpr (!(ext::oneapi::experimental::detail::
@@ -192,9 +188,7 @@ void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
     detail::submit_kernel_direct(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<2>(r, size), std::forward<KernelType>(k));
-  } else
-#endif
-  {
+  } else {
     submit(
         q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
         codeLoc);
@@ -205,7 +199,6 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
-#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   // TODO The handler-less path does not support kernel function properties
   // and kernel functions with the kernel_handler type argument yet.
   if constexpr (!(ext::oneapi::experimental::detail::
@@ -216,9 +209,7 @@ void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
     detail::submit_kernel_direct(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<3>(r, size), std::forward<KernelType>(k));
-  } else
-#endif
-  {
+  } else {
     submit(
         q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
         codeLoc);

sycl/include/sycl/queue.hpp

@@ -3275,7 +3275,6 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
     constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
-#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
     using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
 
     // TODO The handler-less path does not support reductions, kernel
@@ -3290,9 +3289,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
       return detail::submit_kernel_direct<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{}, Range,
           Rest..., TlsCodeLocCapture.query());
-    } else
-#endif
-    {
+    } else {
       return submit(
           [&](handler &CGH) {
             CGH.template parallel_for<KernelName>(Range, Rest...);

sycl/test-e2e/Basic/test_num_kernel_copies.cpp

@@ -29,7 +29,8 @@ int main(int argc, char **argv) {
 
   kernel<1> krn1;
   q.parallel_for(sycl::nd_range<1>{1, 1}, krn1);
-  assert(copy_count == 1);
+  // The kernel is copied on the scheduler-based path only
+  assert(copy_count == 0);
   assert(move_count == 0);
   copy_count = 0;
 

sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionCommandsEvents.cpp

@@ -227,7 +227,6 @@ TEST_F(FreeFunctionCommandsEventsTests, LaunchGroupedShortcutNoEvent) {
   ASSERT_EQ(counter_urEnqueueKernelLaunch, size_t{1});
 }
 
-#if __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
 TEST_F(FreeFunctionCommandsEventsTests,
        LaunchGroupedShortcutMoveKernelNoEvent) {
   mock::getCallbacks().set_replace_callback("urEnqueueKernelLaunch",
@@ -252,6 +251,11 @@ TEST_F(FreeFunctionCommandsEventsTests,
   // to force the scheduler-based submission. In this case, the HostKernel
   // should be constructed.
 
+  // Replace the callback with an event based one, since the scheduler
+  // needs to create an event internally
+  mock::getCallbacks().set_replace_callback(
+      "urEnqueueKernelLaunch", &redefined_urEnqueueKernelLaunchWithEvent);
+
   Queue.submit([&](sycl::handler &CGH) {
     CGH.host_task([&] {
       std::unique_lock<std::mutex> lk(CvMutex);
@@ -274,9 +278,8 @@ TEST_F(FreeFunctionCommandsEventsTests,
   // HostKernel. Copy ctor is called by InstantiateKernelOnHost, can't delete
   // it.
   ASSERT_EQ(TestMoveFunctor::MoveCtorCalls, 1);
-  ASSERT_EQ(counter_urEnqueueKernelLaunch, size_t{2});
+  ASSERT_EQ(counter_urEnqueueKernelLaunchWithEvent, size_t{1});
 }
-#endif
 
 TEST_F(FreeFunctionCommandsEventsTests, SubmitLaunchGroupedKernelNoEvent) {
   mock::getCallbacks().set_replace_callback("urEnqueueKernelLaunch",

sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp

@@ -26,53 +26,48 @@ inline ur_result_t after_urKernelGetInfo(void *pParams) {
 static thread_local size_t counter_urEnqueueKernelLaunch = 0;
 inline ur_result_t redefined_urEnqueueKernelLaunch(void *pParams) {
   ++counter_urEnqueueKernelLaunch;
-// TODO The no-handler scheduler submission includes a fix for the event return,
-// where the event is returned by the scheduler on every submission. This fix
-// is not yet applied to the handler-based path.
-#ifndef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   auto params = *static_cast<ur_enqueue_kernel_launch_params_t *>(pParams);
   EXPECT_EQ(*params.pphEvent, nullptr);
-#endif
+  return UR_RESULT_SUCCESS;
+}
+
+static thread_local size_t counter_urEnqueueKernelLaunchWithEvent = 0;
+inline ur_result_t redefined_urEnqueueKernelLaunchWithEvent(void *pParams) {
+  ++counter_urEnqueueKernelLaunchWithEvent;
+  auto params = *static_cast<ur_enqueue_kernel_launch_params_t *>(pParams);
+  EXPECT_NE(*params.pphEvent, nullptr);
   return UR_RESULT_SUCCESS;
 }
 
 static thread_local size_t counter_urUSMEnqueueMemcpy = 0;
 inline ur_result_t redefined_urUSMEnqueueMemcpy(void *pParams) {
   ++counter_urUSMEnqueueMemcpy;
-#ifndef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   auto params = *static_cast<ur_enqueue_usm_memcpy_params_t *>(pParams);
   EXPECT_EQ(*params.pphEvent, nullptr);
-#endif
   return UR_RESULT_SUCCESS;
 }
 
 static thread_local size_t counter_urUSMEnqueueFill = 0;
 inline ur_result_t redefined_urUSMEnqueueFill(void *pParams) {
   ++counter_urUSMEnqueueFill;
-#ifndef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   auto params = *static_cast<ur_enqueue_usm_fill_params_t *>(pParams);
   EXPECT_EQ(*params.pphEvent, nullptr);
-#endif
   return UR_RESULT_SUCCESS;
 }
 
 static thread_local size_t counter_urUSMEnqueuePrefetch = 0;
 inline ur_result_t redefined_urUSMEnqueuePrefetch(void *pParams) {
   ++counter_urUSMEnqueuePrefetch;
-#ifndef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   auto params = *static_cast<ur_enqueue_usm_prefetch_params_t *>(pParams);
   EXPECT_EQ(*params.pphEvent, nullptr);
-#endif
   return UR_RESULT_SUCCESS;
 }
 
 static thread_local size_t counter_urUSMEnqueueMemAdvise = 0;
 inline ur_result_t redefined_urUSMEnqueueMemAdvise(void *pParams) {
   ++counter_urUSMEnqueueMemAdvise;
-#ifndef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   auto params = *static_cast<ur_enqueue_usm_advise_params_t *>(pParams);
   EXPECT_EQ(*params.pphEvent, nullptr);
-#endif
   return UR_RESULT_SUCCESS;
 }
 

sycl/unittests/compression/CMakeLists.txt

@@ -2,5 +2,4 @@ add_sycl_unittest(CompressionTests OBJECT
   CompressionTests.cpp
 )
 target_compile_definitions(CompressionTests_non_preview PRIVATE SYCL_RT_ZSTD_AVAILABLE)
-target_compile_definitions(CompressionTests_no_cgh PRIVATE SYCL_RT_ZSTD_AVAILABLE)
 target_compile_definitions(CompressionTests_preview PRIVATE SYCL_RT_ZSTD_AVAILABLE __INTEL_PREVIEW_BREAKING_CHANGES)

sycl/unittests/xpti_trace/CMakeLists.txt

@@ -8,4 +8,3 @@ add_sycl_unittest(XptiTraceTests OBJECT
 )
 target_link_libraries(XptiTraceTests_non_preview PRIVATE xpti xptitest_subscriber)
 target_link_libraries(XptiTraceTests_preview PRIVATE xpti xptitest_subscriber)
-target_link_libraries(XptiTraceTests_no_cgh PRIVATE xpti xptitest_subscriber)