From 45ed09d0cb160d42f7a7cb337a9d5d1be60127ce Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Mon, 23 Jan 2023 17:40:30 -0500 Subject: [PATCH] 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: b23c75974684df8e93a61abb5c5077a57c4df9b3] --- projects/clr/rocclr/device/rocm/rocblit.cpp | 24 +++---------------- .../clr/rocclr/device/rocm/rocvirtual.cpp | 15 ++++++++++-- .../clr/rocclr/device/rocm/rocvirtual.hpp | 3 ++- 3 files changed, 18 insertions(+), 24 deletions(-) diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 2f664b43b9..a7a3efd26e 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/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(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(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); diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 7747d14725..0976b18376 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/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(kernel.getDeviceKernel(dev())); Kernel& gpuKernel = static_cast(*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, diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index bf0bb4cef0..74d624a111 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/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);