[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());
+  }
+
+  const void *Arguments = nullptr;
+  int64_t ArgumentsSize = 0;
+  if (TypelessArgs.getArgCount()) {
+    // without decomposition and free functions extension we always expect 1
+    // argument to the kernel - lambda capture.
+    assert(TypelessArgs.getArgCount() == 1 &&
+           "No arg decomposition or extensions are supported now.");
+    // TODO: liboffload doesn't support more than 1 argument without copy now.
+    // It doesn't expect array of arguments, it requires a contiguous memory
+    // with args. While we have only 1 argument we don't need extra handling
+    // here, we just pass the first argument directly.
+    Arguments = TypelessArgs.getArgPtrArray()[0];
+    ArgumentsSize = TypelessArgs.getSizesArray()[0];
+  }
+
+  // ol_kernel_launch_prop_t Props[2];
+  // Props[0].type = OL_KERNEL_LAUNCH_PROP_TYPE_SIZE;
+  // Props[0].data = &ArgumentsSize;
+  // Props[1] = OL_KERNEL_LAUNCH_PROP_END;
----------------
sergey-semenov wrote:

Leftover code?

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