P4 to Git Change 1567349 by jatang@jatang_rocm_lc2 on 2018/06/12 16:09:03

SWDEV-126897 - Support OpenCL Device Enqueue on ROCm.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#23 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.hpp#10 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdefs.hpp#17 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#88 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#35 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.hpp#22 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#32 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsched.hpp#1 add
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocschedcl.cpp#1 add
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.cpp#33 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.hpp#13 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#56 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.hpp#16 edit


[ROCm/clr commit: 6321912698]

projects/clr/rocclr/runtime/device/rocm/rocblit.cpp

@@ -5,7 +5,8 @@
 #include "device/rocm/rocdevice.hpp"
 #include "device/rocm/rocblit.hpp"
 #include "device/rocm/rocmemory.hpp"
-#include "device/rocm/rocvirtual.hpp"
+#include "device/rocm/rockernel.hpp"
+#include "device/rocm/rocsched.hpp"
 #include "utils/debug.hpp"
 #include <algorithm>
 
@@ -2108,4 +2109,69 @@ address KernelBlitManager::captureArguments(const amd::Kernel* kernel) const {
 void KernelBlitManager::releaseArguments(address args) const {
 }
 
+bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
+                                     hsa_queue_t* schedulerQueue,
+                                     hsa_signal_t& schedulerSignal,
+                                     uint threads) {
+  size_t globalWorkOffset[1] = {0};
+  size_t globalWorkSize[1] = {threads};
+  size_t localWorkSize[1] = {1};
+
+  amd::NDRangeContainer ndrange(1, globalWorkOffset, globalWorkSize, localWorkSize);
+
+  device::Kernel* devKernel = const_cast<device::Kernel*>(kernels_[Scheduler]->getDeviceKernel(dev()));
+  Kernel& gpuKernel = static_cast<Kernel&>(*devKernel);
+
+  SchedulerParam* sp = reinterpret_cast<SchedulerParam*>(schedulerParam->getHostMem());
+  memset(sp, 0, sizeof(SchedulerParam));
+
+  Memory* schedulerMem = dev().getRocMemory(schedulerParam);
+  sp->kernarg_address = reinterpret_cast<uint64_t>(schedulerMem->getDeviceMemory());
+
+  sp->hidden_global_offset_x = 0;
+  sp->hidden_global_offset_y = 0;
+  sp->hidden_global_offset_z = 0;
+  sp->thread_counter = 0;
+  sp->child_queue = reinterpret_cast<uint64_t>(schedulerQueue);
+  sp->complete_signal = schedulerSignal;
+
+  hsa_signal_t  signal1;
+  signal1 = schedulerSignal;
+  hsa_signal_store_relaxed(signal1, 1);
+
+  sp->scheduler_aql.header = 0;
+  sp->scheduler_aql.setup = 1;
+  sp->scheduler_aql.workgroup_size_x = 1;
+  sp->scheduler_aql.workgroup_size_y = 1;
+  sp->scheduler_aql.workgroup_size_z = 1;
+  sp->scheduler_aql.grid_size_x = threads;
+  sp->scheduler_aql.grid_size_y = 1;
+  sp->scheduler_aql.grid_size_z = 1;
+  sp->scheduler_aql.kernel_object = gpuKernel.KernelCodeHandle();
+  sp->scheduler_aql.kernarg_address = (void*)sp->kernarg_address;
+  sp->scheduler_aql.private_segment_size = 0;
+  sp->scheduler_aql.group_segment_size = 0;
+  sp->vqueue_header = vqVM;
+
+  sp->parentAQL = sp->kernarg_address + sizeof(SchedulerParam);
+
+  cl_mem mem = as_cl<amd::Memory>(schedulerParam);
+  setArgument(kernels_[Scheduler], 0, sizeof(cl_mem), &mem);
+
+  address parameters = captureArguments(kernels_[Scheduler]);
+
+  bool result = false;
+  result = gpu().submitKernelInternal(ndrange, *kernels_[Scheduler], parameters, nullptr);
+  releaseArguments(parameters);
+  synchronize();
+
+  if (hsa_signal_wait_acquire(signal1, HSA_SIGNAL_CONDITION_LT, 1, (-1),
+                                HSA_WAIT_STATE_BLOCKED) != 0) {
+    LogWarning("Failed schedulerSignal wait");
+    return false;
+  }
+
+  return true;
+}
+
 }  // namespace pal

projects/clr/rocclr/runtime/device/rocm/rocblit.hpp

@@ -9,6 +9,7 @@
 #include "device/device.hpp"
 #include "device/blit.hpp"
 #include "device/rocm/rocdefs.hpp"
+#include "device/rocm/rocsched.hpp"
 
 /*! \addtogroup ROC Blit Implementation
  *  @{
@@ -232,6 +233,7 @@ class KernelBlitManager : public DmaBlitManager {
     BlitCopyBufferAligned,
     FillBuffer,
     FillImage,
+    Scheduler,
     BlitTotal
   };
 
@@ -366,6 +368,12 @@ class KernelBlitManager : public DmaBlitManager {
                          bool entire = false          //!< Entire buffer will be updated
                          ) const;
 
+  bool runScheduler(uint64_t vqVM,
+                    amd::Memory* schedulerParam,
+                    hsa_queue_t* schedulerQueue,
+                    hsa_signal_t& schedulerSignal,
+                    uint threads);
+
  private:
   static const size_t MaxXferBuffers = 2;
   static const uint TransferSplitSize = 1;
@@ -426,7 +434,7 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = {
     "copyImage",         "copyImage1DA",      "copyImageToBuffer",
     "copyBufferToImage", "copyBufferRect",    "copyBufferRectAligned",
     "copyBuffer",        "copyBufferAligned", "fillBuffer",
-    "fillImage",
+    "fillImage",         "scheduler",
 };
 
 inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index, size_t size, const void* value) const {

projects/clr/rocclr/runtime/device/rocm/rocdefs.hpp

@@ -7,6 +7,9 @@ namespace roc {
 //! Alignment restriciton for the pinned memory
 const static size_t PinnedMemoryAlignment = 4 * Ki;
 
+//! Specific defines for images for Dynamic Parallelism
+const static uint DeviceQueueMaskSize = 32;
+
 typedef uint HsaDeviceId;
 
 struct AMDDeviceInfo {

projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp

@@ -560,6 +560,8 @@ bool Device::init() {
   return true;
 }
 
+extern const char* SchedulerSourceCode;
+
 void Device::tearDown() {
   NullDevice::tearDown();
   hsa_shut_down();
@@ -611,7 +613,12 @@ bool Device::create() {
     return false;
   }
 
+  const char* scheduler = nullptr;
+
 #if defined(WITH_LIGHTNING_COMPILER)
+  std::string sch = SchedulerSourceCode;
+  scheduler = sch.c_str();
+
   //  create compilation object with cache support
   int gfxipMajor = deviceInfo_.gfxipVersion_ / 100;
   int gfxipMinor = deviceInfo_.gfxipVersion_ / 10 % 10;
@@ -647,7 +654,7 @@ bool Device::create() {
 
   blitProgram_ = new BlitProgram(context_);
   // Create blit programs
-  if (blitProgram_ == nullptr || !blitProgram_->create(this)) {
+  if (blitProgram_ == nullptr || !blitProgram_->create(this, scheduler)) {
     delete blitProgram_;
     blitProgram_ = nullptr;
     LogError("Couldn't create blit kernels!");
@@ -1025,7 +1032,7 @@ bool Device::populateOCLDeviceConstants() {
   ss <<  ")";
 
   strcpy(info_.driverVersion_, ss.str().c_str());
-  info_.version_ = "OpenCL " /*OPENCL_VERSION_STR*/"1.2" " ";
+  info_.version_ = "OpenCL " OPENCL_VERSION_STR " ";
 
   info_.builtInKernels_ = "";
   info_.linkerAvailable_ = true;
@@ -1196,6 +1203,13 @@ bool Device::populateOCLDeviceConstants() {
   info_.maxPipeActiveReservations_ = 16;
   info_.maxPipeArgs_ = 16;
 
+  info_.queueOnDeviceProperties_ =
+      CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE;
+  info_.queueOnDevicePreferredSize_ = 256 * Ki;
+  info_.queueOnDeviceMaxSize_ = 8 * Mi;
+  info_.maxOnDeviceQueues_ = 1;
+  info_.maxOnDeviceEvents_ = settings().numDeviceEvents_;
+
   return true;
 }
 

projects/clr/rocclr/runtime/device/rocm/rockernel.cpp

@@ -30,6 +30,8 @@ static inline ROC_ARG_TYPE GetKernelArgType(const KernelArgMD& lcArg) {
       return ROC_ARGTYPE_IMAGE;
     case ValueKind::Sampler:
       return ROC_ARGTYPE_SAMPLER;
+    case ValueKind::Queue:
+      return ROC_ARGTYPE_QUEUE;
     case ValueKind::HiddenGlobalOffsetX:
       return ROC_ARGTYPE_HIDDEN_GLOBAL_OFFSET_X;
     case ValueKind::HiddenGlobalOffsetY:
@@ -78,6 +80,8 @@ static inline ROC_ARG_TYPE GetKernelArgType(const aclArgData* argInfo) {
       return ROC_ARGTYPE_IMAGE;
     case ARG_TYPE_SAMPLER:
       return ROC_ARGTYPE_SAMPLER;
+    case ARG_TYPE_QUEUE:
+      return ROC_ARGTYPE_QUEUE;
     case ARG_TYPE_ERROR:
     default:
       return ROC_ARGTYPE_ERROR;
@@ -417,6 +421,8 @@ static inline clk_value_type_t GetOclType(const Kernel::Argument* arg) {
     }
   } else if (arg->type_ == ROC_ARGTYPE_SAMPLER) {
     return T_SAMPLER;
+  } else if (arg->type_ == ROC_ARGTYPE_QUEUE) {
+    return T_QUEUE;
   } else {
     return T_VOID;
   }
@@ -718,6 +724,43 @@ bool LightningKernel::init() {
     workGroupInfo_.compileVecTypeHint_ = kernelMD->mAttrs.mVecTypeHint.c_str();
   }
 
+  if (!kernelMD->mAttrs.mRuntimeHandle.empty()) {
+    hsa_agent_t             agent = program_->hsaDevice();
+    hsa_executable_symbol_t kernelSymbol;
+    hsa_status_t            status;
+    int                     variable_size;
+    uint64_t                variable_address;
+
+    // Only kernels that could be enqueued by another kernel has the RuntimeHandle metadata. The RuntimeHandle
+    // metadata is a string that represents a variable from which the library code can retrieve the kernel code
+    // object handle of such a kernel. The address of the variable and the kernel code object handle are known
+    // only after the hsa executable is loaded. The below code copies the kernel code object handle to the
+    // address of the variable.
+
+    status = hsa_executable_get_symbol_by_name(program_->hsaExecutable(), kernelMD->mAttrs.mRuntimeHandle.c_str(),
+                                               &agent, &kernelSymbol);
+    if (status != HSA_STATUS_SUCCESS) {
+      return false;
+    }
+
+    status = hsa_executable_symbol_get_info(kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
+                                            &variable_size);
+    if (status != HSA_STATUS_SUCCESS) {
+      return false;
+    }
+
+    status = hsa_executable_symbol_get_info(kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
+                                            &variable_address);
+    if (status != HSA_STATUS_SUCCESS) {
+      return false;
+    }
+
+    status = hsa_memory_copy(reinterpret_cast<void*>(variable_address), &kernelCodeHandle_, variable_size);
+    if (status != HSA_STATUS_SUCCESS) {
+      return false;
+    }
+  }
+
   uint32_t wavefront_size = 0;
   if (hsa_agent_get_info(program_->hsaDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size) !=
       HSA_STATUS_SUCCESS) {

projects/clr/rocclr/runtime/device/rocm/rockernel.hpp

@@ -22,6 +22,7 @@ enum ROC_ARG_TYPE {
   ROC_ARGTYPE_REFERENCE,
   ROC_ARGTYPE_IMAGE,
   ROC_ARGTYPE_SAMPLER,
+  ROC_ARGTYPE_QUEUE,
   ROC_ARGTYPE_HIDDEN_GLOBAL_OFFSET_X,
   ROC_ARGTYPE_HIDDEN_GLOBAL_OFFSET_Y,
   ROC_ARGTYPE_HIDDEN_GLOBAL_OFFSET_Z,
@@ -121,6 +122,12 @@ class Kernel : public device::Kernel {
   //! Return printf info array
   const std::vector<PrintfInfo>& printfInfo() const { return printf_; }
 
+  //! Returns TRUE if kernel uses dynamic parallelism
+  bool dynamicParallelism() const { return (flags_.dynamicParallelism_) ? true : false; }
+
+  //! set dynamic parallelism flag
+  void setDynamicParallelFlag(bool flag) { flags_.dynamicParallelism_ = flag; }
+
   //! Return TRUE if kernel is internal blit kernel
   bool isInternalKernel() const { return (flags_.internalKernel_) ? true : false; }
 
@@ -139,6 +146,7 @@ class Kernel : public device::Kernel {
       uint internalKernel_ : 1;  //!< Is a blit kernel?
       uint imageEnable_    : 1;  //!< Kernel uses images
       uint imageWrite_     : 1;  //!< Kernel writes images
+      uint dynamicParallelism_ : 1;  //!< Dynamic parallelism enabled
     };
     uint value_;
     Flags() : value_(0) {}

projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp

@@ -51,6 +51,8 @@ class Program : public device::Program {
   //! Returns the hsaBinary associated with the program
   hsa_agent_t hsaDevice() const { return dev().getBackendDevice(); }
 
+  hsa_executable_t hsaExecutable() const { return hsaExecutable_; }
+
   bool hasGlobalStores() const { return hasGlobalStores_; }
 
  protected:

projects/clr/rocclr/runtime/device/rocm/rocsched.hpp

@@ -0,0 +1,75 @@
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+#pragma once
+
+namespace roc {
+
+//! AmdAqlWrap slot state
+enum AqlWrapState {
+  AQL_WRAP_FREE = 0,
+  AQL_WRAP_RESERVED,
+  AQL_WRAP_READY,
+  AQL_WRAP_MARKER,
+  AQL_WRAP_BUSY,
+  AQL_WRAP_DONE
+};
+
+struct AmdVQueueHeader {
+  uint32_t aql_slot_num;     //!< [LRO/SRO] The total number of the AQL slots (multiple of 64).
+  uint32_t event_slot_num;   //!< [LRO] The number of kernel events in the events buffer
+  uint64_t event_slot_mask;  //!< [LRO] A pointer to the allocation bitmask array for the events
+  uint64_t event_slots;      //!< [LRO] Pointer to a buffer for the events.
+                             // Array of event_slot_num entries of AmdEvent
+  uint64_t aql_slot_mask;    //!< [LRO/SRO]A pointer to the allocation bitmask for aql_warp slots
+  uint32_t command_counter;  //!< [LRW] The global counter for the submitted commands into the queue
+  uint32_t wait_size;        //!< [LRO] The wait list size (in clk_event_t)
+  uint32_t arg_size;         //!< [LRO] The size of argument buffer (in bytes)
+  uint32_t mask_groups;      //!< Processed mask groups by one thread
+  uint64_t kernel_table;     //!< [LRO] Pointer to an array with all kernel objects (ulong for each entry)
+  uint32_t reserved[2];      //!< For the future usage
+};
+
+struct AmdAqlWrap {
+  uint32_t state;          //!< [LRW/SRW] The current state of the AQL wrapper:  FREE, RESERVED, READY,
+                           // MARKER, BUSY and DONE. The block could be returned back to a free state.
+  uint32_t enqueue_flags;  //!< [LWO/SRO] Contains the flags for the kernel execution start
+  uint32_t command_id;     //!< [LWO/SRO] The unique command ID
+  uint32_t child_counter;  //!< [LRW/SRW] Counter that determine the launches of child kernels.
+                           // It’s incremented on the
+                           // start and decremented on the finish. The parent kernel can be considered as
+                           // done when the value is 0 and the state is DONE
+  uint64_t completion;     //!< [LWO/SRO] CL event for the current execution (clk_event_t)
+  uint64_t parent_wrap;    //!< [LWO/SRO] Pointer to the parent AQL wrapper (AmdAqlWrap*)
+  uint64_t wait_list;      //!< [LRO/SRO] Pointer to an array of clk_event_t objects (64 bytes default)
+  uint32_t wait_num;       //!< [LWO/SRO] The number of cl_event_wait objects
+  uint32_t reserved[5];              //!< For the future usage
+  hsa_kernel_dispatch_packet_t aql;  //!< [LWO/SRO] AQL packet – 64 bytes AQL packet
+};
+
+struct AmdEvent {
+  uint32_t state;        //!< [LRO/SRW] Event state: START, END, COMPLETE
+  uint32_t counter;      //!< [LRW] Event retain/release counter. 0 means the event is free
+  uint64_t timer[3];     //!< [LRO/SWO] Timer values for profiling for each state
+  uint64_t captureInfo;  //!< [LRW/SRO] Profiling capture info for CLK_PROFILING_COMMAND_EXEC_TIME
+};
+
+struct SchedulerParam {
+  uint64_t        kernarg_address;
+  uint64_t        hidden_global_offset_x;
+  uint64_t        hidden_global_offset_y;
+  uint64_t        hidden_global_offset_z;
+  uint64_t        thread_counter;
+  uint64_t        child_queue;
+  hsa_kernel_dispatch_packet_t    scheduler_aql;
+  hsa_signal_t    complete_signal;
+  uint64_t        vqueue_header;
+  uint32_t        signal;                   //!< Signal to stop the child queue
+  uint32_t        eng_clk;                  //!< Engine clock in Mhz
+  uint32_t        releaseHostCP;            //!< Releases CP on the host queue
+  uint64_t        parentAQL;                //!< Host parent AmdAqlWrap packet
+  uint32_t        dedicatedQueue;           //!< Scheduler uses a dedicated queue
+  uint32_t        reserved[2];              //!< Processed mask groups by one thread
+};
+
+}

projects/clr/rocclr/runtime/device/rocm/rocschedcl.cpp

@@ -0,0 +1,17 @@
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+namespace roc {
+
+#define SCHEDULER_KERNEL(...) #__VA_ARGS__
+
+const char* SchedulerSourceCode = SCHEDULER_KERNEL(
+\n
+extern void __amd_scheduler(__global void*);
+\n
+__kernel void scheduler(__global void* params) {
+  __amd_scheduler(params);
+}
+\n);
+
+}  // namespace roc

projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp

@@ -74,6 +74,10 @@ Settings::Settings() {
   singleFpDenorm_ = false;
 
   apuSystem_ = false;
+
+  // Device enqueuing settings
+  numDeviceEvents_ = 1024;
+  numWaitEvents_ = 8;
 }
 
 bool Settings::create(bool fullProfile, int gfxipVersion) {

projects/clr/rocclr/runtime/device/rocm/rocsettings.hpp

@@ -52,6 +52,8 @@ class Settings : public device::Settings {
 
   uint kernargPoolSize_;
   uint signalPoolSize_;
+  uint numDeviceEvents_;      //!< The number of device events
+  uint numWaitEvents_;        //!< The number of wait events for device enqueue
 
   size_t xferBufSize_;        //!< Transfer buffer size for image copy optimization
   size_t stagedXferSize_;     //!< Staged buffer size

projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp

@@ -449,6 +449,13 @@ bool VirtualGPU::releaseGpuMemoryFence() {
 VirtualGPU::VirtualGPU(Device& device)
     : device::VirtualDevice(device),
       roc_device_(device),
+      virtualQueue_(nullptr),
+      deviceQueueSize_(0),
+      maskGroups_(0),
+      schedulerThreads_(0),
+      schedulerParam_(nullptr),
+      schedulerQueue_(nullptr),
+      schedulerSignal_({0}),
       index_(device.numOfVgpus_++)  // Virtual gpu unique index incrementing
 {
   gpu_device_ = device.getBackendDevice();
@@ -478,6 +485,22 @@ VirtualGPU::~VirtualGPU() {
     printfdbg_ = nullptr;
   }
 
+  if (0 != schedulerSignal_.handle) {
+    hsa_signal_destroy(schedulerSignal_);
+  }
+
+  if (nullptr != schedulerQueue_) {
+    hsa_queue_destroy(schedulerQueue_);
+  }
+
+  if (nullptr != schedulerParam_) {
+    schedulerParam_->release();
+  }
+
+  if (nullptr != virtualQueue_) {
+    virtualQueue_->release();
+  }
+
   --roc_device_.numOfVgpus_;  // Virtual gpu unique index decrementing
 }
 
@@ -1558,6 +1581,186 @@ static void fillSampleDescriptor(hsa_ext_sampler_descriptor_t& samplerDescriptor
   }
 }
 
+bool VirtualGPU::createSchedulerParam()
+{
+  if (nullptr != schedulerParam_) {
+    return true;
+  }
+
+  while(true) {
+    schedulerParam_ = new (dev().context()) amd::Buffer(dev().context(), CL_MEM_ALLOC_HOST_PTR, sizeof(SchedulerParam) + sizeof(AmdAqlWrap));
+
+    if ((nullptr != schedulerParam_) && !schedulerParam_->create(nullptr)) {
+      break;
+    }
+
+    // The queue is written by multiple threads of the scheduler kernel
+    if (HSA_STATUS_SUCCESS != hsa_queue_create(gpu_device(), 2048, HSA_QUEUE_TYPE_MULTI,
+        nullptr, nullptr, std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
+        &schedulerQueue_)) {
+      break;
+    }
+
+    hsa_signal_t  signal0 = {0};
+
+    if (HSA_STATUS_SUCCESS != hsa_signal_create(0, 0, nullptr, &signal0)) {
+      break;
+    }
+
+    schedulerSignal_ = signal0;
+
+    Memory* schedulerMem = dev().getRocMemory(schedulerParam_);
+
+    if (nullptr == schedulerMem) {
+      break;
+    }
+
+    schedulerParam_->setVirtualDevice(this);
+    return true;
+  }
+
+  if (0 != schedulerSignal_.handle) {
+    hsa_signal_destroy(schedulerSignal_);
+    schedulerSignal_.handle = 0;
+  }
+
+  if (nullptr != schedulerQueue_) {
+    hsa_queue_destroy(schedulerQueue_);
+    schedulerQueue_ = nullptr;
+  }
+
+  if (nullptr != schedulerParam_) {
+    schedulerParam_->release();
+    schedulerParam_ = nullptr;
+  }
+
+  return false;
+}
+
+uint64_t VirtualGPU::getVQVirtualAddress()
+{
+  Memory* vqMem = dev().getRocMemory(virtualQueue_);
+  return reinterpret_cast<uint64_t>(vqMem->getDeviceMemory());
+}
+
+bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
+{
+  uint MinDeviceQueueSize = 16 * 1024;
+  deviceQueueSize = std::max(deviceQueueSize, MinDeviceQueueSize);
+
+  maskGroups_ = deviceQueueSize / (512 * Ki);
+  maskGroups_ = (maskGroups_ == 0) ? 1 : maskGroups_;
+
+  // Align the queue size for the multiple dispatch scheduler.
+  // Each thread works with 32 entries * maskGroups
+  uint extra = deviceQueueSize % (sizeof(AmdAqlWrap) * DeviceQueueMaskSize * maskGroups_);
+  if (extra != 0) {
+    deviceQueueSize += (sizeof(AmdAqlWrap) * DeviceQueueMaskSize * maskGroups_) - extra;
+  }
+
+  if (deviceQueueSize_ == deviceQueueSize) {
+    return true;
+  } else {
+    if (0 != deviceQueueSize_) {
+      virtualQueue_->release();
+      virtualQueue_ = nullptr;
+      deviceQueueSize_ = 0;
+      schedulerThreads_ = 0;
+    }
+  }
+
+  uint numSlots = deviceQueueSize / sizeof(AmdAqlWrap);
+  uint allocSize = deviceQueueSize;
+
+  // Add the virtual queue header
+  allocSize += sizeof(AmdVQueueHeader);
+  allocSize = amd::alignUp(allocSize, sizeof(AmdAqlWrap));
+
+  uint argOffs = allocSize;
+
+  // Add the kernel arguments and wait events
+  uint singleArgSize = amd::alignUp(
+      dev().info().maxParameterSize_ + 64 + dev().settings().numWaitEvents_ * sizeof(uint64_t),
+      sizeof(AmdAqlWrap));
+  allocSize += singleArgSize * numSlots;
+
+  uint eventsOffs = allocSize;
+  // Add the device events
+  allocSize += dev().settings().numDeviceEvents_ * sizeof(AmdEvent);
+
+  uint eventMaskOffs = allocSize;
+  // Add mask array for events
+  allocSize += amd::alignUp(dev().settings().numDeviceEvents_, DeviceQueueMaskSize) / 8;
+
+  uint slotMaskOffs = allocSize;
+  // Add mask array for AmdAqlWrap slots
+  allocSize += amd::alignUp(numSlots, DeviceQueueMaskSize) / 8;
+
+  // CL_MEM_ALLOC_HOST_PTR/CL_MEM_READ_WRITE
+  virtualQueue_ = new (dev().context()) amd::Buffer(dev().context(), CL_MEM_READ_WRITE, allocSize);
+
+  if ((nullptr != virtualQueue_) && !virtualQueue_->create(nullptr)) {
+    virtualQueue_->release();
+    return false;
+  }
+
+  Memory* vqMem = dev().getRocMemory(virtualQueue_);
+
+  if (nullptr == vqMem) {
+    return false;
+  }
+
+  uint64_t vqVA = reinterpret_cast<uint64_t>(vqMem->getDeviceMemory());
+  uint64_t pattern = 0;
+  amd::Coord3D origin(0, 0, 0);
+  amd::Coord3D region(virtualQueue_->getSize(), 0, 0);
+
+  if (!dev().xferMgr().fillBuffer(*vqMem, &pattern, sizeof(pattern), origin, region)) {
+    return false;
+  }
+
+  AmdVQueueHeader header = {};
+  // Initialize the virtual queue header
+  header.aql_slot_num = numSlots;
+  header.event_slot_num = dev().settings().numDeviceEvents_;
+  header.event_slot_mask = vqVA + eventMaskOffs;
+  header.event_slots = vqVA + eventsOffs;
+  header.aql_slot_mask = vqVA + slotMaskOffs;
+  header.wait_size = dev().settings().numWaitEvents_;
+  header.arg_size = dev().info().maxParameterSize_ + 64;
+  header.mask_groups = maskGroups_;
+
+  amd::Coord3D origin_header(0);
+  amd::Coord3D region_header(sizeof(AmdVQueueHeader));
+
+  if (!dev().xferMgr().writeBuffer(&header, *vqMem, origin_header, region_header)) {
+    return false;
+  }
+
+  // Go over all slots and perform initialization
+  AmdAqlWrap slot = {};
+  size_t offset = sizeof(AmdVQueueHeader);
+  for (uint i = 0; i < numSlots; ++i) {
+    uint64_t argStart = vqVA + argOffs + i * singleArgSize;
+    amd::Coord3D origin_slot(offset);
+    amd::Coord3D region_slot(sizeof(AmdAqlWrap));
+
+    slot.aql.kernarg_address = reinterpret_cast<void*>(argStart);
+    slot.wait_list = argStart + dev().info().maxParameterSize_ + 64;
+
+    if (!dev().xferMgr().writeBuffer(&slot, *vqMem, origin_slot, region_slot)) {
+      return false;
+    }
+
+    offset += sizeof(AmdAqlWrap);
+  }
+
+  deviceQueueSize_ = deviceQueueSize;
+  schedulerThreads_ = numSlots / (DeviceQueueMaskSize * maskGroups_);
+
+  return true;
+}
+
 bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel,
                                       const_address parameters, void* eventHandle) {
   device::Kernel* devKernel = const_cast<device::Kernel*>(kernel.getDeviceKernel(dev()));
@@ -1667,8 +1870,45 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
           argPtr = addArg(argPtr, &bufferPtr, arg->size_, arg->alignment_);
           break;
         }
-        case ROC_ARGTYPE_HIDDEN_DEFAULT_QUEUE:
-        case ROC_ARGTYPE_HIDDEN_COMPLETION_ACTION:
+        case ROC_ARGTYPE_QUEUE: {
+          uint32_t index = signature.at(arg->index_).info_.arrayIndex_;
+          const amd::DeviceQueue* queue = reinterpret_cast<amd::DeviceQueue* const*>(parameters +
+            kernelParams.samplerObjOffset())[index];
+          if (queue == nullptr) {
+            return false;
+          }
+
+          if (!createVirtualQueue(queue->size()) || !createSchedulerParam()) {
+            return false;
+          }
+          gpuKernel.setDynamicParallelFlag(true);
+          uint64_t vqVA = getVQVirtualAddress();
+          argPtr = addArg(argPtr, &vqVA, arg->size_, arg->alignment_);
+          break;
+        }
+        case ROC_ARGTYPE_HIDDEN_DEFAULT_QUEUE: {
+
+          amd::DeviceQueue* defQueue = kernel.program().context().defDeviceQueue(dev());
+
+          if (!createVirtualQueue(defQueue->size()) || !createSchedulerParam()) {
+            return false;
+          }
+          gpuKernel.setDynamicParallelFlag(true);
+          uint64_t vqVA = getVQVirtualAddress();
+          argPtr = addArg(argPtr, &vqVA, arg->size_, arg->alignment_);
+          break;
+        }
+        case ROC_ARGTYPE_HIDDEN_COMPLETION_ACTION: {
+
+          Memory* schedulerMem = dev().getRocMemory(schedulerParam_);
+          AmdAqlWrap* wrap = reinterpret_cast<AmdAqlWrap*>(reinterpret_cast<uint64_t>(schedulerParam_->getHostMem()) + sizeof(SchedulerParam));
+          memset(wrap, 0, sizeof(AmdAqlWrap));
+          wrap->state = AQL_WRAP_DONE;
+
+          uint64_t spVA = reinterpret_cast<uint64_t>(schedulerMem->getDeviceMemory()) + sizeof(SchedulerParam);
+          argPtr = addArg(argPtr, &spVA, arg->size_, arg->alignment_);
+          break;
+        }
         case ROC_ARGTYPE_HIDDEN_NONE: {
           void* zero = 0;
           assert(arg->size_ <= sizeof(zero) && "check the sizes");
@@ -1846,6 +2086,13 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
     LogError("\nCould not print data from the printf buffer!");
     return false;
   }
+
+  if (gpuKernel.dynamicParallelism()) {
+    dispatchBarrierPacket(&barrier_packet_);
+    static_cast<KernelBlitManager&>(blitMgr()).runScheduler(
+      getVQVirtualAddress(), schedulerParam_, schedulerQueue_, schedulerSignal_, schedulerThreads_);
+  }
+
   return true;
 }
 /**

projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp

@@ -11,6 +11,7 @@
 #include "hsa_ext_amd.h"
 #include "rocprintf.hpp"
 #include "hsa_ven_amd_aqlprofile.h"
+#include "rocsched.hpp"
 
 namespace roc {
 class Device;
@@ -253,6 +254,13 @@ class VirtualGPU : public device::VirtualDevice {
   void* allocKernArg(size_t size, size_t alignment);
   void resetKernArgPool() { kernarg_pool_cur_offset_ = 0; }
 
+  uint64_t getVQVirtualAddress();
+
+  bool createSchedulerParam();
+
+  //! Returns TRUE if virtual queue was successfully allocatted
+  bool createVirtualQueue(uint deviceQueueSize);
+
   //! Updates AQL header for the upcomming dispatch
   void setAqlHeader(uint16_t header) { aqlHeader_ = header; }
 
@@ -281,6 +289,15 @@ class VirtualGPU : public device::VirtualDevice {
   MemoryDependency memoryDependency_;  //!< Memory dependency class
   uint16_t aqlHeader_;                 //!< AQL header for dispatch
 
+  amd::Memory* virtualQueue_;     //!< Virtual device queue
+  uint deviceQueueSize_;          //!< Device queue size
+  uint maskGroups_;               //!< The number of mask groups processed in the scheduler by one thread
+  uint schedulerThreads_;         //!< The number of scheduler threads
+
+  amd::Memory* schedulerParam_;
+  hsa_queue_t* schedulerQueue_;
+  hsa_signal_t schedulerSignal_;
+
   char* kernarg_pool_base_;
   size_t kernarg_pool_size_;
   uint kernarg_pool_cur_offset_;