graph_impl.hpp

Go to the documentation of this file.

//==--------- graph_impl.hpp --- SYCL graph extension ---------------------==//
//
// 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
//
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/detail/cg_types.hpp>
#include <sycl/detail/os_util.hpp>
#include <sycl/ext/oneapi/experimental/graph.hpp>
#include <sycl/handler.hpp>
 
#include <detail/accessor_impl.hpp>
#include <detail/cg.hpp>
#include <detail/event_impl.hpp>
#include <detail/host_task.hpp>
#include <detail/kernel_impl.hpp>
#include <detail/sycl_mem_obj_t.hpp>
 
#include <cstring>
#include <deque>
#include <fstream>
#include <functional>
#include <iomanip>
#include <list>
#include <set>
#include <shared_mutex>
 
namespace sycl {
inline namespace _V1 {
 
namespace detail {
class SYCLMemObjT;
}
 
namespace ext {
namespace oneapi {
namespace experimental {
namespace detail {
 
inline node_type getNodeTypeFromCG(sycl::detail::CGType CGType) {
  using sycl::detail::CG;
 
  switch (CGType) {
  case sycl::detail::CGType::None:
    return node_type::empty;
  case sycl::detail::CGType::Kernel:
    return node_type::kernel;
  case sycl::detail::CGType::CopyAccToPtr:
  case sycl::detail::CGType::CopyPtrToAcc:
  case sycl::detail::CGType::CopyAccToAcc:
  case sycl::detail::CGType::CopyUSM:
    return node_type::memcpy;
  case sycl::detail::CGType::Memset2DUSM:
    return node_type::memset;
  case sycl::detail::CGType::Fill:
  case sycl::detail::CGType::FillUSM:
    return node_type::memfill;
  case sycl::detail::CGType::PrefetchUSM:
    return node_type::prefetch;
  case sycl::detail::CGType::AdviseUSM:
    return node_type::memadvise;
  case sycl::detail::CGType::Barrier:
  case sycl::detail::CGType::BarrierWaitlist:
    return node_type::ext_oneapi_barrier;
  case sycl::detail::CGType::CodeplayHostTask:
    return node_type::host_task;
  case sycl::detail::CGType::ExecCommandBuffer:
    return node_type::subgraph;
  default:
    assert(false && "Invalid Graph Node Type");
    return node_type::empty;
  }
}
 
class node_impl {
public:
  using id_type = uint64_t;
 
  id_type MID = getNextNodeID();
  std::vector<std::weak_ptr<node_impl>> MSuccessors;
  std::vector<std::weak_ptr<node_impl>> MPredecessors;
  sycl::detail::CGType MCGType = sycl::detail::CGType::None;
  node_type MNodeType = node_type::empty;
  std::unique_ptr<sycl::detail::CG> MCommandGroup;
  std::shared_ptr<exec_graph_impl> MSubGraphImpl;
 
  bool MVisited = false;
 
  int MPartitionNum = -1;
 
  bool MNDRangeUsed = false;
 
  void registerSuccessor(const std::shared_ptr<node_impl> &Node,
                         const std::shared_ptr<node_impl> &Prev) {
    if (std::find_if(MSuccessors.begin(), MSuccessors.end(),
                     [Node](const std::weak_ptr<node_impl> &Ptr) {
                       return Ptr.lock() == Node;
                     }) != MSuccessors.end()) {
      return;
    }
    MSuccessors.push_back(Node);
    Node->registerPredecessor(Prev);
  }
 
  void registerPredecessor(const std::shared_ptr<node_impl> &Node) {
    if (std::find_if(MPredecessors.begin(), MPredecessors.end(),
                     [&Node](const std::weak_ptr<node_impl> &Ptr) {
                       return Ptr.lock() == Node;
                     }) != MPredecessors.end()) {
      return;
    }
    MPredecessors.push_back(Node);
  }
 
  node_impl() {}
 
  node_impl(node_type NodeType,
            std::unique_ptr<sycl::detail::CG> &&CommandGroup)
      : MCGType(CommandGroup->getType()), MNodeType(NodeType),
        MCommandGroup(std::move(CommandGroup)) {
    if (NodeType == node_type::subgraph) {
      MSubGraphImpl =
          static_cast<sycl::detail::CGExecCommandBuffer *>(MCommandGroup.get())
              ->MExecGraph;
    }
  }
 
  node_impl(node_impl &Other)
      : MSuccessors(Other.MSuccessors), MPredecessors(Other.MPredecessors),
        MCGType(Other.MCGType), MNodeType(Other.MNodeType),
        MCommandGroup(Other.getCGCopy()), MSubGraphImpl(Other.MSubGraphImpl) {}
 
  node_impl &operator=(node_impl &Other) {
    if (this != &Other) {
      MSuccessors = Other.MSuccessors;
      MPredecessors = Other.MPredecessors;
      MCGType = Other.MCGType;
      MNodeType = Other.MNodeType;
      MCommandGroup = Other.getCGCopy();
      MSubGraphImpl = Other.MSubGraphImpl;
    }
    return *this;
  }
  bool hasRequirementDependency(sycl::detail::AccessorImplHost *IncomingReq) {
    if (!MCommandGroup)
      return false;
 
    access_mode InMode = IncomingReq->MAccessMode;
    switch (InMode) {
    case access_mode::read:
    case access_mode::read_write:
    case access_mode::atomic:
      break;
    // These access modes don't care about existing buffer data, so we don't
    // need a dependency.
    case access_mode::write:
    case access_mode::discard_read_write:
    case access_mode::discard_write:
      return false;
    }
 
    for (sycl::detail::AccessorImplHost *CurrentReq :
         MCommandGroup->getRequirements()) {
      if (IncomingReq->MSYCLMemObj == CurrentReq->MSYCLMemObj) {
        access_mode CurrentMode = CurrentReq->MAccessMode;
        // Since we have an incoming read requirement, we only care
        // about requirements on this node if they are write
        if (CurrentMode != access_mode::read) {
          return true;
        }
      }
    }
    // No dependency necessary
    return false;
  }
 
  bool isEmpty() const {
    return ((MCGType == sycl::detail::CGType::None) ||
            (MCGType == sycl::detail::CGType::Barrier));
  }
 
  std::unique_ptr<sycl::detail::CG> getCGCopy() const {
    switch (MCGType) {
    case sycl::detail::CGType::Kernel: {
      auto CGCopy = createCGCopy<sycl::detail::CGExecKernel>();
      rebuildArgStorage(CGCopy->MArgs, MCommandGroup->getArgsStorage(),
                        CGCopy->getArgsStorage());
      return std::move(CGCopy);
    }
    case sycl::detail::CGType::CopyAccToPtr:
    case sycl::detail::CGType::CopyPtrToAcc:
    case sycl::detail::CGType::CopyAccToAcc:
      return createCGCopy<sycl::detail::CGCopy>();
    case sycl::detail::CGType::Fill:
      return createCGCopy<sycl::detail::CGFill>();
    case sycl::detail::CGType::UpdateHost:
      return createCGCopy<sycl::detail::CGUpdateHost>();
    case sycl::detail::CGType::CopyUSM:
      return createCGCopy<sycl::detail::CGCopyUSM>();
    case sycl::detail::CGType::FillUSM:
      return createCGCopy<sycl::detail::CGFillUSM>();
    case sycl::detail::CGType::PrefetchUSM:
      return createCGCopy<sycl::detail::CGPrefetchUSM>();
    case sycl::detail::CGType::AdviseUSM:
      return createCGCopy<sycl::detail::CGAdviseUSM>();
    case sycl::detail::CGType::Copy2DUSM:
      return createCGCopy<sycl::detail::CGCopy2DUSM>();
    case sycl::detail::CGType::Fill2DUSM:
      return createCGCopy<sycl::detail::CGFill2DUSM>();
    case sycl::detail::CGType::Memset2DUSM:
      return createCGCopy<sycl::detail::CGMemset2DUSM>();
    case sycl::detail::CGType::EnqueueNativeCommand:
    case sycl::detail::CGType::CodeplayHostTask: {
      // The unique_ptr to the `sycl::detail::HostTask`, which is also used for
      // a EnqueueNativeCommand command, in the HostTask CG prevents from
      // copying the CG. We overcome this restriction by creating a new CG with
      // the same data.
      auto CommandGroupPtr =
          static_cast<sycl::detail::CGHostTask *>(MCommandGroup.get());
      sycl::detail::HostTask HostTask = *CommandGroupPtr->MHostTask.get();
      auto HostTaskSPtr = std::make_shared<sycl::detail::HostTask>(HostTask);
 
      sycl::detail::CG::StorageInitHelper Data(
          CommandGroupPtr->getArgsStorage(), CommandGroupPtr->getAccStorage(),
          CommandGroupPtr->getSharedPtrStorage(),
          CommandGroupPtr->getRequirements(), CommandGroupPtr->getEvents());
 
      std::vector<sycl::detail::ArgDesc> NewArgs = CommandGroupPtr->MArgs;
 
      rebuildArgStorage(NewArgs, CommandGroupPtr->getArgsStorage(),
                        Data.MArgsStorage);
 
      sycl::detail::code_location Loc(CommandGroupPtr->MFileName.data(),
                                      CommandGroupPtr->MFunctionName.data(),
                                      CommandGroupPtr->MLine,
                                      CommandGroupPtr->MColumn);
 
      return std::make_unique<sycl::detail::CGHostTask>(
          sycl::detail::CGHostTask(
              std::move(HostTaskSPtr), CommandGroupPtr->MQueue,
              CommandGroupPtr->MContext, std::move(NewArgs), std::move(Data),
              CommandGroupPtr->getType(), Loc));
    }
    case sycl::detail::CGType::Barrier:
    case sycl::detail::CGType::BarrierWaitlist:
      // Barrier nodes are stored in the graph with only the base CG class,
      // since they are treated internally as empty nodes.
      return createCGCopy<sycl::detail::CG>();
    case sycl::detail::CGType::CopyToDeviceGlobal:
      return createCGCopy<sycl::detail::CGCopyToDeviceGlobal>();
    case sycl::detail::CGType::CopyFromDeviceGlobal:
      return createCGCopy<sycl::detail::CGCopyFromDeviceGlobal>();
    case sycl::detail::CGType::ReadWriteHostPipe:
      return createCGCopy<sycl::detail::CGReadWriteHostPipe>();
    case sycl::detail::CGType::CopyImage:
      return createCGCopy<sycl::detail::CGCopyImage>();
    case sycl::detail::CGType::SemaphoreSignal:
      return createCGCopy<sycl::detail::CGSemaphoreSignal>();
    case sycl::detail::CGType::SemaphoreWait:
      return createCGCopy<sycl::detail::CGSemaphoreWait>();
    case sycl::detail::CGType::ProfilingTag:
      return createCGCopy<sycl::detail::CGProfilingTag>();
    case sycl::detail::CGType::ExecCommandBuffer:
      return createCGCopy<sycl::detail::CGExecCommandBuffer>();
    case sycl::detail::CGType::None:
      return nullptr;
    }
    return nullptr;
  }
 
  bool isSimilar(const std::shared_ptr<node_impl> &Node,
                 bool CompareContentOnly = false) const {
    if (!CompareContentOnly) {
      if (MSuccessors.size() != Node->MSuccessors.size())
        return false;
 
      if (MPredecessors.size() != Node->MPredecessors.size())
        return false;
    }
    if (MCGType != Node->MCGType)
      return false;
 
    switch (MCGType) {
    case sycl::detail::CGType::Kernel: {
      sycl::detail::CGExecKernel *ExecKernelA =
          static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get());
      sycl::detail::CGExecKernel *ExecKernelB =
          static_cast<sycl::detail::CGExecKernel *>(Node->MCommandGroup.get());
      return ExecKernelA->MKernelName.compare(ExecKernelB->MKernelName) == 0;
    }
    case sycl::detail::CGType::CopyUSM: {
      sycl::detail::CGCopyUSM *CopyA =
          static_cast<sycl::detail::CGCopyUSM *>(MCommandGroup.get());
      sycl::detail::CGCopyUSM *CopyB =
          static_cast<sycl::detail::CGCopyUSM *>(Node->MCommandGroup.get());
      return (CopyA->getSrc() == CopyB->getSrc()) &&
             (CopyA->getDst() == CopyB->getDst()) &&
             (CopyA->getLength() == CopyB->getLength());
    }
    case sycl::detail::CGType::CopyAccToAcc:
    case sycl::detail::CGType::CopyAccToPtr:
    case sycl::detail::CGType::CopyPtrToAcc: {
      sycl::detail::CGCopy *CopyA =
          static_cast<sycl::detail::CGCopy *>(MCommandGroup.get());
      sycl::detail::CGCopy *CopyB =
          static_cast<sycl::detail::CGCopy *>(Node->MCommandGroup.get());
      return (CopyA->getSrc() == CopyB->getSrc()) &&
             (CopyA->getDst() == CopyB->getDst());
    }
    default:
      assert(false && "Unexpected command group type!");
      return false;
    }
  }
 
  void printDotRecursive(std::fstream &Stream,
                         std::vector<node_impl *> &Visited, bool Verbose) {
    // if Node has been already visited, we skip it
    if (std::find(Visited.begin(), Visited.end(), this) != Visited.end())
      return;
 
    Visited.push_back(this);
 
    printDotCG(Stream, Verbose);
    for (const auto &Dep : MPredecessors) {
      auto NodeDep = Dep.lock();
      Stream << "  \"" << NodeDep.get() << "\" -> \"" << this << "\""
             << std::endl;
    }
 
    for (std::weak_ptr<node_impl> Succ : MSuccessors) {
      if (MPartitionNum == Succ.lock()->MPartitionNum)
        Succ.lock()->printDotRecursive(Stream, Visited, Verbose);
    }
  }
 
  bool isNDCopyNode() const {
    if ((MCGType != sycl::detail::CGType::CopyAccToAcc) &&
        (MCGType != sycl::detail::CGType::CopyAccToPtr) &&
        (MCGType != sycl::detail::CGType::CopyPtrToAcc)) {
      return false;
    }
 
    auto Copy = static_cast<sycl::detail::CGCopy *>(MCommandGroup.get());
    auto ReqSrc = static_cast<sycl::detail::Requirement *>(Copy->getSrc());
    auto ReqDst = static_cast<sycl::detail::Requirement *>(Copy->getDst());
    return (ReqSrc->MDims > 1) || (ReqDst->MDims > 1);
  }
 
  void updateAccessor(int ArgIndex, const sycl::detail::AccessorBaseHost *Acc) {
    auto &Args =
        static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get())->MArgs;
    auto NewAccImpl = sycl::detail::getSyclObjImpl(*Acc);
    for (auto &Arg : Args) {
      if (Arg.MIndex != ArgIndex) {
        continue;
      }
      assert(Arg.MType == sycl::detail::kernel_param_kind_t::kind_accessor);
 
      // Find old accessor in accessor storage and replace with new one
      if (static_cast<sycl::detail::SYCLMemObjT *>(NewAccImpl->MSYCLMemObj)
              ->needsWriteBack()) {
        throw sycl::exception(
            make_error_code(errc::invalid),
            "Accessors to buffers which have write_back enabled "
            "are not allowed to be used in command graphs.");
      }
 
      // All accessors passed to this function will be placeholders, so we must
      // perform steps similar to what happens when handler::require() is
      // called here.
      sycl::detail::Requirement *NewReq = NewAccImpl.get();
      if (NewReq->MAccessMode != sycl::access_mode::read) {
        auto SYCLMemObj =
            static_cast<sycl::detail::SYCLMemObjT *>(NewReq->MSYCLMemObj);
        SYCLMemObj->handleWriteAccessorCreation();
      }
 
      for (auto &Acc : MCommandGroup->getAccStorage()) {
        if (auto OldAcc =
                static_cast<sycl::detail::AccessorImplHost *>(Arg.MPtr);
            Acc.get() == OldAcc) {
          Acc = NewAccImpl;
        }
      }
 
      for (auto &Req : MCommandGroup->getRequirements()) {
        if (auto OldReq =
                static_cast<sycl::detail::AccessorImplHost *>(Arg.MPtr);
            Req == OldReq) {
          Req = NewReq;
        }
      }
      Arg.MPtr = NewAccImpl.get();
      break;
    }
  }
 
  void updateArgValue(int ArgIndex, const void *NewValue, size_t Size) {
 
    auto &Args =
        static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get())->MArgs;
    for (auto &Arg : Args) {
      if (Arg.MIndex != ArgIndex) {
        continue;
      }
      assert(Arg.MSize == static_cast<int>(Size));
      // MPtr may be a pointer into arg storage so we memcpy the contents of
      // NewValue rather than assign it directly
      std::memcpy(Arg.MPtr, NewValue, Size);
      break;
    }
  }
 
  template <int Dimensions>
  void updateNDRange(nd_range<Dimensions> ExecutionRange) {
    if (MCGType != sycl::detail::CGType::Kernel) {
      throw sycl::exception(
          sycl::errc::invalid,
          "Cannot update execution range of nodes which are not kernel nodes");
    }
    if (!MNDRangeUsed) {
      throw sycl::exception(sycl::errc::invalid,
                            "Cannot update node which was created with a "
                            "sycl::range with a sycl::nd_range");
    }
 
    auto &NDRDesc =
        static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get())
            ->MNDRDesc;
 
    if (NDRDesc.Dims != Dimensions) {
      throw sycl::exception(sycl::errc::invalid,
                            "Cannot update execution range of a node with an "
                            "execution range of different dimensions than what "
                            "the node was originall created with.");
    }
 
    NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
  }
 
  template <int Dimensions> void updateRange(range<Dimensions> ExecutionRange) {
    if (MCGType != sycl::detail::CGType::Kernel) {
      throw sycl::exception(
          sycl::errc::invalid,
          "Cannot update execution range of nodes which are not kernel nodes");
    }
    if (MNDRangeUsed) {
      throw sycl::exception(sycl::errc::invalid,
                            "Cannot update node which was created with a "
                            "sycl::nd_range with a sycl::range");
    }
 
    auto &NDRDesc =
        static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get())
            ->MNDRDesc;
 
    if (NDRDesc.Dims != Dimensions) {
      throw sycl::exception(sycl::errc::invalid,
                            "Cannot update execution range of a node with an "
                            "execution range of different dimensions than what "
                            "the node was originall created with.");
    }
 
    NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
  }
 
  void updateFromOtherNode(const std::shared_ptr<node_impl> &Other) {
    auto ExecCG =
        static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get());
    auto OtherExecCG =
        static_cast<sycl::detail::CGExecKernel *>(Other->MCommandGroup.get());
 
    ExecCG->MArgs = OtherExecCG->MArgs;
    ExecCG->MNDRDesc = OtherExecCG->MNDRDesc;
    ExecCG->getAccStorage() = OtherExecCG->getAccStorage();
    ExecCG->getRequirements() = OtherExecCG->getRequirements();
 
    auto &OldArgStorage = OtherExecCG->getArgsStorage();
    auto &NewArgStorage = ExecCG->getArgsStorage();
    // Rebuild the arg storage and update the args
    rebuildArgStorage(ExecCG->MArgs, OldArgStorage, NewArgStorage);
  }
 
  id_type getID() const { return MID; }
 
private:
  void rebuildArgStorage(std::vector<sycl::detail::ArgDesc> &Args,
                         const std::vector<std::vector<char>> &OldArgStorage,
                         std::vector<std::vector<char>> &NewArgStorage) const {
    // Clear the arg storage so we can rebuild it
    NewArgStorage.clear();
 
    // Loop over all the args, any std_layout ones need their pointers updated
    // to point to the new arg storage.
    for (auto &Arg : Args) {
      if (Arg.MType != sycl::detail::kernel_param_kind_t::kind_std_layout) {
        continue;
      }
      // Find which ArgStorage Arg.MPtr is pointing to
      for (auto &ArgStorage : OldArgStorage) {
        if (ArgStorage.data() != Arg.MPtr) {
          continue;
        }
        NewArgStorage.emplace_back(Arg.MSize);
        // Memcpy contents from old storage to new storage
        std::memcpy(NewArgStorage.back().data(), ArgStorage.data(), Arg.MSize);
        // Update MPtr to point to the new storage instead of the old
        Arg.MPtr = NewArgStorage.back().data();
 
        break;
      }
    }
  }
  // Gets the next unique identifier for a node, should only be used when
  // constructing nodes.
  static id_type getNextNodeID() {
    static id_type nextID = 0;
 
    // Return the value then increment the next ID
    return nextID++;
  }
 
  void printDotCG(std::ostream &Stream, bool Verbose) {
    Stream << "\"" << this << "\" [style=bold, label=\"";
 
    Stream << "ID = " << this << "\\n";
    Stream << "TYPE = ";
 
    switch (MCGType) {
    case sycl::detail::CGType::None:
      Stream << "None \\n";
      break;
    case sycl::detail::CGType::Kernel: {
      Stream << "CGExecKernel \\n";
      sycl::detail::CGExecKernel *Kernel =
          static_cast<sycl::detail::CGExecKernel *>(MCommandGroup.get());
      Stream << "NAME = " << Kernel->MKernelName << "\\n";
      if (Verbose) {
        Stream << "ARGS = \\n";
        for (size_t i = 0; i < Kernel->MArgs.size(); i++) {
          auto Arg = Kernel->MArgs[i];
          std::string Type = "Undefined";
          if (Arg.MType == sycl::detail::kernel_param_kind_t::kind_accessor) {
            Type = "Accessor";
          } else if (Arg.MType ==
                     sycl::detail::kernel_param_kind_t::kind_std_layout) {
            Type = "STD_Layout";
          } else if (Arg.MType ==
                     sycl::detail::kernel_param_kind_t::kind_sampler) {
            Type = "Sampler";
          } else if (Arg.MType ==
                     sycl::detail::kernel_param_kind_t::kind_pointer) {
            Type = "Pointer";
            auto Fill = Stream.fill();
            Stream << i << ") Type: " << Type << " Ptr: " << Arg.MPtr << "(0x"
                   << std::hex << std::setfill('0');
            for (int i = Arg.MSize - 1; i >= 0; --i) {
              Stream << std::setw(2)
                     << static_cast<int16_t>(
                            (static_cast<unsigned char *>(Arg.MPtr))[i]);
            }
            Stream.fill(Fill);
            Stream << std::dec << ")\\n";
            continue;
          } else if (Arg.MType == sycl::detail::kernel_param_kind_t::
                                      kind_specialization_constants_buffer) {
            Type = "Specialization Constants Buffer";
          } else if (Arg.MType ==
                     sycl::detail::kernel_param_kind_t::kind_stream) {
            Type = "Stream";
          } else if (Arg.MType ==
                     sycl::detail::kernel_param_kind_t::kind_invalid) {
            Type = "Invalid";
          }
          Stream << i << ") Type: " << Type << " Ptr: " << Arg.MPtr << "\\n";
        }
      }
      break;
    }
    case sycl::detail::CGType::CopyAccToPtr:
      Stream << "CGCopy Device-to-Host \\n";
      if (Verbose) {
        sycl::detail::CGCopy *Copy =
            static_cast<sycl::detail::CGCopy *>(MCommandGroup.get());
        Stream << "Src: " << Copy->getSrc() << " Dst: " << Copy->getDst()
               << "\\n";
      }
      break;
    case sycl::detail::CGType::CopyPtrToAcc:
      Stream << "CGCopy Host-to-Device \\n";
      if (Verbose) {
        sycl::detail::CGCopy *Copy =
            static_cast<sycl::detail::CGCopy *>(MCommandGroup.get());
        Stream << "Src: " << Copy->getSrc() << " Dst: " << Copy->getDst()
               << "\\n";
      }
      break;
    case sycl::detail::CGType::CopyAccToAcc:
      Stream << "CGCopy Device-to-Device \\n";
      if (Verbose) {
        sycl::detail::CGCopy *Copy =
            static_cast<sycl::detail::CGCopy *>(MCommandGroup.get());
        Stream << "Src: " << Copy->getSrc() << " Dst: " << Copy->getDst()
               << "\\n";
      }
      break;
    case sycl::detail::CGType::Fill:
      Stream << "CGFill \\n";
      if (Verbose) {
        sycl::detail::CGFill *Fill =
            static_cast<sycl::detail::CGFill *>(MCommandGroup.get());
        Stream << "Ptr: " << Fill->MPtr << "\\n";
      }
      break;
    case sycl::detail::CGType::UpdateHost:
      Stream << "CGCUpdateHost \\n";
      if (Verbose) {
        sycl::detail::CGUpdateHost *Host =
            static_cast<sycl::detail::CGUpdateHost *>(MCommandGroup.get());
        Stream << "Ptr: " << Host->getReqToUpdate() << "\\n";
      }
      break;
    case sycl::detail::CGType::CopyUSM:
      Stream << "CGCopyUSM \\n";
      if (Verbose) {
        sycl::detail::CGCopyUSM *CopyUSM =
            static_cast<sycl::detail::CGCopyUSM *>(MCommandGroup.get());
        Stream << "Src: " << CopyUSM->getSrc() << " Dst: " << CopyUSM->getDst()
               << " Length: " << CopyUSM->getLength() << "\\n";
      }
      break;
    case sycl::detail::CGType::FillUSM:
      Stream << "CGFillUSM \\n";
      if (Verbose) {
        sycl::detail::CGFillUSM *FillUSM =
            static_cast<sycl::detail::CGFillUSM *>(MCommandGroup.get());
        Stream << "Dst: " << FillUSM->getDst()
               << " Length: " << FillUSM->getLength() << " Pattern: ";
        for (auto byte : FillUSM->getPattern())
          Stream << byte;
        Stream << "\\n";
      }
      break;
    case sycl::detail::CGType::PrefetchUSM:
      Stream << "CGPrefetchUSM \\n";
      if (Verbose) {
        sycl::detail::CGPrefetchUSM *Prefetch =
            static_cast<sycl::detail::CGPrefetchUSM *>(MCommandGroup.get());
        Stream << "Dst: " << Prefetch->getDst()
               << " Length: " << Prefetch->getLength() << "\\n";
      }
      break;
    case sycl::detail::CGType::AdviseUSM:
      Stream << "CGAdviseUSM \\n";
      if (Verbose) {
        sycl::detail::CGAdviseUSM *AdviseUSM =
            static_cast<sycl::detail::CGAdviseUSM *>(MCommandGroup.get());
        Stream << "Dst: " << AdviseUSM->getDst()
               << " Length: " << AdviseUSM->getLength() << "\\n";
      }
      break;
    case sycl::detail::CGType::CodeplayHostTask:
      Stream << "CGHostTask \\n";
      break;
    case sycl::detail::CGType::Barrier:
      Stream << "CGBarrier \\n";
      break;
    case sycl::detail::CGType::Copy2DUSM:
      Stream << "CGCopy2DUSM \\n";
      if (Verbose) {
        sycl::detail::CGCopy2DUSM *Copy2DUSM =
            static_cast<sycl::detail::CGCopy2DUSM *>(MCommandGroup.get());
        Stream << "Src:" << Copy2DUSM->getSrc()
               << " Dst: " << Copy2DUSM->getDst() << "\\n";
      }
      break;
    case sycl::detail::CGType::Fill2DUSM:
      Stream << "CGFill2DUSM \\n";
      if (Verbose) {
        sycl::detail::CGFill2DUSM *Fill2DUSM =
            static_cast<sycl::detail::CGFill2DUSM *>(MCommandGroup.get());
        Stream << "Dst: " << Fill2DUSM->getDst() << "\\n";
      }
      break;
    case sycl::detail::CGType::Memset2DUSM:
      Stream << "CGMemset2DUSM \\n";
      if (Verbose) {
        sycl::detail::CGMemset2DUSM *Memset2DUSM =
            static_cast<sycl::detail::CGMemset2DUSM *>(MCommandGroup.get());
        Stream << "Dst: " << Memset2DUSM->getDst() << "\\n";
      }
      break;
    case sycl::detail::CGType::ReadWriteHostPipe:
      Stream << "CGReadWriteHostPipe \\n";
      break;
    case sycl::detail::CGType::CopyToDeviceGlobal:
      Stream << "CGCopyToDeviceGlobal \\n";
      if (Verbose) {
        sycl::detail::CGCopyToDeviceGlobal *CopyToDeviceGlobal =
            static_cast<sycl::detail::CGCopyToDeviceGlobal *>(
                MCommandGroup.get());
        Stream << "Src: " << CopyToDeviceGlobal->getSrc()
               << " Dst: " << CopyToDeviceGlobal->getDeviceGlobalPtr() << "\\n";
      }
      break;
    case sycl::detail::CGType::CopyFromDeviceGlobal:
      Stream << "CGCopyFromDeviceGlobal \\n";
      if (Verbose) {
        sycl::detail::CGCopyFromDeviceGlobal *CopyFromDeviceGlobal =
            static_cast<sycl::detail::CGCopyFromDeviceGlobal *>(
                MCommandGroup.get());
        Stream << "Src: " << CopyFromDeviceGlobal->getDeviceGlobalPtr()
               << " Dst: " << CopyFromDeviceGlobal->getDest() << "\\n";
      }
      break;
    case sycl::detail::CGType::ExecCommandBuffer:
      Stream << "CGExecCommandBuffer \\n";
      break;
    default:
      Stream << "Other \\n";
      break;
    }
    Stream << "\"];" << std::endl;
  }
 
  template <typename CGT> std::unique_ptr<CGT> createCGCopy() const {
    return std::make_unique<CGT>(*static_cast<CGT *>(MCommandGroup.get()));
  }
};
 
class partition {
public:
  partition() : MSchedule(), MCommandBuffers() {}
 
  std::set<std::weak_ptr<node_impl>, std::owner_less<std::weak_ptr<node_impl>>>
      MRoots;
  std::list<std::shared_ptr<node_impl>> MSchedule;
  std::unordered_map<sycl::device, ur_exp_command_buffer_handle_t>
      MCommandBuffers;
  std::vector<std::shared_ptr<partition>> MPredecessors;
  bool MIsInOrderGraph = false;
 
  bool isHostTask() const {
    return (MRoots.size() && ((*MRoots.begin()).lock()->MCGType ==
                              sycl::detail::CGType::CodeplayHostTask));
  }
 
  bool checkIfGraphIsSinglePath() {
    if (MRoots.size() > 1) {
      return false;
    }
    for (const auto &Node : MSchedule) {
      // In version 1.3.28454 of the L0 driver, 2D Copy ops cannot not
      // be enqueued in an in-order cmd-list (causing execution to stall).
      // The 2D Copy test should be removed from here when the bug is fixed.
      if ((Node->MSuccessors.size() > 1) || (Node->isNDCopyNode())) {
        return false;
      }
    }
 
    return true;
  }
 
  void schedule();
};
 
class graph_impl {
public:
  using ReadLock = std::shared_lock<std::shared_mutex>;
  using WriteLock = std::unique_lock<std::shared_mutex>;
 
  mutable std::shared_mutex MMutex;
 
  graph_impl(const sycl::context &SyclContext, const sycl::device &SyclDevice,
             const sycl::property_list &PropList = {})
      : MContext(SyclContext), MDevice(SyclDevice), MRecordingQueues(),
        MEventsMap(), MInorderQueueMap() {
    if (PropList.has_property<property::graph::no_cycle_check>()) {
      MSkipCycleChecks = true;
    }
    if (PropList
            .has_property<property::graph::assume_buffer_outlives_graph>()) {
      MAllowBuffers = true;
    }
 
    if (!SyclDevice.has(aspect::ext_oneapi_limited_graph) &&
        !SyclDevice.has(aspect::ext_oneapi_graph)) {
      std::stringstream Stream;
      Stream << SyclDevice.get_backend();
      std::string BackendString = Stream.str();
      throw sycl::exception(
          sycl::make_error_code(errc::invalid),
          BackendString + " backend is not supported by SYCL Graph extension.");
    }
  }
 
  ~graph_impl();
 
  void removeRoot(const std::shared_ptr<node_impl> &Root);
 
  std::shared_ptr<node_impl>
  add(node_type NodeType, std::unique_ptr<sycl::detail::CG> CommandGroup,
      const std::vector<std::shared_ptr<node_impl>> &Dep = {});
 
  std::shared_ptr<node_impl>
  add(const std::shared_ptr<graph_impl> &Impl,
      std::function<void(handler &)> CGF,
      const std::vector<sycl::detail::ArgDesc> &Args,
      const std::vector<std::shared_ptr<node_impl>> &Dep = {});
 
  std::shared_ptr<node_impl>
  add(const std::shared_ptr<graph_impl> &Impl,
      const std::vector<std::shared_ptr<node_impl>> &Dep = {});
 
  std::shared_ptr<node_impl>
  add(const std::shared_ptr<graph_impl> &Impl,
      const std::vector<sycl::detail::EventImplPtr> Events);
 
  void
  addQueue(const std::shared_ptr<sycl::detail::queue_impl> &RecordingQueue) {
    MRecordingQueues.insert(RecordingQueue);
  }
 
  void
  removeQueue(const std::shared_ptr<sycl::detail::queue_impl> &RecordingQueue) {
    MRecordingQueues.erase(RecordingQueue);
  }
 
  bool clearQueues();
 
  void addEventForNode(std::shared_ptr<graph_impl> GraphImpl,
                       std::shared_ptr<sycl::detail::event_impl> EventImpl,
                       std::shared_ptr<node_impl> NodeImpl) {
    if (!(EventImpl->getCommandGraph()))
      EventImpl->setCommandGraph(GraphImpl);
    MEventsMap[EventImpl] = NodeImpl;
  }
 
  std::shared_ptr<sycl::detail::event_impl>
  getEventForNode(std::shared_ptr<node_impl> NodeImpl) const {
    ReadLock Lock(MMutex);
    if (auto EventImpl = std::find_if(
            MEventsMap.begin(), MEventsMap.end(),
            [NodeImpl](auto &it) { return it.second == NodeImpl; });
        EventImpl != MEventsMap.end()) {
      return EventImpl->first;
    }
 
    throw sycl::exception(
        sycl::make_error_code(errc::invalid),
        "No event has been recorded for the specified graph node");
  }
 
  std::shared_ptr<node_impl>
  getNodeForEvent(std::shared_ptr<sycl::detail::event_impl> EventImpl) {
    ReadLock Lock(MMutex);
 
    if (auto NodeFound = MEventsMap.find(EventImpl);
        NodeFound != std::end(MEventsMap)) {
      return NodeFound->second;
    }
 
    throw sycl::exception(
        sycl::make_error_code(errc::invalid),
        "No node in this graph is associated with this event");
  }
 
  sycl::context getContext() const { return MContext; }
 
  sycl::device getDevice() const { return MDevice; }
 
  std::set<std::weak_ptr<node_impl>, std::owner_less<std::weak_ptr<node_impl>>>
      MRoots;
 
  std::vector<std::shared_ptr<node_impl>> MNodeStorage;
 
  std::shared_ptr<node_impl>
  getLastInorderNode(std::shared_ptr<sycl::detail::queue_impl> Queue) {
    std::weak_ptr<sycl::detail::queue_impl> QueueWeakPtr(Queue);
    if (0 == MInorderQueueMap.count(QueueWeakPtr)) {
      return {};
    }
    return MInorderQueueMap[QueueWeakPtr];
  }
 
  void setLastInorderNode(std::shared_ptr<sycl::detail::queue_impl> Queue,
                          std::shared_ptr<node_impl> Node) {
    std::weak_ptr<sycl::detail::queue_impl> QueueWeakPtr(Queue);
    MInorderQueueMap[QueueWeakPtr] = Node;
  }
 
  void printGraphAsDot(const std::string FilePath, bool Verbose) const {
    std::vector<node_impl *> VisitedNodes;
 
    std::fstream Stream(FilePath, std::ios::out);
    Stream << "digraph dot {" << std::endl;
 
    for (std::weak_ptr<node_impl> Node : MRoots)
      Node.lock()->printDotRecursive(Stream, VisitedNodes, Verbose);
 
    Stream << "}" << std::endl;
 
    Stream.close();
  }
 
  void makeEdge(std::shared_ptr<node_impl> Src,
                std::shared_ptr<node_impl> Dest);
 
  void throwIfGraphRecordingQueue(const std::string ExceptionMsg) const {
    if (MRecordingQueues.size()) {
      throw sycl::exception(make_error_code(sycl::errc::invalid),
                            ExceptionMsg +
                                " cannot be called when a queue "
                                "is currently recording commands to a graph.");
    }
  }
 
  static bool checkNodeRecursive(const std::shared_ptr<node_impl> &NodeA,
                                 const std::shared_ptr<node_impl> &NodeB) {
    size_t FoundCnt = 0;
    for (std::weak_ptr<node_impl> &SuccA : NodeA->MSuccessors) {
      for (std::weak_ptr<node_impl> &SuccB : NodeB->MSuccessors) {
        if (NodeA->isSimilar(NodeB) &&
            checkNodeRecursive(SuccA.lock(), SuccB.lock())) {
          FoundCnt++;
          break;
        }
      }
    }
    if (FoundCnt != NodeA->MSuccessors.size()) {
      return false;
    }
 
    return true;
  }
 
  bool hasSimilarStructure(std::shared_ptr<detail::graph_impl> Graph,
                           bool DebugPrint = false) const {
    if (this == Graph.get())
      return true;
 
    if (MContext != Graph->MContext) {
      if (DebugPrint) {
        throw sycl::exception(sycl::make_error_code(errc::invalid),
                              "MContext are not the same.");
      }
      return false;
    }
 
    if (MDevice != Graph->MDevice) {
      if (DebugPrint) {
        throw sycl::exception(sycl::make_error_code(errc::invalid),
                              "MDevice are not the same.");
      }
      return false;
    }
 
    if (MEventsMap.size() != Graph->MEventsMap.size()) {
      if (DebugPrint) {
        throw sycl::exception(sycl::make_error_code(errc::invalid),
                              "MEventsMap sizes are not the same.");
      }
      return false;
    }
 
    if (MInorderQueueMap.size() != Graph->MInorderQueueMap.size()) {
      if (DebugPrint) {
        throw sycl::exception(sycl::make_error_code(errc::invalid),
                              "MInorderQueueMap sizes are not the same.");
      }
      return false;
    }
 
    if (MRoots.size() != Graph->MRoots.size()) {
      if (DebugPrint) {
        throw sycl::exception(sycl::make_error_code(errc::invalid),
                              "MRoots sizes are not the same.");
      }
      return false;
    }
 
    size_t RootsFound = 0;
    for (std::weak_ptr<node_impl> NodeA : MRoots) {
      for (std::weak_ptr<node_impl> NodeB : Graph->MRoots) {
        auto NodeALocked = NodeA.lock();
        auto NodeBLocked = NodeB.lock();
 
        if (NodeALocked->isSimilar(NodeBLocked)) {
          if (checkNodeRecursive(NodeALocked, NodeBLocked)) {
            RootsFound++;
            break;
          }
        }
      }
    }
 
    if (RootsFound != MRoots.size()) {
      if (DebugPrint) {
        throw sycl::exception(sycl::make_error_code(errc::invalid),
                              "Root Nodes do NOT match.");
      }
      return false;
    }
 
    return true;
  }
 
  size_t getNumberOfNodes() const { return MNodeStorage.size(); }
 
  std::vector<sycl::detail::EventImplPtr>
  getExitNodesEvents(std::weak_ptr<sycl::detail::queue_impl> Queue);
 
  void setBarrierDep(std::weak_ptr<sycl::detail::queue_impl> Queue,
                     std::shared_ptr<node_impl> BarrierNodeImpl) {
    MBarrierDependencyMap[Queue] = BarrierNodeImpl;
  }
 
  std::shared_ptr<node_impl>
  getBarrierDep(std::weak_ptr<sycl::detail::queue_impl> Queue) {
    return MBarrierDependencyMap[Queue];
  }
 
private:
  void
  searchDepthFirst(std::function<bool(std::shared_ptr<node_impl> &,
                                      std::deque<std::shared_ptr<node_impl>> &)>
                       NodeFunc);
 
  bool checkForCycles();
 
  void addRoot(const std::shared_ptr<node_impl> &Root);
 
  std::shared_ptr<node_impl>
  addNodesToExits(const std::shared_ptr<graph_impl> &Impl,
                  const std::list<std::shared_ptr<node_impl>> &NodeList);
 
  void addDepsToNode(std::shared_ptr<node_impl> Node,
                     const std::vector<std::shared_ptr<node_impl>> &Deps) {
    if (!Deps.empty()) {
      for (auto &N : Deps) {
        N->registerSuccessor(Node, N);
        this->removeRoot(Node);
      }
    } else {
      this->addRoot(Node);
    }
  }
 
  sycl::context MContext;
  sycl::device MDevice;
  std::set<std::weak_ptr<sycl::detail::queue_impl>,
           std::owner_less<std::weak_ptr<sycl::detail::queue_impl>>>
      MRecordingQueues;
  std::unordered_map<std::shared_ptr<sycl::detail::event_impl>,
                     std::shared_ptr<node_impl>>
      MEventsMap;
  std::map<std::weak_ptr<sycl::detail::queue_impl>, std::shared_ptr<node_impl>,
           std::owner_less<std::weak_ptr<sycl::detail::queue_impl>>>
      MInorderQueueMap;
  bool MSkipCycleChecks = false;
  std::set<sycl::detail::SYCLMemObjT *> MMemObjs;
 
  bool MAllowBuffers = false;
 
  std::map<std::weak_ptr<sycl::detail::queue_impl>, std::shared_ptr<node_impl>,
           std::owner_less<std::weak_ptr<sycl::detail::queue_impl>>>
      MBarrierDependencyMap;
};
 
class exec_graph_impl {
public:
  using ReadLock = std::shared_lock<std::shared_mutex>;
  using WriteLock = std::unique_lock<std::shared_mutex>;
 
  mutable std::shared_mutex MMutex;
 
  exec_graph_impl(sycl::context Context,
                  const std::shared_ptr<graph_impl> &GraphImpl,
                  const property_list &PropList);
 
  ~exec_graph_impl();
 
  void makePartitions();
 
  sycl::event enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue,
                      sycl::detail::CG::StorageInitHelper CGData);
 
  void createCommandBuffers(sycl::device Device,
                            std::shared_ptr<partition> &Partition);
 
  sycl::device getDevice() const { return MDevice; }
 
  sycl::context getContext() const { return MContext; }
 
  const std::list<std::shared_ptr<node_impl>> &getSchedule() const {
    return MSchedule;
  }
 
  const std::shared_ptr<graph_impl> &getGraphImpl() const { return MGraphImpl; }
 
  const std::vector<std::shared_ptr<partition>> &getPartitions() const {
    return MPartitions;
  }
 
  bool previousSubmissionCompleted() const {
    for (auto Event : MExecutionEvents) {
      if (!Event->isCompleted()) {
        return false;
      }
    }
    return true;
  }
 
  std::vector<sycl::detail::AccessorImplHost *> getRequirements() const {
    return MRequirements;
  }
 
  void update(std::shared_ptr<graph_impl> GraphImpl);
  void update(std::shared_ptr<node_impl> Node);
  void update(const std::vector<std::shared_ptr<node_impl>> Nodes);
 
  void updateImpl(std::shared_ptr<node_impl> NodeImpl);
 
private:
  ur_exp_command_buffer_sync_point_t
  enqueueNode(sycl::context Ctx, sycl::detail::DeviceImplPtr DeviceImpl,
              ur_exp_command_buffer_handle_t CommandBuffer,
              std::shared_ptr<node_impl> Node);
 
  ur_exp_command_buffer_sync_point_t
  enqueueNodeDirect(sycl::context Ctx, sycl::detail::DeviceImplPtr DeviceImpl,
                    ur_exp_command_buffer_handle_t CommandBuffer,
                    std::shared_ptr<node_impl> Node);
 
  void findRealDeps(std::vector<ur_exp_command_buffer_sync_point_t> &Deps,
                    std::shared_ptr<node_impl> CurrentNode,
                    int ReferencePartitionNum);
 
  void duplicateNodes();
 
  void printGraphAsDot(const std::string FilePath, bool Verbose) const {
    std::vector<node_impl *> VisitedNodes;
 
    std::fstream Stream(FilePath, std::ios::out);
    Stream << "digraph dot {" << std::endl;
 
    std::vector<std::shared_ptr<node_impl>> Roots;
    for (auto &Node : MNodeStorage) {
      if (Node->MPredecessors.size() == 0) {
        Roots.push_back(Node);
      }
    }
 
    for (std::shared_ptr<node_impl> Node : Roots)
      Node->printDotRecursive(Stream, VisitedNodes, Verbose);
 
    Stream << "}" << std::endl;
 
    Stream.close();
  }
 
  std::list<std::shared_ptr<node_impl>> MSchedule;
  std::shared_ptr<graph_impl> MGraphImpl;
  std::unordered_map<std::shared_ptr<node_impl>,
                     ur_exp_command_buffer_sync_point_t>
      MSyncPoints;
  std::unordered_map<std::shared_ptr<node_impl>, int> MPartitionNodes;
  sycl::device MDevice;
  sycl::context MContext;
  std::vector<sycl::detail::AccessorImplHost *> MRequirements;
  std::vector<sycl::detail::AccessorImplPtr> MAccessors;
  std::vector<sycl::detail::EventImplPtr> MExecutionEvents;
  std::vector<std::shared_ptr<partition>> MPartitions;
  std::vector<std::shared_ptr<node_impl>> MNodeStorage;
  std::unordered_map<std::shared_ptr<node_impl>,
                     ur_exp_command_buffer_command_handle_t>
      MCommandMap;
  bool MIsUpdatable;
  bool MEnableProfiling;
 
  // Stores a cache of node ids from modifiable graph nodes to the companion
  // node(s) in this graph. Used for quick access when updating this graph.
  std::multimap<node_impl::id_type, std::shared_ptr<node_impl>> MIDCache;
};
 
class dynamic_parameter_impl {
public:
  dynamic_parameter_impl(std::shared_ptr<graph_impl> GraphImpl,
                         size_t ParamSize, const void *Data)
      : MGraph(GraphImpl), MValueStorage(ParamSize) {
    std::memcpy(MValueStorage.data(), Data, ParamSize);
  }
 
  void registerNode(std::shared_ptr<node_impl> NodeImpl, int ArgIndex) {
    MNodes.emplace_back(NodeImpl, ArgIndex);
  }
 
  void *getValue() { return MValueStorage.data(); }
 
  void updateValue(const void *NewValue, size_t Size) {
    for (auto &[NodeWeak, ArgIndex] : MNodes) {
      auto NodeShared = NodeWeak.lock();
      if (NodeShared) {
        NodeShared->updateArgValue(ArgIndex, NewValue, Size);
      }
    }
    std::memcpy(MValueStorage.data(), NewValue, Size);
  }
 
  void updateAccessor(const sycl::detail::AccessorBaseHost *Acc) {
    for (auto &[NodeWeak, ArgIndex] : MNodes) {
      auto NodeShared = NodeWeak.lock();
      // Should we fail here if the node isn't alive anymore?
      if (NodeShared) {
        NodeShared->updateAccessor(ArgIndex, Acc);
      }
    }
    std::memcpy(MValueStorage.data(), Acc,
                sizeof(sycl::detail::AccessorBaseHost));
  }
 
  // Weak ptrs to node_impls which will be updated
  std::vector<std::pair<std::weak_ptr<node_impl>, int>> MNodes;
 
  std::shared_ptr<graph_impl> MGraph;
  std::vector<std::byte> MValueStorage;
};
 
} // namespace detail
} // namespace experimental
} // namespace oneapi
} // namespace ext
} // namespace _V1
} // namespace sycl