commands.hpp

Go to the documentation of this file.

//==-------------- commands.hpp - SYCL standard header file ----------------==//
//
// 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 <atomic>
#include <cstdint>
#include <deque>
#include <memory>
#include <optional>
#include <set>
#include <unordered_set>
#include <vector>
 
#include <detail/accessor_impl.hpp>
#include <detail/cg.hpp>
#include <detail/event_impl.hpp>
#include <detail/program_manager/program_manager.hpp>
#include <sycl/access/access.hpp>
 
namespace sycl {
inline namespace _V1 {
 
namespace ext::oneapi::experimental::detail {
class exec_graph_impl;
class node_impl;
} // namespace ext::oneapi::experimental::detail
namespace detail {
 
#ifdef XPTI_ENABLE_INSTRUMENTATION
void emitInstrumentationGeneral(uint32_t StreamID, uint64_t InstanceID,
                                xpti_td *TraceEvent, uint16_t Type,
                                const void *Addr);
#endif
RTDeviceBinaryImage *
retrieveAMDGCNOrNVPTXKernelBinary(const DeviceImplPtr DeviceImpl,
                                  const std::string &KernelName);
 
class queue_impl;
class event_impl;
class context_impl;
class DispatchHostTask;
 
using QueueImplPtr = std::shared_ptr<detail::queue_impl>;
using EventImplPtr = std::shared_ptr<detail::event_impl>;
using ContextImplPtr = std::shared_ptr<detail::context_impl>;
using StreamImplPtr = std::shared_ptr<detail::stream_impl>;
 
class Command;
class AllocaCommand;
class AllocaCommandBase;
class ReleaseCommand;
class ExecCGCommand;
class EmptyCommand;
 
enum BlockingT { NON_BLOCKING = 0, BLOCKING };
 
struct EnqueueResultT {
  enum ResultT {
    SyclEnqueueReady,
    SyclEnqueueSuccess,
    SyclEnqueueBlocked,
    SyclEnqueueFailed
  };
  EnqueueResultT(ResultT Result = SyclEnqueueSuccess, Command *Cmd = nullptr,
                 ur_result_t ErrCode = UR_RESULT_SUCCESS)
      : MResult(Result), MCmd(Cmd), MErrCode(ErrCode) {}
  ResultT MResult;
  Command *MCmd;
  ur_result_t MErrCode;
};
 
struct DepDesc {
  DepDesc(Command *DepCommand, const Requirement *Req,
          AllocaCommandBase *AllocaCmd)
      : MDepCommand(DepCommand), MDepRequirement(Req), MAllocaCmd(AllocaCmd) {}
 
  friend bool operator<(const DepDesc &Lhs, const DepDesc &Rhs) {
    return std::tie(Lhs.MDepRequirement, Lhs.MDepCommand) <
           std::tie(Rhs.MDepRequirement, Rhs.MDepCommand);
  }
 
  Command *MDepCommand = nullptr;
  const Requirement *MDepRequirement = nullptr;
  AllocaCommandBase *MAllocaCmd = nullptr;
};
 
class Command {
public:
  enum CommandType {
    RUN_CG,
    COPY_MEMORY,
    ALLOCA,
    ALLOCA_SUB_BUF,
    RELEASE,
    MAP_MEM_OBJ,
    UNMAP_MEM_OBJ,
    UPDATE_REQUIREMENT,
    EMPTY_TASK,
    HOST_TASK,
    FUSION,
    EXEC_CMD_BUFFER,
    UPDATE_CMD_BUFFER
  };
 
  Command(
      CommandType Type, QueueImplPtr Queue,
      ur_exp_command_buffer_handle_t CommandBuffer = nullptr,
      const std::vector<ur_exp_command_buffer_sync_point_t> &SyncPoints = {});
 
  [[nodiscard]] Command *addDep(DepDesc NewDep,
                                std::vector<Command *> &ToCleanUp);
 
  [[nodiscard]] Command *addDep(EventImplPtr Event,
                                std::vector<Command *> &ToCleanUp);
 
  void addUser(Command *NewUser) { MUsers.insert(NewUser); }
 
  CommandType getType() const { return MType; }
 
  virtual bool enqueue(EnqueueResultT &EnqueueResult, BlockingT Blocking,
                       std::vector<Command *> &ToCleanUp);
 
  bool isFinished();
 
  bool isSuccessfullyEnqueued() const {
    return MEnqueueStatus == EnqueueResultT::SyclEnqueueSuccess;
  }
 
  // Shows that command could not be enqueued, now it may be true for empty task
  // only
  bool isEnqueueBlocked() const {
    return MIsBlockable && MEnqueueStatus == EnqueueResultT::SyclEnqueueBlocked;
  }
  // Shows that command could be enqueued, but is blocking enqueue of all
  // commands depending on it. Regular usage - host task.
  bool isBlocking() const { return isHostTask() && !MEvent->isCompleted(); }
 
  void addBlockedUserUnique(const EventImplPtr &NewUser) {
    if (std::find(MBlockedUsers.begin(), MBlockedUsers.end(), NewUser) !=
        MBlockedUsers.end())
      return;
    MBlockedUsers.push_back(NewUser);
  }
 
  const QueueImplPtr &getQueue() const { return MQueue; }
 
  const EventImplPtr &getEvent() const { return MEvent; }
 
  // Methods needed to support SYCL instrumentation
 
  void emitInstrumentationDataProxy();
  virtual void emitInstrumentationData() = 0;
  void resolveReleaseDependencies(std::set<Command *> &list);
  void emitEdgeEventForCommandDependence(
      Command *Cmd, void *ObjAddr, bool IsCommand,
      std::optional<access::mode> AccMode = std::nullopt);
  void emitEdgeEventForEventDependence(Command *Cmd,
                                       ur_event_handle_t &EventAddr);
  void emitEnqueuedEventSignal(ur_event_handle_t &UrEventAddr);
  uint64_t makeTraceEventProlog(void *MAddress);
  void makeTraceEventEpilog();
  void emitInstrumentation(uint16_t Type, const char *Txt = nullptr);
 
  // End Methods needed to support SYCL instrumentation
 
  virtual void printDot(std::ostream &Stream) const = 0;
 
  virtual const Requirement *getRequirement() const {
    assert(false && "Internal Error. The command has no stored requirement");
    return nullptr;
  }
 
  virtual ~Command() { MEvent->cleanDepEventsThroughOneLevel(); }
 
  const char *getBlockReason() const;
 
  virtual ContextImplPtr getWorkerContext() const;
 
  virtual bool producesPiEvent() const;
 
  virtual bool supportsPostEnqueueCleanup() const;
 
  virtual bool readyForCleanup() const;
 
  std::vector<ur_event_handle_t>
  getUrEvents(const std::vector<EventImplPtr> &EventImpls) const;
  std::vector<ur_event_handle_t>
  getUrEventsBlocking(const std::vector<EventImplPtr> &EventImpls) const;
 
  bool isHostTask() const;
 
  bool isFusable() const;
 
protected:
  QueueImplPtr MQueue;
  EventImplPtr MEvent;
  QueueImplPtr MWorkerQueue;
 
  std::vector<EventImplPtr> &MPreparedDepsEvents;
  std::vector<EventImplPtr> &MPreparedHostDepsEvents;
 
  void waitForEvents(QueueImplPtr Queue, std::vector<EventImplPtr> &RawEvents,
                     ur_event_handle_t &Event);
 
  void waitForPreparedHostEvents() const;
 
  [[nodiscard]] Command *processDepEvent(EventImplPtr DepEvent,
                                         const DepDesc &Dep,
                                         std::vector<Command *> &ToCleanUp);
 
  virtual ur_result_t enqueueImp() = 0;
 
  CommandType MType;
  std::mutex MEnqueueMtx;
 
  friend class DispatchHostTask;
 
public:
  const std::vector<EventImplPtr> &getPreparedHostDepsEvents() const {
    return MPreparedHostDepsEvents;
  }
 
  const std::vector<EventImplPtr> &getPreparedDepsEvents() const {
    return MPreparedDepsEvents;
  }
 
  // XPTI instrumentation. Copy code location details to the internal struct.
  // Memory is allocated in this method and released in destructor.
  void copySubmissionCodeLocation();
 
  void clearAllDependencies() {
    MPreparedDepsEvents.clear();
    MPreparedHostDepsEvents.clear();
    MDeps.clear();
  }
 
  std::vector<DepDesc> MDeps;
  std::unordered_set<Command *> MUsers;
  bool MIsBlockable = false;
  unsigned MLeafCounter = 0;
 
  struct Marks {
    bool MVisited = false;
    bool MToBeDeleted = false;
  };
  Marks MMarks;
 
  enum class BlockReason : int { HostAccessor = 0, HostTask };
 
  // Only have reasonable value while MIsBlockable is true
  BlockReason MBlockReason;
 
  std::atomic<EnqueueResultT::ResultT> MEnqueueStatus;
 
  // All member variables defined here are needed for the SYCL instrumentation
  // layer. Do not guard these variables below with XPTI_ENABLE_INSTRUMENTATION
  // to ensure we have the same object layout when the macro in the library and
  // SYCL app are not the same.
 
  void *MTraceEvent = nullptr;
  int32_t MStreamID = -1;
  void *MAddress = nullptr;
  std::string MAddressString;
  std::string MCommandNodeType;
  std::string MCommandName;
  bool MTraceEventPrologComplete = false;
  bool MFirstInstance = false;
  uint64_t MInstanceID = 0;
  code_location MSubmissionCodeLocation;
  std::string MSubmissionFileName;
  std::string MSubmissionFunctionName;
 
  // This flag allows to control whether event should be set complete
  // after successfull enqueue of command. Event is considered as "host" event
  // if there is no backend representation of event (i.e. getHandleRef() return
  // reference to nullptr value). By default the flag is set to true due to most
  // of host operations are synchronous. The only asynchronous operation
  // currently is host-task.
  bool MShouldCompleteEventIfPossible = true;
 
  bool MMarkedForCleanup = false;
 
  std::vector<EventImplPtr> MBlockedUsers;
  std::mutex MBlockedUsersMutex;
 
protected:
  ur_exp_command_buffer_handle_t getCommandBuffer() const {
    return MCommandBuffer;
  }
 
  ur_exp_command_buffer_handle_t MCommandBuffer;
  std::vector<ur_exp_command_buffer_sync_point_t> MSyncPointDeps;
};
 
class EmptyCommand : public Command {
public:
  EmptyCommand();
 
  void printDot(std::ostream &Stream) const final;
  const Requirement *getRequirement() const final { return &MRequirements[0]; }
  void addRequirement(Command *DepCmd, AllocaCommandBase *AllocaCmd,
                      const Requirement *Req);
 
  void emitInstrumentationData() override;
 
  bool producesPiEvent() const final;
 
private:
  ur_result_t enqueueImp() final;
 
  // Employing deque here as it allows to push_back/emplace_back without
  // invalidation of pointer or reference to stored data item regardless of
  // iterator invalidation.
  std::deque<Requirement> MRequirements;
};
 
class ReleaseCommand : public Command {
public:
  ReleaseCommand(QueueImplPtr Queue, AllocaCommandBase *AllocaCmd);
 
  void printDot(std::ostream &Stream) const final;
  void emitInstrumentationData() override;
  bool producesPiEvent() const final;
  bool supportsPostEnqueueCleanup() const final;
  bool readyForCleanup() const final;
 
private:
  ur_result_t enqueueImp() final;
 
  AllocaCommandBase *MAllocaCmd = nullptr;
};
 
class AllocaCommandBase : public Command {
public:
  AllocaCommandBase(CommandType Type, QueueImplPtr Queue, Requirement Req,
                    AllocaCommandBase *LinkedAllocaCmd, bool IsConst);
 
  ReleaseCommand *getReleaseCmd() { return &MReleaseCmd; }
 
  SYCLMemObjI *getSYCLMemObj() const { return MRequirement.MSYCLMemObj; }
 
  virtual void *getMemAllocation() const = 0;
 
  const Requirement *getRequirement() const final { return &MRequirement; }
 
  void emitInstrumentationData() override;
 
  bool producesPiEvent() const final;
 
  bool supportsPostEnqueueCleanup() const final;
 
  bool readyForCleanup() const final;
 
  void *MMemAllocation = nullptr;
 
  AllocaCommandBase *MLinkedAllocaCmd = nullptr;
  bool MIsActive = true;
 
  bool MIsLeaderAlloca = true;
  // Indicates that the data in this allocation must not be modified
  bool MIsConst = false;
 
protected:
  Requirement MRequirement;
  ReleaseCommand MReleaseCmd;
};
 
class AllocaCommand : public AllocaCommandBase {
public:
  AllocaCommand(QueueImplPtr Queue, Requirement Req,
                bool InitFromUserData = true,
                AllocaCommandBase *LinkedAllocaCmd = nullptr,
                bool IsConst = false);
 
  void *getMemAllocation() const final { return MMemAllocation; }
  void printDot(std::ostream &Stream) const final;
  void emitInstrumentationData() override;
 
private:
  ur_result_t enqueueImp() final;
 
  bool MInitFromUserData = false;
};
 
class AllocaSubBufCommand : public AllocaCommandBase {
public:
  AllocaSubBufCommand(QueueImplPtr Queue, Requirement Req,
                      AllocaCommandBase *ParentAlloca,
                      std::vector<Command *> &ToEnqueue,
                      std::vector<Command *> &ToCleanUp);
 
  void *getMemAllocation() const final;
  void printDot(std::ostream &Stream) const final;
  AllocaCommandBase *getParentAlloca() { return MParentAlloca; }
  void emitInstrumentationData() override;
 
private:
  ur_result_t enqueueImp() final;
 
  AllocaCommandBase *MParentAlloca = nullptr;
};
 
class MapMemObject : public Command {
public:
  MapMemObject(AllocaCommandBase *SrcAllocaCmd, Requirement Req, void **DstPtr,
               QueueImplPtr Queue, access::mode MapMode);
 
  void printDot(std::ostream &Stream) const final;
  const Requirement *getRequirement() const final { return &MSrcReq; }
  void emitInstrumentationData() override;
 
private:
  ur_result_t enqueueImp() final;
 
  AllocaCommandBase *MSrcAllocaCmd = nullptr;
  Requirement MSrcReq;
  void **MDstPtr = nullptr;
  access::mode MMapMode;
};
 
class UnMapMemObject : public Command {
public:
  UnMapMemObject(AllocaCommandBase *DstAllocaCmd, Requirement Req,
                 void **SrcPtr, QueueImplPtr Queue);
 
  void printDot(std::ostream &Stream) const final;
  const Requirement *getRequirement() const final { return &MDstReq; }
  void emitInstrumentationData() override;
  bool producesPiEvent() const final;
 
private:
  ur_result_t enqueueImp() final;
 
  AllocaCommandBase *MDstAllocaCmd = nullptr;
  Requirement MDstReq;
  void **MSrcPtr = nullptr;
};
 
class MemCpyCommand : public Command {
public:
  MemCpyCommand(Requirement SrcReq, AllocaCommandBase *SrcAllocaCmd,
                Requirement DstReq, AllocaCommandBase *DstAllocaCmd,
                QueueImplPtr SrcQueue, QueueImplPtr DstQueue);
 
  void printDot(std::ostream &Stream) const final;
  const Requirement *getRequirement() const final { return &MDstReq; }
  void emitInstrumentationData() final;
  ContextImplPtr getWorkerContext() const final;
  bool producesPiEvent() const final;
 
private:
  ur_result_t enqueueImp() final;
 
  QueueImplPtr MSrcQueue;
  Requirement MSrcReq;
  AllocaCommandBase *MSrcAllocaCmd = nullptr;
  Requirement MDstReq;
  AllocaCommandBase *MDstAllocaCmd = nullptr;
};
 
class MemCpyCommandHost : public Command {
public:
  MemCpyCommandHost(Requirement SrcReq, AllocaCommandBase *SrcAllocaCmd,
                    Requirement DstReq, void **DstPtr, QueueImplPtr SrcQueue,
                    QueueImplPtr DstQueue);
 
  void printDot(std::ostream &Stream) const final;
  const Requirement *getRequirement() const final { return &MDstReq; }
  void emitInstrumentationData() final;
  ContextImplPtr getWorkerContext() const final;
 
private:
  ur_result_t enqueueImp() final;
 
  QueueImplPtr MSrcQueue;
  Requirement MSrcReq;
  AllocaCommandBase *MSrcAllocaCmd = nullptr;
  Requirement MDstReq;
  void **MDstPtr = nullptr;
};
 
ur_result_t enqueueReadWriteHostPipe(const QueueImplPtr &Queue,
                                     const std::string &PipeName, bool blocking,
                                     void *ptr, size_t size,
                                     std::vector<ur_event_handle_t> &RawEvents,
                                     const detail::EventImplPtr &OutEventImpl,
                                     bool read);
 
void enqueueImpKernel(
    const QueueImplPtr &Queue, NDRDescT &NDRDesc, std::vector<ArgDesc> &Args,
    const std::shared_ptr<detail::kernel_bundle_impl> &KernelBundleImplPtr,
    const std::shared_ptr<detail::kernel_impl> &MSyclKernel,
    const std::string &KernelName, std::vector<ur_event_handle_t> &RawEvents,
    const detail::EventImplPtr &Event,
    const std::function<void *(Requirement *Req)> &getMemAllocationFunc,
    ur_kernel_cache_config_t KernelCacheConfig, bool KernelIsCooperative,
    const bool KernelUsesClusterLaunch,
    const RTDeviceBinaryImage *BinImage = nullptr);
 
class KernelFusionCommand;
 
class ExecCGCommand : public Command {
public:
  ExecCGCommand(
      std::unique_ptr<detail::CG> CommandGroup, QueueImplPtr Queue,
      bool EventNeeded, ur_exp_command_buffer_handle_t CommandBuffer = nullptr,
      const std::vector<ur_exp_command_buffer_sync_point_t> &Dependencies = {});
 
  std::vector<std::shared_ptr<const void>> getAuxiliaryResources() const;
 
  void clearAuxiliaryResources();
 
  void printDot(std::ostream &Stream) const final;
  void emitInstrumentationData() final;
  std::string_view getTypeString() const;
 
  detail::CG &getCG() const { return *MCommandGroup; }
 
  // MEmptyCmd is only employed if this command refers to host-task.
  // The mechanism of lookup for single EmptyCommand amongst users of
  // host-task-representing command is unreliable. This unreliability roots in
  // the cleanup process.
  EmptyCommand *MEmptyCmd = nullptr;
 
  // MFusionCommand is employed to mark a CG command as part of a kernel fusion
  // and allows to refer back to the corresponding KernelFusionCommand if
  // necessary.
  KernelFusionCommand *MFusionCmd = nullptr;
 
  // MEventNeeded is true if the command needs to produce a valid event. The
  // implementation may elect to not produce events (native or SYCL) if this
  // is false.
  bool MEventNeeded = true;
 
  bool producesPiEvent() const final;
 
  bool supportsPostEnqueueCleanup() const final;
 
  bool readyForCleanup() const final;
 
private:
  ur_result_t enqueueImp() final;
  ur_result_t enqueueImpCommandBuffer();
  ur_result_t enqueueImpQueue();
 
  AllocaCommandBase *getAllocaForReq(Requirement *Req);
 
  std::unique_ptr<detail::CG> MCommandGroup;
 
  friend class Command;
};
 
// For XPTI instrumentation only.
// Method used to emit data in cases when we do not create node in graph.
// Very close to ExecCGCommand::emitInstrumentationData content.
#ifdef XPTI_ENABLE_INSTRUMENTATION
std::pair<xpti_td *, uint64_t> emitKernelInstrumentationData(
    int32_t StreamID, const std::shared_ptr<detail::kernel_impl> &SyclKernel,
    const detail::code_location &CodeLoc, const std::string &SyclKernelName,
    const QueueImplPtr &Queue, const NDRDescT &NDRDesc,
    const std::shared_ptr<detail::kernel_bundle_impl> &KernelBundleImplPtr,
    std::vector<ArgDesc> &CGArgs);
#endif
 
class UpdateHostRequirementCommand : public Command {
public:
  UpdateHostRequirementCommand(QueueImplPtr Queue, Requirement Req,
                               AllocaCommandBase *SrcAllocaCmd, void **DstPtr);
 
  void printDot(std::ostream &Stream) const final;
  const Requirement *getRequirement() const final { return &MDstReq; }
  void emitInstrumentationData() final;
 
private:
  ur_result_t enqueueImp() final;
 
  AllocaCommandBase *MSrcAllocaCmd = nullptr;
  Requirement MDstReq;
  void **MDstPtr = nullptr;
};
 
class KernelFusionCommand : public Command {
public:
  enum class FusionStatus { ACTIVE, CANCELLED, COMPLETE, DELETED };
 
  explicit KernelFusionCommand(QueueImplPtr Queue);
 
  void printDot(std::ostream &Stream) const final;
  void emitInstrumentationData() final;
  bool producesPiEvent() const final;
 
  std::vector<Command *> &auxiliaryCommands();
 
  void addToFusionList(ExecCGCommand *Kernel);
 
  std::vector<ExecCGCommand *> &getFusionList();
 
  void setFusionStatus(FusionStatus Status);
 
  void resetQueue();
 
  bool isActive() const { return MStatus == FusionStatus::ACTIVE; }
 
  bool readyForDeletion() const { return MStatus == FusionStatus::DELETED; }
 
private:
  ur_result_t enqueueImp() final;
 
  std::vector<ExecCGCommand *> MFusionList;
 
  std::vector<Command *> MAuxiliaryCommands;
 
  FusionStatus MStatus;
};
 
class UpdateCommandBufferCommand : public Command {
public:
  explicit UpdateCommandBufferCommand(
      QueueImplPtr Queue,
      ext::oneapi::experimental::detail::exec_graph_impl *Graph,
      std::vector<std::shared_ptr<ext::oneapi::experimental::detail::node_impl>>
          Nodes);
 
  void printDot(std::ostream &Stream) const final;
  void emitInstrumentationData() final;
  bool producesPiEvent() const final;
 
private:
  ur_result_t enqueueImp() final;
 
  ext::oneapi::experimental::detail::exec_graph_impl *MGraph;
  std::vector<std::shared_ptr<ext::oneapi::experimental::detail::node_impl>>
      MNodes;
};
 
// Enqueues a given kernel to a ur_exp_command_buffer_handle_t
ur_result_t enqueueImpCommandBufferKernel(
    context Ctx, DeviceImplPtr DeviceImpl,
    ur_exp_command_buffer_handle_t CommandBuffer,
    const CGExecKernel &CommandGroup,
    std::vector<ur_exp_command_buffer_sync_point_t> &SyncPoints,
    ur_exp_command_buffer_sync_point_t *OutSyncPoint,
    ur_exp_command_buffer_command_handle_t *OutCommand,
    const std::function<void *(Requirement *Req)> &getMemAllocationFunc);
 
// Sets arguments for a given kernel and device based on the argument type.
// Refactored from SetKernelParamsAndLaunch to allow it to be used in the graphs
// extension.
void SetArgBasedOnType(
    const detail::PluginPtr &Plugin, ur_kernel_handle_t Kernel,
    const std::shared_ptr<device_image_impl> &DeviceImageImpl,
    const std::function<void *(Requirement *Req)> &getMemAllocationFunc,
    const sycl::context &Context, detail::ArgDesc &Arg, size_t NextTrueIndex);
 
void applyFuncOnFilteredArgs(
    const KernelArgMask *EliminatedArgMask, std::vector<ArgDesc> &Args,
    std::function<void(detail::ArgDesc &Arg, int NextTrueIndex)> Func);
 
void ReverseRangeDimensionsForKernel(NDRDescT &NDR);
 
} // namespace detail
} // namespace _V1
} // namespace sycl