SWDEV-372790 - Copy AQL packet from runtime setup

Scheduler in device queue requires relaunching itself. Make sure
scheduler uses exactly the same AQL packet as the host launch.

Change-Id: I4eb03c4c91bf2408a6d4607731f081a2e2c2c8ae


[ROCm/clr commit: b23c759746]

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

@@ -2181,7 +2181,7 @@ bool KernelBlitManager::fillBuffer3D(device::Memory& memory, const void* pattern
 // ================================================================================================
 bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& dstMemory,
                                    const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin,
-                                   const amd::Coord3D& sizeIn, bool entire, 
+                                   const amd::Coord3D& sizeIn, bool entire,
                                    amd::CopyMetadata copyMetadata) const {
   amd::ScopedLock k(lockXferOps_);
   bool result = false;
@@ -2681,31 +2681,13 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* 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_store_relaxed(schedulerSignal, kInitSignalValueOne);
 
-  sp->scheduler_aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
-                             (1 << HSA_PACKET_HEADER_BARRIER) |
-                             (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
-                             (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
-  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);
@@ -2720,7 +2702,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
   address parameters = captureArguments(kernels_[Scheduler]);
 
   if (!gpu().submitKernelInternal(ndrange, *kernels_[Scheduler],
-                                  parameters, nullptr)) {
+                                  parameters, nullptr, 0, nullptr, &sp->scheduler_aql)) {
     return false;
   }
   releaseArguments(parameters);

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

@@ -1817,7 +1817,7 @@ bool VirtualGPU::copyMemory(cl_command_type type, amd::Memory& srcMem, amd::Memo
         realSize.c[0] *= elemSize;
       }
 
-      result = blitMgr().copyBuffer(*srcDevMem, *dstDevMem, realSrcOrigin, realDstOrigin, 
+      result = blitMgr().copyBuffer(*srcDevMem, *dstDevMem, realSrcOrigin, realDstOrigin,
                                     realSize, entire, copyMetadata);
       break;
     }
@@ -2781,7 +2781,8 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
 // ================================================================================================
 bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
     const amd::Kernel& kernel, const_address parameters, void* eventHandle,
-    uint32_t sharedMemBytes, amd::NDRangeKernelCommand* vcmd) {
+    uint32_t sharedMemBytes, amd::NDRangeKernelCommand* vcmd,
+    hsa_kernel_dispatch_packet_t* aql_packet) {
   device::Kernel* devKernel = const_cast<device::Kernel*>(kernel.getDeviceKernel(dev()));
   Kernel& gpuKernel = static_cast<Kernel&>(*devKernel);
   size_t ldsUsage = gpuKernel.WorkgroupGroupSegmentByteSize();
@@ -3108,6 +3109,16 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
       dispatchPacket.reserved2 = vcmd->profilingInfo().correlation_id_;
     }
 
+    // Copy scheduler's AQL packet for possible relaunch from the scheduler itself
+    if (aql_packet != nullptr) {
+      *aql_packet = dispatchPacket;
+      aql_packet->header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+                           (1 << HSA_PACKET_HEADER_BARRIER) |
+                           (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+                           (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,

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

@@ -312,7 +312,8 @@ class VirtualGPU : public device::VirtualDevice {
                             const_address parameters,            //!< Parameters for the kernel
                             void* event_handle,  //!< Handle to OCL event for debugging
                             uint32_t sharedMemBytes = 0, //!< Shared memory size
-                            amd::NDRangeKernelCommand* vcmd = nullptr //!< Original launch command
+                            amd::NDRangeKernelCommand* vcmd = nullptr, //!< Original launch command
+                            hsa_kernel_dispatch_packet_t* aql_packet = nullptr  //!< Scheduler launch
                             );
   void submitNativeFn(amd::NativeFnCommand& cmd);
   void submitMarker(amd::Marker& cmd);