SWDEV-422207 - Capture AQL Packets for graph Kernel nodes during graph Inst. And enqueue AQL packet during launch

Change-Id: I1e5f7f9e2a70bd500d190193cb6ba0867f5a63e7

hipamd/src/hip_graph.cpp

@@ -1177,6 +1177,10 @@ hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
   hip::GraphExec* ge;
   hipError_t status = ihipGraphInstantiate(&ge, reinterpret_cast<hip::Graph*>(graph));
   *pGraphExec = reinterpret_cast<hipGraphExec_t>(ge);
+  if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) {
+    // For graph nodes capture AQL packets to dispatch them directly during graph launch.
+    status = ge->CaptureAQLPackets();
+  }
   HIP_RETURN(status);
 }
 

hipamd/src/hip_graph_internal.cpp

@@ -480,7 +480,7 @@ hipError_t GraphExec::CreateStreams(uint32_t num_streams) {
 }
 
 hipError_t GraphExec::Init() {
-  hipError_t status;
+  hipError_t status = hipSuccess;
   size_t min_num_streams = 1;
 
   for (auto& node : topoOrder_) {
@@ -493,11 +493,63 @@ hipError_t GraphExec::Init() {
   return status;
 }
 
+hipError_t GraphExec::CaptureAQLPackets() {
+  hipError_t status = hipSuccess;
+  size_t KernArgSizeForGraph = 0;
+  bool GraphHasOnlyKerns = true;
+  // GPU packet capture is enabled for kernel nodes. Calculate the kernel arg size required for all
+  // graph kernel nodes to allocate
+  for (const auto& list : parallelLists_) {
+    hip::Stream* stream = GetAvailableStreams();
+    for (auto& node : list) {
+      node->SetStream(stream, this);
+      if (node->GetType() == hipGraphNodeTypeKernel) {
+        KernArgSizeForGraph += reinterpret_cast<hip::GraphKernelNode*>(node)->GetKerArgSize();
+      } else {
+        GraphHasOnlyKerns = false;
+      }
+    }
+  }
+
+  auto device = g_devices[ihipGetDevice()]->devices()[0];
+  const auto& info = device->info();
+  // Enable allocating kerns on device memory if graph as only kernels. memcpy nodes require hdp
+  // flush. ToDo: Work on enabling device kern args later for all type of nodes for large bar
+  if (GraphHasOnlyKerns == true && info.largeBar_) {
+    kernarg_pool_graph_ = reinterpret_cast<address>(device->deviceLocalAlloc(KernArgSizeForGraph));
+    device_kernarg_pool_ = true;
+  } else {
+    kernarg_pool_graph_ = reinterpret_cast<address>(
+        device->hostAlloc(KernArgSizeForGraph, 0, amd::Device::MemorySegment::kKernArg));
+  }
+
+  if (kernarg_pool_graph_ == nullptr) {
+    return hipErrorMemoryAllocation;
+  }
+  kernarg_pool_size_graph_ = KernArgSizeForGraph;
+
+  for (auto& node : topoOrder_) {
+    if (node->GetType() == hipGraphNodeTypeKernel) {
+      auto kernelnode = reinterpret_cast<hip::GraphKernelNode*>(node);
+      status = node->CreateCommand(node->GetQueue());
+      // From the kernel pool allocate the kern arg size required for the current kernel node.
+      address kernArgOffset = allocKernArg(kernelnode->GetKernargSegmentByteSize(),
+                                           kernelnode->GetKernargSegmentAlignment());
+      if (kernArgOffset == nullptr) {
+        return hipErrorMemoryAllocation;
+      }
+      // Enable GPU packet capture for the kernel node.
+      kernelnode->EnableCapturing(kernArgOffset);
+    }
+  }
+  return status;
+}
+
 hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
                         std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
                         std::vector<Node>& topoOrder, Graph* clonedGraph,
                         amd::Command*& graphStart, amd::Command*& graphEnd, hip::Stream* stream) {
-  hipError_t status;
+  hipError_t status = hipSuccess;
   for (auto& node : topoOrder) {
     // TODO: clone commands from next launch
     status = node->CreateCommand(node->GetQueue());
@@ -578,7 +630,7 @@ void UpdateStream(std::vector<std::vector<Node>>& parallelLists, hip::Stream* st
 }
 
 hipError_t GraphExec::Run(hipStream_t stream) {
-  hipError_t status;
+  hipError_t status = hipSuccess;
 
   if (hip::getStream(stream) == nullptr) {
     return hipErrorInvalidResourceHandle;
@@ -603,19 +655,30 @@ hipError_t GraphExec::Run(hipStream_t stream) {
     repeatLaunch_ = true;
   }
   if (parallelLists_.size() == 1) {
+    if (device_kernarg_pool_) {
+      // If kernelArgs are in device memory flush the HDP.
+      amd::Command* startCommand = nullptr;
+      startCommand = new amd::Marker(*hip_stream, false);
+      startCommand->enqueue();
+      startCommand->release();
+    }
     for (int i = 0; i < topoOrder_.size(); i++) {
-      topoOrder_[i]->SetStream(hip_stream, this);
-      status = topoOrder_[i]->CreateCommand(topoOrder_[i]->GetQueue());
-      if (DEBUG_CLR_GRAPH_ENABLE_BUFFERING) {
-        // Enable buffering for graph with single branch
-        // Peep through the next node. If current and next node are kernel then enable AQL
-        // buffering
-        if (((i + 1) != topoOrder_.size()) && topoOrder_[i]->GetType() == hipGraphNodeTypeKernel &&
-            topoOrder_[i + 1]->GetType() == hipGraphNodeTypeKernel) {
-          topoOrder_[i]->EnableBuffering();
-        }
+      if (DEBUG_CLR_GRAPH_PACKET_CAPTURE && topoOrder_[i]->GetType() == hipGraphNodeTypeKernel) {
+        hip_stream->vdev()->dispatchAqlPacket(topoOrder_[i]->GetAqlPacket());
+      } else {
+        topoOrder_[i]->SetStream(hip_stream, this);
+        status = topoOrder_[i]->CreateCommand(topoOrder_[i]->GetQueue());
+        topoOrder_[i]->EnqueueCommands(stream);
       }
-      topoOrder_[i]->EnqueueCommands(stream);
+    }
+    if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) {
+      amd::Command* endCommand = nullptr;
+      endCommand = new amd::Marker(*hip_stream, false);
+      // Since the end command is for graph completion tracking,
+      // it may not need release scopes
+      endCommand->setEventScope(amd::Device::kCacheStateIgnore);
+      endCommand->enqueue();
+      endCommand->release();
     }
   } else {
     UpdateStream(parallelLists_, hip_stream, this);

hipamd/src/hip_graph_internal.hpp

@@ -182,6 +182,7 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
   static std::unordered_set<GraphNode*> nodeSet_;
   static amd::Monitor nodeSetLock_;
   unsigned int isEnabled_;
+  uint8_t gpuPacket_[64];  //!< GPU Packet to enqueue during graph launch
 
  public:
   GraphNode(hipGraphNodeType type, std::string style = "", std::string shape = "",
@@ -229,7 +230,8 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
     }
     return true;
   }
-
+  // Return gpu packet address to update with actual packet under capture.
+  uint8_t* GetAqlPacket() { return gpuPacket_; }
   hip::Stream* GetQueue() { return stream_; }
 
   virtual void SetStream(hip::Stream* stream, GraphExec* ptr = nullptr) {
@@ -336,11 +338,6 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
       command->release();
     }
   }
-  virtual void EnableBuffering() {
-    for (auto& command : commands_) {
-      command->setBufferingState(true);
-    }
-  }
   Graph* GetParentGraph() { return parentGraph_; }
   virtual Graph* GetChildGraph() { return nullptr; }
   void SetParentGraph(Graph* graph) { parentGraph_ = graph; }
@@ -567,7 +564,11 @@ struct GraphExec {
   static amd::Monitor graphExecSetLock_;
   uint64_t flags_ = 0;
   bool repeatLaunch_ = false;
-
+  // Graph Kernel arg vars
+  bool device_kernarg_pool_ = false;
+  address kernarg_pool_graph_ = nullptr;
+  uint32_t kernarg_pool_size_graph_ = 0;
+  uint32_t kernarg_pool_cur_graph_offset_ = 0;
  public:
   GraphExec(std::vector<Node>& topoOrder, std::vector<std::vector<Node>>& lists,
             std::unordered_map<Node, std::vector<Node>>& nodeWaitLists, struct Graph*& clonedGraph,
@@ -592,6 +593,11 @@ struct GraphExec {
         hip::Stream::Destroy(stream);
       }
     }
+    // Release the kernel arg memory.
+    auto device = g_devices[ihipGetDevice()]->devices()[0];
+    if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) {
+      device->hostFree(kernarg_pool_graph_, kernarg_pool_size_graph_);
+    }
     amd::ScopedLock lock(graphExecSetLock_);
     graphExecSet_.erase(this);
     delete clonedGraph_;
@@ -606,7 +612,16 @@ struct GraphExec {
     }
     return clonedNode;
   }
-
+  address allocKernArg(size_t size, size_t alignment) {
+    assert(alignment != 0);
+    address result = nullptr;
+    result = amd::alignUp(kernarg_pool_graph_ + kernarg_pool_cur_graph_offset_, alignment);
+    const size_t pool_new_usage = (result + size) - kernarg_pool_graph_;
+    if (pool_new_usage <= kernarg_pool_size_graph_) {
+      kernarg_pool_cur_graph_offset_ = pool_new_usage;
+    }
+    return result;
+  }
   // check executable graphs validity
   static bool isGraphExecValid(GraphExec* pGraphExec);
 
@@ -617,6 +632,8 @@ struct GraphExec {
   hipError_t Init();
   hipError_t CreateStreams(uint32_t num_streams);
   hipError_t Run(hipStream_t stream);
+  // Capture GPU Packets from graph commands
+  hipError_t CaptureAQLPackets();
 };
 
 struct ChildGraphNode : public GraphNode {
@@ -742,31 +759,48 @@ struct ChildGraphNode : public GraphNode {
 };
 
 class GraphKernelNode : public GraphNode {
-  hipKernelNodeParams kernelParams_;
-  unsigned int numParams_;
-  hipKernelNodeAttrValue kernelAttr_;
-  unsigned int kernelAttrInUse_;
-  ihipExtKernelEvents kernelEvents_;
+  hipKernelNodeParams kernelParams_;   //!< Kernel node parameters
+  unsigned int numParams_;             //!< No. of kernel params as part of signature
+  hipKernelNodeAttrValue kernelAttr_;  //!< Kernel node attributes
+  unsigned int kernelAttrInUse_;       //!< Kernel attributes in use
+  ihipExtKernelEvents kernelEvents_;   //!< Events for Ext launch kernel
+  size_t alignedKernArgSize_;          //!< Aligned size required for kernel args
+  size_t kernargSegmentByteSize_;      //!< Kernel arg segment byte size
+  size_t kernargSegmentAlignment_;     //!< Kernel arg segment alignment
 
  public:
-    void PrintAttributes(std::ostream& out, hipGraphDebugDotFlags flag) {
-      out << "[";
-      out << "style";
-      out << "=\"";
-      out << style_;
-      (flag == hipGraphDebugDotFlagsKernelNodeParams ||
-       flag == hipGraphDebugDotFlagsKernelNodeAttributes) ?
-       out << "\n" : out << "\"";
-      out << "shape";
-      out << "=\"";
-      out << GetShape(flag);
-      out << "\"";
-      out << "label";
-      out << "=\"";
-      out << GetLabel(flag);
-      out << "\"";
-      out << "];";
+  size_t GetKerArgSize() const { return alignedKernArgSize_; }
+  size_t GetKernargSegmentByteSize() const { return kernargSegmentByteSize_; }
+  size_t GetKernargSegmentAlignment() const { return kernargSegmentAlignment_; }
+  void PrintAttributes(std::ostream& out, hipGraphDebugDotFlags flag) {
+    out << "[";
+    out << "style";
+    out << "=\"";
+    out << style_;
+    (flag == hipGraphDebugDotFlagsKernelNodeParams ||
+      flag == hipGraphDebugDotFlagsKernelNodeAttributes) ?
+      out << "\n" : out << "\"";
+    out << "shape";
+    out << "=\"";
+    out << GetShape(flag);
+    out << "\"";
+    out << "label";
+    out << "=\"";
+    out << GetLabel(flag);
+    out << "\"";
+    out << "];";
+    }
+
+    void EnableCapturing(address kernArgOffset) {
+      for (auto& command : commands_) {
+        reinterpret_cast<amd::NDRangeKernelCommand*>(command)->setCapturingState(
+            true, GetAqlPacket(), kernArgOffset);
+        // Enqueue command to capture GPU Packet. Packet is not sent to hardware queue.
+
+        command->submit(*(command->queue())->vdev());
+        command->release();
       }
+    }
 
   std::string GetLabel(hipGraphDebugDotFlags flag) {
     hipFunction_t func = getFunc(kernelParams_, ihipGetDevice());
@@ -842,6 +876,14 @@ class GraphKernelNode : public GraphNode {
     }
     hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func);
     amd::Kernel* kernel = function->kernel();
+    if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) {
+      auto device = g_devices[ihipGetDevice()]->devices()[0];
+      device::Kernel* devKernel = const_cast<device::Kernel*>(kernel->getDeviceKernel(*device));
+      kernargSegmentByteSize_ = devKernel->KernargSegmentByteSize();
+      kernargSegmentAlignment_ = devKernel->KernargSegmentAlignment();
+      alignedKernArgSize_ =
+          amd::alignUp(devKernel->KernargSegmentByteSize(), devKernel->KernargSegmentAlignment());
+    }
     const amd::KernelSignature& signature = kernel->signature();
     numParams_ = signature.numParameters();
 

rocclr/device/device.hpp

@@ -1279,6 +1279,7 @@ class VirtualDevice : public amd::HeapObject {
 
   //! Returns fence state of the VirtualGPU
   virtual bool isFenceDirty() const = 0;
+  virtual bool dispatchAqlPacket(uint8_t* aqlpacket) = 0;
 
   //! Resets fence state of the VirtualGPU
   virtual void resetFenceDirty() = 0;
@@ -1733,6 +1734,12 @@ class Device : public RuntimeObject {
     return NULL;
   }
 
+  virtual void* deviceLocalAlloc(size_t size, bool atomics = false,
+                                 bool pseudo_fine_grain = false) const {
+    ShouldNotCallThis();
+    return NULL;
+  }
+
   virtual bool deviceAllowAccess(void* dst) const {
     ShouldNotCallThis();
     return true;

rocclr/device/pal/palvirtual.hpp

@@ -344,6 +344,8 @@ class VirtualGPU : public device::VirtualDevice {
 
   bool isFenceDirty() const { return false; }
 
+  inline bool dispatchAqlPacket(uint8_t* aqlpacket) { return false; }
+
   void resetFenceDirty() {}
 
   //! Returns GPU device object associated with this kernel

rocclr/device/rocm/rocvirtual.cpp

@@ -815,74 +815,6 @@ static inline void packet_store_release(uint32_t* packet, uint16_t header, uint1
   __atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
 }
 
-bool VirtualGPU::dispatchAqlBuffer() {
-  size_t size = aqlBuffer_.size();
-  if (size > 0) {
-    const uint32_t queueSize = gpu_queue_->size;
-    const uint32_t queueMask = queueSize - 1;
-    const uint32_t sw_queue_size = queueMask;
-
-    // Check for queue full and wait if needed.
-    uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size);
-    uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_);
-
-    while (index - hsa_queue_load_read_index_scacquire(gpu_queue_) >= sw_queue_size - size) {
-      amd::Os::yield();
-    }
-
-    if (timestamp_ != nullptr) {
-      for (uint i = 0; i < size; i++) {
-        // Get active signal for current dispatch if profiling is necessary
-        aqlBuffer_[i].completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
-      }
-    }
-    for (uint i = 0; i < size; i++) {
-      ClPrint(
-          amd::LOG_DEBUG, amd::LOG_AQL,
-          "HWq=0x%zx, Dispatch AQL Buffer Header = "
-          "0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
-          "setup=%d, grid=[%zu, %zu, %zu], workgroup=[%zu, %zu, %zu], private_seg_size=%zu, "
-          "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx",
-          gpu_queue_->base_address, aqlBuffer_[i].header,
-          extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_TYPE,
-                         HSA_PACKET_HEADER_WIDTH_TYPE),
-          extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_BARRIER,
-                         HSA_PACKET_HEADER_WIDTH_BARRIER),
-          extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
-                         HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
-          extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
-                         HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
-          aqlBuffer_[i].setup, (aqlBuffer_[i]).grid_size_x, (aqlBuffer_[i]).grid_size_y,
-          (aqlBuffer_[i]).grid_size_z, (aqlBuffer_[i]).workgroup_size_x,
-          (aqlBuffer_[i]).workgroup_size_y, (aqlBuffer_[i]).workgroup_size_z,
-          (aqlBuffer_[i]).private_segment_size, (aqlBuffer_[i]).group_segment_size,
-          (aqlBuffer_[i]).kernel_object, (aqlBuffer_[i]).kernarg_address,
-          (aqlBuffer_[i]).completion_signal);
-    }
-    uint16_t firstPacketHeader = aqlBuffer_.front().header;
-    aqlBuffer_.front().header = kInvalidAql;
-    hsa_kernel_dispatch_packet_t* aql_loc =
-        &((hsa_kernel_dispatch_packet_t*)(gpu_queue_->base_address))[index & queueMask];
-    size_t size_before_wrap = queueSize - (index % queueSize);
-    if (size_before_wrap < size) {
-      amd::Os::fastMemcpy(aql_loc, &aqlBuffer_[0], sizeof(hsa_kernel_dispatch_packet_t) * size_before_wrap);
-      hsa_kernel_dispatch_packet_t* aql_loc_0 =
-          &((hsa_kernel_dispatch_packet_t*)(gpu_queue_->base_address))[0];
-      amd::Os::fastMemcpy(aql_loc_0, &aqlBuffer_[size_before_wrap],
-             sizeof(hsa_kernel_dispatch_packet_t) * (size - size_before_wrap));
-    } else {
-      amd::Os::fastMemcpy(aql_loc, &aqlBuffer_[0], sizeof(hsa_kernel_dispatch_packet_t) * size);
-    }
-
-    packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), firstPacketHeader,
-                         aqlBuffer_.front().setup);
-    hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index + size - 1);
-
-    aqlBuffer_.clear();
-  }
-  return true;
-}
-
 // ================================================================================================
 template <typename AqlPacket>
 bool VirtualGPU::dispatchGenericAqlPacket(
@@ -990,36 +922,23 @@ void VirtualGPU::dispatchBlockingWait() {
     }
   }
 }
-
 // ================================================================================================
 bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header,
-                                   uint16_t rest, bool blocking, bool buffering) {
-  static size_t initialAQLBufferSize = 1;
+                                   uint16_t rest, bool blocking, bool capturing,
+                                   const uint8_t* aqlPacket) {
   dispatchBlockingWait();
-  if (buffering == true) {
-    aqlBuffer_.push_back(*packet);
-    aqlBuffer_.back().header = header;
-    aqlBuffer_.back().setup = rest;
-    if (!(aqlBuffer_.size() >=
-          initialAQLBufferSize)) {  // Buffer maximum of AQL Buffer size packets once it
-                                    // exceeds send them for dispatch
-      return true;
+  if (capturing == true) {
+    packet->header = header;
+    packet->setup = rest;
+    if (timestamp_ != nullptr) {
+      // Get active signal for current dispatch if profiling is necessary
+      packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
     }
-  } else if (!aqlBuffer_.empty()) {  // If buffering is disabled and AQLBuffer is not empty then
-                                     // make sure current packet is added to the buffer for dispatch
-    aqlBuffer_.push_back(*packet);
-    aqlBuffer_.back().header = header;
-    aqlBuffer_.back().setup = rest;
-  }
-  if (aqlBuffer_.empty()) {
-    return dispatchGenericAqlPacket(packet, header, rest, blocking);
+    amd::Os::fastMemcpy(const_cast<uint8_t*>(aqlPacket), packet,
+                        sizeof(hsa_kernel_dispatch_packet_t));
+    return true;
   } else {
-    ClPrint(amd::LOG_DEBUG, amd::LOG_CODE, "Dispath AQL Buffer size:%ld", aqlBuffer_.size());
-    // Increment buffer size ^2 until DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE
-    if (initialAQLBufferSize < DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE) {
-      initialAQLBufferSize = initialAQLBufferSize << 1;
-    }
-    return dispatchAqlBuffer();
+    return dispatchGenericAqlPacket(packet, header, rest, blocking);
   }
 }
 // ================================================================================================
@@ -1129,6 +1048,18 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
   barrier_packet_.dep_signal[4] = hsa_signal_t{};
 }
 
+inline bool VirtualGPU::dispatchAqlPacket(uint8_t* aqlpacket) {
+  auto packet = reinterpret_cast<hsa_kernel_dispatch_packet_t*>(aqlpacket);
+  // If rocprof tracing is enabled, store the correlation ID in the dispatch packet.
+  // The profiler can retrieve this correlation ID to attribute waves to specific dispatch
+  // locations.
+  if (activity_prof::IsEnabled(OP_ID_DISPATCH)) {
+    packet->reserved2 = activity_prof::correlation_id;
+  }
+  dispatchGenericAqlPacket(packet, packet->header, packet->setup, false);
+  return true;
+}
+
 // ================================================================================================
 void VirtualGPU::dispatchBarrierValuePacket(uint16_t packetHeader, bool resolveDepSignal,
                                             hsa_signal_t signal, hsa_signal_value_t value,
@@ -1449,15 +1380,9 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) {
     kernarg_pool_cur_offset_ = pool_new_usage;
     return result;
   } else {
-    //! We run out of the arguments space!
     //! That means the app didn't call clFlush/clFinish for very long time.
     // Reset the signal for the barrier packet
     hsa_signal_silent_store_relaxed(kernarg_pool_signal_[active_chunk_], kInitSignalValueOne);
-    // dispatch any buffered AQL packets
-    bool status = dispatchAqlBuffer();
-    if (!status) {
-      LogError("dispatch Aql Buffer failed!");
-    }
     // Dispatch a barrier packet into the queue
     dispatchBarrierPacket(kBarrierPacketHeader, true, kernarg_pool_signal_[active_chunk_]);
     // Get the next chunk
@@ -3196,8 +3121,12 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
     // Find all parameters for the current kernel
     if (!kernel.parameters().deviceKernelArgs() || gpuKernel.isInternalKernel()) {
       // Allocate buffer to hold kernel arguments
-      argBuffer = reinterpret_cast<address>(allocKernArg(gpuKernel.KernargSegmentByteSize(),
+      if(vcmd != nullptr && vcmd->getCapturingState()) {
+        argBuffer = vcmd->getKernArgOffset();
+      } else {
+        argBuffer = reinterpret_cast<address>(allocKernArg(gpuKernel.KernargSegmentByteSize(),
                                             gpuKernel.KernargSegmentAlignment()));
+      }
       // Load all kernel arguments
       nontemporalMemcpy(argBuffer, parameters,
                         std::min(gpuKernel.KernargSegmentByteSize(),
@@ -3272,13 +3201,20 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
                            (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
       aql_packet->setup = sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
     }
-
-    // Dispatch the packet
-    if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder,
-                           (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
-                           GPU_FLUSH_ON_EXECUTION,
-                           (vcmd != nullptr) ? vcmd->getBufferingState() : false)) {
-      return false;
+    if (vcmd == nullptr) {
+      // Dispatch the packet
+      if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder,
+                             (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
+                             GPU_FLUSH_ON_EXECUTION)) {
+        return false;
+      }
+    } else {
+      if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder,
+                             (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
+                             GPU_FLUSH_ON_EXECUTION, vcmd->getCapturingState(),
+                             vcmd->getAqlPacket())) {
+        return false;
+      }
     }
   }
 

rocclr/device/rocm/rocvirtual.hpp

@@ -424,9 +424,10 @@ class VirtualGPU : public device::VirtualDevice {
   //! Dispatches a barrier with blocking HSA signals
   void dispatchBlockingWait();
 
-  bool dispatchAqlBuffer();
-  bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header,
-                         uint16_t rest, bool blocking = true, bool buffering = false);
+  inline bool dispatchAqlPacket(uint8_t* aqlpacket);
+  bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest,
+                         bool blocking = true, bool capturing = false,
+                         const uint8_t* aqlPacket = nullptr);
   bool dispatchAqlPacket(hsa_barrier_and_packet_t* packet, uint16_t header,
                         uint16_t rest, bool blocking = true);
   template <typename AqlPacket> bool dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header,
@@ -566,6 +567,5 @@ class VirtualGPU : public device::VirtualDevice {
   int fence_state_;                     //!< Fence scope
                                         //!< kUnknown/kFlushedToDevice/kFlushedToSystem
   bool fence_dirty_;                    //!< Fence modified flag
-  std::vector<hsa_kernel_dispatch_packet_t> aqlBuffer_;  //!< AQL packet buffer for graphs
 };
 }

rocclr/platform/command.cpp

@@ -317,7 +317,6 @@ Command::Command(HostQueue& queue, cl_command_type type, const EventWaitList& ev
       type_(type),
       data_(nullptr),
       waitingEvent_(waitingEvent),
-      buffering_(false),
       eventWaitList_(eventWaitList),
       commandWaitBits_(commandWaitBits) {
   // Retain the commands from the event wait list.
@@ -354,7 +353,7 @@ void Command::enqueue() {
 
     // Notify all commands about the waiter. Barrier will be sent in order to obtain
     // HSA signal for a wait on the current queue
-    for (const auto &event: eventWaitList()) {
+    for (const auto& event : eventWaitList()) {
       event->notifyCmdQueue(!kCpuWait);
     }
 

rocclr/platform/command.hpp

@@ -251,13 +251,12 @@ union CopyMetadata {
  */
 class Command : public Event {
  private:
-  HostQueue* queue_;              //!< The command queue this command is enqueue into
-  Command* next_;                 //!< Next GPU command in the queue list
-  Command* batch_head_ = nullptr; //!< The head of the batch commands
-  cl_command_type     type_;      //!< This command's OpenCL type.
+  HostQueue* queue_;               //!< The command queue this command is enqueue into
+  Command* next_;                  //!< Next GPU command in the queue list
+  Command* batch_head_ = nullptr;  //!< The head of the batch commands
+  cl_command_type type_;           //!< This command's OpenCL type.
   void* data_;
-  const Event* waitingEvent_;     //!< Waiting event associated with the marker
-  bool buffering_;                //!< Flag to enable/disable AQL buffering
+  const Event* waitingEvent_;  //!< Waiting event associated with the marker
 
  protected:
   bool cpu_wait_ = false;         //!< If true, then the command was issued for CPU/GPU sync
@@ -281,7 +280,6 @@ class Command : public Event {
         type_(type),
         data_(nullptr),
         waitingEvent_(nullptr),
-        buffering_(false),
         eventWaitList_(nullWaitList),
         commandWaitBits_(0) {}
 
@@ -296,12 +294,6 @@ class Command : public Event {
   }
 
  public:
-  //! Returns AQL buffer state
-  bool getBufferingState() const { return buffering_; }
-
-  //! Sets AQL buffer state
-  void setBufferingState(bool state) { buffering_ = state; }
-
   //! Return the queue this command is enqueued into.
   HostQueue* queue() const { return queue_; }
 
@@ -1083,6 +1075,10 @@ class NDRangeKernelCommand : public Command {
   uint32_t firstDevice_;    //!< Device index of the first device in the gridc
   uint32_t numWorkgroups_;  //!< Total number of workgroups in the current launch
 
+  bool capturing_ = false;           //!< Flag to enable/disable graph gpu packet capture
+  uint8_t* gpuPacket_ = nullptr;     //!< GPU packet to capture, when graph capturing is enabled
+  address kernArgOffset_ = nullptr;  //!< KernelArg buffer to used when graph capturing is enabled
+
  public:
   enum {
     CooperativeGroups = 0x01,
@@ -1090,6 +1086,22 @@ class NDRangeKernelCommand : public Command {
     AnyOrderLaunch = 0x04,
   };
 
+  //! Returns AQL buffer state
+  bool getCapturingState() const { return capturing_; }
+
+  //! Sets AQL capture state, aql packet to capture and where to copy kernArgs
+  void setCapturingState(bool state, uint8_t* packet, address kernArgOffset) {
+    capturing_ = state;
+    gpuPacket_ = packet;
+    kernArgOffset_ = kernArgOffset;
+  }
+
+  //! returns the graph executable object command belongs to.
+  const uint8_t* getAqlPacket() const { return gpuPacket_; }
+
+  //! returns the graph executable object command belongs to.
+  const address getKernArgOffset() const { return kernArgOffset_; }
+
   //! Construct an ExecuteKernel command
   NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList, Kernel& kernel,
                        const NDRangeContainer& sizes, uint32_t sharedMemBytes = 0,

rocclr/utils/flags.hpp

@@ -230,10 +230,8 @@ release(size_t, HIP_INITIAL_DM_SIZE, 8 * Mi,                                  \
         "Set initial heap size for device malloc.")                           \
 release(bool, HIP_FORCE_DEV_KERNARG, 0,                                       \
          "Force device mem for kernel args.")                                 \
-release(uint, DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE, 32,                        \
-         "Size of AQL buffering queue")                                       \
-release(bool, DEBUG_CLR_GRAPH_ENABLE_BUFFERING, false,                        \
-         "Enable/Disable graph AQL buffering")                                \
+release(bool, DEBUG_CLR_GRAPH_PACKET_CAPTURE, false,                          \
+         "Enable/Disable graph packet capturing")                             \
 release(cstring, HIPRTC_COMPILE_OPTIONS_APPEND, "",                           \
         "Set compile options needed for hiprtc compilation")                  \
 release(cstring, HIPRTC_LINK_OPTIONS_APPEND, "",                              \