[llvm-branch-commits] [llvm] [libsycl] add single_task (PR #188797)

[Sergey Semenov via llvm-branch-commits]

================
@@ -34,5 +60,91 @@ backend QueueImpl::getBackend() const noexcept { return MDevice.getBackend(); }
 
 void QueueImpl::wait() { callAndThrow(olSyncQueue, MOffloadQueue); }
 
+static bool checkEventsPlatformMatch(std::vector<EventImplPtr> &Events,
+                                     const PlatformImpl &QueuePlatform) {
+  // liboffload limitation to olWaitEvents. We can't do any extra handling for
+  // cross context/platform events without host task support now.
+  //   "The input events can be from any queue on any device provided by the
+  //   same platform as `Queue`."
+  return std::all_of(Events.cbegin(), Events.cend(),
+                     [&QueuePlatform](const EventImplPtr &Event) {
+                       return &Event->getPlatformImpl() == &QueuePlatform;
+                     });
+}
+
+void QueueImpl::setKernelParameters(std::vector<EventImplPtr> &&Events,
+                                    const detail::UnifiedRangeView &Range) {
+  if (!checkEventsPlatformMatch(Events, MDevice.getPlatformImpl()))
+    throw sycl::exception(
+        sycl::make_error_code(sycl::errc::feature_not_supported),
+        "libsycl doesn't support cross-context/platform event dependencies "
+        "now.");
+
+  // TODO: this convertion and storing only offload events is possible only
+  // while we don't have host tasks (and features based on host tasks, like
+  // streams). With them - it is very likely we should copy EventImplPtr
+  // (shared_ptr) and keep it here. Although it may differ if host tasks will be
+  // implemented on offload level (no data now).
+  assert(MCurrentSubmitInfo.DepEvents.empty() &&
+         "Kernel submission must clean up dependencies.");
+  MCurrentSubmitInfo.DepEvents.reserve(Events.size());
+  for (auto &Event : Events) {
+    assert(Event && "Event impl object can't be nullptr");
+    MCurrentSubmitInfo.DepEvents.push_back(Event->getHandle());
+  }
+  setKernelLaunchArgs(Range, MCurrentSubmitInfo.Range);
+}
+
+void QueueImpl::submitKernelImpl(const char *KernelName,
+                                 detail::ArgCollection &TypelessArgs) {
+  ol_symbol_handle_t Kernel =
+      detail::ProgramManager::getInstance().getOrCreateKernel(KernelName,
+                                                              MDevice);
+  assert(Kernel);
+
+  ol_event_handle_t NewEvent{};
+  if (!MCurrentSubmitInfo.DepEvents.empty()) {
+    callAndThrow(olWaitEvents, MOffloadQueue,
+                 MCurrentSubmitInfo.DepEvents.data(),
+                 MCurrentSubmitInfo.DepEvents.size());
+  }
----------------
sergey-semenov wrote:

I think this part warrants a comment about why we currently have to wait before submission.

https://github.com/llvm/llvm-project/pull/188797