[SYCL][L0] Implementation of ext_oneapi_prod API

sycl/doc/extensions/proposed/sycl_ext_oneapi_prod.asciidoc → sycl/doc/extensions/supported/sycl_ext_oneapi_prod.asciidoc

@@ -36,17 +36,13 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 5 specification.  All
+This extension is written against the SYCL 2020 revision 8 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This extension is implemented and fully supported by DPC++.
 
 == Overview
 

sycl/include/sycl/queue.hpp

@@ -2590,6 +2590,11 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
         CodeLoc);
   }
 
+  /// Provides a hint to the  runtime that previously issued commands to this
+  /// queue should begin executing once their prerequisites have been satisfied.
+  ///
+  void ext_oneapi_prod();
+
   /// Returns whether the queue is in order or OoO
   ///
   /// Equivalent to has_property<property::queue::in_order>()

sycl/source/detail/queue_impl.hpp

@@ -388,6 +388,20 @@ class queue_impl {
   template <typename Param>
   typename Param::return_type get_backend_info() const;
 
+  /// Provides a hint to the backend to execute previously issued commands on
+  /// this queue. Overrides normal batching behaviour. Note that this is merely
+  /// a hint and not a guarantee.
+  void flush() {
+    if (MGraph.lock()) {
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "flush cannot be called for a queue which is "
+                            "recording to a command graph.");
+    }
+    for (const auto &queue : MQueues) {
+      getPlugin()->call<PiApiKind::piQueueFlush>(queue);
+    }
+  }
+
   using SubmitPostProcessF = std::function<void(bool, bool, event &)>;
 
   /// Submits a command group function object to the queue, in order to be

sycl/source/feature_test.hpp.in

@@ -106,6 +106,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_PRIVATE_ALLOCA 1
 #define SYCL_EXT_ONEAPI_FORWARD_PROGRESS 1
 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_KERNELS 1
+#define SYCL_EXT_ONEAPI_PROD 1
 
 #ifndef __has_include
 #define __has_include(x) 0

sycl/source/queue.cpp

@@ -323,6 +323,8 @@ backend queue::get_backend() const noexcept { return getImplBackend(impl); }
 
 bool queue::ext_oneapi_empty() const { return impl->ext_oneapi_empty(); }
 
+void queue::ext_oneapi_prod() { impl->flush(); }
+
 pi_native_handle queue::getNative(int32_t &NativeHandleDesc) const {
   return impl->getNative(NativeHandleDesc);
 }

sycl/test-e2e/QueueFlushing/queue_flushing.cpp

@@ -0,0 +1,85 @@
+// RUN: %{build} -o %t.out
+// RUN: %t.out
+
+#include <stdlib.h>
+#include <sycl/sycl.hpp>
+#include <vector>
+
+using namespace sycl;
+using namespace std;
+
+// This test checks the ext_oneapi_prod extension. This extension introduces
+// only one new function, queue::ext_oneapi_prod(), which serves as a hint to
+// the compiler to flush the queue. Since it is simply a hint, we cannot really
+// test what the backend is doing but we can at least make sure that code
+// involving this function compiles and runs successfully in various contexts.
+int main() {
+  sycl::queue q;
+
+  // Test on an empty queue multiple times.
+  q.ext_oneapi_prod();
+  q.ext_oneapi_prod();
+
+  // Test on a queue after we've submitted a kernel in various contexts.
+  q.single_task([]() {});
+  q.ext_oneapi_prod();
+
+  q.parallel_for(range<1>{}, [=](auto &idx) {});
+  q.ext_oneapi_prod();
+  q.wait();
+
+  // Test that the result of an in-progress addition kernel is not affected by
+  // calling ext_oneapi_prod.
+  srand(time(0));
+  constexpr int N = 16;
+  int A[N];
+  int B[N];
+  int add[N];
+  int mult[N];
+  for (int i = 0; i < N; ++i) {
+    A[i] = rand();
+    B[i] = rand();
+  }
+  {
+    buffer<int> bufA{A, N};
+    buffer<int> bufB{B, N};
+    buffer<int> bufadd{add, N};
+
+    q.submit([&](handler &cgh) {
+      accessor accA{bufA, cgh};
+      accessor accB{bufB, cgh};
+      accessor accadd{bufadd, cgh};
+      cgh.parallel_for(N, [=](id<1> i) { accadd[i] = accA[i] + accB[i]; });
+    });
+    q.ext_oneapi_prod();
+  }
+  for (int i = 0; i < N; ++i) {
+    assert(add[i] == (A[i] + B[i]));
+  }
+
+  // Test that the result of an in-progress addition and multiplication kernel
+  // is not affected by calling ext_oneapi_prod
+  {
+    buffer<int> bufA{A, N};
+    buffer<int> bufB{B, N};
+    buffer<int> bufadd{add, N};
+    buffer<int> bufmult{mult, N};
+    q.submit([&](handler &cgh) {
+      accessor accA{bufA, cgh};
+      accessor accB{bufB, cgh};
+      accessor accadd{bufadd, cgh};
+      cgh.parallel_for(N, [=](id<1> i) { accadd[i] = accA[i] + accB[i]; });
+    });
+
+    q.submit([&](handler &cgh) {
+      accessor accA{bufA, cgh};
+      accessor accB{bufB, cgh};
+      accessor accmult{bufmult, cgh};
+      cgh.parallel_for(N, [=](id<1> i) { accmult[i] = accA[i] * accB[i]; });
+    });
+    q.ext_oneapi_prod();
+  }
+  for (int i = 0; i < N; ++i) {
+    assert(add[i] == (A[i] + B[i]) && mult[i] == (A[i] * B[i]));
+  }
+}

sycl/test/abi/sycl_symbols_linux.dump

@@ -3747,6 +3747,7 @@ _ZNK4sycl3_V15queue12has_propertyINS0_8property5queue16enable_profilingEEEbv
 _ZNK4sycl3_V15queue12has_propertyINS0_8property5queue4cuda18use_default_streamEEEbv
 _ZNK4sycl3_V15queue12has_propertyINS0_8property5queue8in_orderEEEbv
 _ZNK4sycl3_V15queue16ext_oneapi_emptyEv
+_ZN4sycl3_V15queue15ext_oneapi_prodEv
 _ZNK4sycl3_V15queue16get_backend_infoINS0_4info6device15backend_versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv
 _ZNK4sycl3_V15queue16get_backend_infoINS0_4info6device7versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv
 _ZNK4sycl3_V15queue16get_backend_infoINS0_4info8platform7versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv

sycl/test/abi/sycl_symbols_windows.dump

@@ -4085,6 +4085,7 @@
 ?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@0AEBUimage_descriptor@56723@@Z
 ?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@PEAXAEBUimage_descriptor@56723@@Z
 ?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@V?$range@$02@23@AEBUimage_descriptor@56723@PEAX111@Z
+?ext_oneapi_prod@queue@_V1@sycl@@QEAAXXZ
 ?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KAEBUcode_location@detail@23@@Z
 ?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z
 ?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KV423@AEBUcode_location@detail@23@@Z

sycl/unittests/Extensions/CMakeLists.txt

@@ -9,6 +9,7 @@ add_sycl_unittest(ExtensionsTests OBJECT
   OneAPISubGroupMask.cpp
   USMP2P.cpp
   CompositeDevice.cpp
+  OneAPIProd.cpp
 )
 
 add_subdirectory(CommandGraph)

sycl/unittests/Extensions/OneAPIProd.cpp

@@ -0,0 +1,42 @@
+//==-------- OneAPIProd.cpp --- sycl_ext_oneapi_prod unit tests ------------==//
+//
+// 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 <sycl/sycl.hpp>
+
+#include <helpers/PiMock.hpp>
+
+#include <gtest/gtest.h>
+
+using namespace sycl;
+
+static bool QueueFlushed = false;
+
+static pi_result redefinedQueueFlush(pi_queue Queue) {
+  QueueFlushed = true;
+  return PI_SUCCESS;
+}
+
+TEST(OneAPIProdTest, PiQueueFlush) {
+  sycl::unittest::PiMock Mock(backend::ext_oneapi_level_zero);
+  sycl::platform Plt = Mock.getPlatform();
+  Mock.redefine<detail::PiApiKind::piQueueFlush>(redefinedQueueFlush);
+  context Ctx{Plt};
+  queue Queue{Ctx, default_selector_v};
+  Queue.ext_oneapi_prod();
+  EXPECT_TRUE(QueueFlushed);
+  sycl::ext::oneapi::experimental::command_graph Graph(Ctx, Queue.get_device());
+  Graph.begin_recording(Queue);
+  try {
+    Queue.ext_oneapi_prod(); // flushing while graph is recording is not allowed
+    FAIL() << "Expected exception when calling ext_oneapi_prod() during graph "
+              "recording not seen.";
+  } catch (exception &ex) {
+    EXPECT_EQ(ex.code(), make_error_code(errc::invalid));
+  }
+  Graph.end_recording(Queue);
+}