From 028339d9be7dfa566867db3065cb28ed0a8e5dc3 Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 12 Jun 2018 16:17:29 -0400 Subject: [PATCH] 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: 63219126984e7dc814bf9ef281067e9681e2627c] --- .../rocclr/runtime/device/rocm/rocblit.cpp | 68 ++++- .../rocclr/runtime/device/rocm/rocblit.hpp | 10 +- .../rocclr/runtime/device/rocm/rocdefs.hpp | 3 + .../rocclr/runtime/device/rocm/rocdevice.cpp | 18 +- .../rocclr/runtime/device/rocm/rockernel.cpp | 43 +++ .../rocclr/runtime/device/rocm/rockernel.hpp | 8 + .../rocclr/runtime/device/rocm/rocprogram.hpp | 2 + .../rocclr/runtime/device/rocm/rocsched.hpp | 75 ++++++ .../rocclr/runtime/device/rocm/rocschedcl.cpp | 17 ++ .../runtime/device/rocm/rocsettings.cpp | 4 + .../runtime/device/rocm/rocsettings.hpp | 2 + .../rocclr/runtime/device/rocm/rocvirtual.cpp | 251 +++++++++++++++++- .../rocclr/runtime/device/rocm/rocvirtual.hpp | 17 ++ 13 files changed, 512 insertions(+), 6 deletions(-) create mode 100644 projects/clr/rocclr/runtime/device/rocm/rocsched.hpp create mode 100644 projects/clr/rocclr/runtime/device/rocm/rocschedcl.cpp diff --git a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp index 98eec0ee95..b47407d2e7 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp +++ b/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 @@ -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(kernels_[Scheduler]->getDeviceKernel(dev())); + Kernel& gpuKernel = static_cast(*devKernel); + + SchedulerParam* sp = reinterpret_cast(schedulerParam->getHostMem()); + memset(sp, 0, sizeof(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_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(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 diff --git a/projects/clr/rocclr/runtime/device/rocm/rocblit.hpp b/projects/clr/rocclr/runtime/device/rocm/rocblit.hpp index 4fbd960fed..fbb1442637 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocblit.hpp +++ b/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 { diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdefs.hpp b/projects/clr/rocclr/runtime/device/rocm/rocdefs.hpp index b13e1235ab..94617bf4ea 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdefs.hpp +++ b/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 { diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp index 0e1f7e22a3..0d83f8cec8 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/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; } diff --git a/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp b/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp index b59847b16c..f0bf4e95c6 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rockernel.cpp +++ b/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(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) { diff --git a/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp b/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp index 926b304d00..f0b8690e71 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rockernel.hpp +++ b/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() 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) {} diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp index 188cef4c0e..9feae26df8 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp +++ b/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: diff --git a/projects/clr/rocclr/runtime/device/rocm/rocsched.hpp b/projects/clr/rocclr/runtime/device/rocm/rocsched.hpp new file mode 100644 index 0000000000..7fdd9d7c48 --- /dev/null +++ b/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 +}; + +} diff --git a/projects/clr/rocclr/runtime/device/rocm/rocschedcl.cpp b/projects/clr/rocclr/runtime/device/rocm/rocschedcl.cpp new file mode 100644 index 0000000000..4249c9b1d3 --- /dev/null +++ b/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 diff --git a/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp b/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp index c2a23ed3ac..a768801d3b 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp +++ b/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) { diff --git a/projects/clr/rocclr/runtime/device/rocm/rocsettings.hpp b/projects/clr/rocclr/runtime/device/rocm/rocsettings.hpp index cb14035a55..1ecd636d2a 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocsettings.hpp +++ b/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 diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp index c94a1b56f3..569a516942 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/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::max(), std::numeric_limits::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(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(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(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(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(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(reinterpret_cast(schedulerParam_->getHostMem()) + sizeof(SchedulerParam)); + memset(wrap, 0, sizeof(AmdAqlWrap)); + wrap->state = AQL_WRAP_DONE; + + uint64_t spVA = reinterpret_cast(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(blitMgr()).runScheduler( + getVQVirtualAddress(), schedulerParam_, schedulerQueue_, schedulerSignal_, schedulerThreads_); + } + return true; } /** diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp index 148699c172..520cc9f515 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp +++ b/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_;