From 517bf08c37c696821e6faca718ce3db346163e33 Mon Sep 17 00:00:00 2001 From: foreman Date: Wed, 12 Jun 2019 10:00:38 -0400 Subject: [PATCH] P4 to Git Change 1809277 by gandryey@gera-win10 on 2019/06/11 17:34:13 SWDEV-180872 - Runtime support changes for Cooperative Group Features - Initial implementation of the core functionality. Disabled by default. Use GPU_ENABLE_COOP_GROUPS=1 to enable the feature. - Runtime uses device queue for cooperative executions with a synchronization on the launched queue. - The current implementation is pure runtime change and it can work if only one app uses this feature. No ROCr/KFD support was added or tested - Only inline assembler was tested Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#20 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#15 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#15 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#17 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#28 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#32 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#338 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#606 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.hpp#171 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palblit.cpp#31 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palblit.hpp#9 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#142 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.hpp#39 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palschedcl.cpp#6 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#135 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.hpp#61 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#32 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.hpp#12 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#127 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#37 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocschedcl.cpp#3 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#75 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.hpp#23 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#94 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/command.hpp#92 edit ... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.hpp#311 edit --- rocclr/runtime/device/device.hpp | 9 +- rocclr/runtime/device/gpu/gpudevice.cpp | 3 +- rocclr/runtime/device/gpu/gpudevice.hpp | 3 +- rocclr/runtime/device/pal/palblit.cpp | 23 ++++ rocclr/runtime/device/pal/palblit.hpp | 7 +- rocclr/runtime/device/pal/paldevice.cpp | 52 ++++++++- rocclr/runtime/device/pal/paldevice.hpp | 6 +- rocclr/runtime/device/pal/palschedcl.cpp | 24 +++- rocclr/runtime/device/pal/palvirtual.cpp | 60 ++++++++-- rocclr/runtime/device/pal/palvirtual.hpp | 11 +- rocclr/runtime/device/rocm/rocblit.cpp | 24 ++++ rocclr/runtime/device/rocm/rocblit.hpp | 7 +- rocclr/runtime/device/rocm/rocdevice.cpp | 43 +++++++ rocclr/runtime/device/rocm/rocdevice.hpp | 17 ++- rocclr/runtime/device/rocm/rocschedcl.cpp | 23 +++- rocclr/runtime/device/rocm/rocvirtual.cpp | 131 +++++++++++++++++++--- rocclr/runtime/device/rocm/rocvirtual.hpp | 12 +- rocclr/runtime/platform/command.cpp | 11 +- rocclr/runtime/platform/command.hpp | 16 ++- rocclr/runtime/utils/flags.hpp | 2 + 20 files changed, 437 insertions(+), 47 deletions(-) diff --git a/rocclr/runtime/device/device.hpp b/rocclr/runtime/device/device.hpp index bec4a1669b..9e7aa6cbd0 100644 --- a/rocclr/runtime/device/device.hpp +++ b/rocclr/runtime/device/device.hpp @@ -507,6 +507,11 @@ struct Info : public amd::EmbeddedObject { //! Max numbers of threads per CU cl_uint maxThreadsPerCU_; + + //! GPU device supports a launch of cooperative groups + cl_bool cooperativeGroups_; + //! GPU device supports a launch of cooperative groups on multiple devices + cl_bool cooperativeMultiDeviceGroups_; }; //! Device settings @@ -1286,7 +1291,9 @@ class Device : public RuntimeObject { virtual void svmFree(void* ptr) const = 0; //! Validate kernel - virtual bool validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev) { + virtual bool validateKernel(const amd::Kernel& kernel, + const device::VirtualDevice* vdev, + bool coop_groups = false) { return true; }; diff --git a/rocclr/runtime/device/gpu/gpudevice.cpp b/rocclr/runtime/device/gpu/gpudevice.cpp index b44312fad9..44de000cd0 100644 --- a/rocclr/runtime/device/gpu/gpudevice.cpp +++ b/rocclr/runtime/device/gpu/gpudevice.cpp @@ -1955,7 +1955,8 @@ bool Device::allocScratch(uint regNum, const VirtualGPU* vgpu) { return true; } -bool Device::validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev) { +bool Device::validateKernel( + const amd::Kernel& kernel, const device::VirtualDevice* vdev, bool coop_groups) { // Find the number of scratch registers used in the kernel const device::Kernel* devKernel = kernel.getDeviceKernel(*this); uint regNum = static_cast(devKernel->workGroupInfo()->scratchRegs_); diff --git a/rocclr/runtime/device/gpu/gpudevice.hpp b/rocclr/runtime/device/gpu/gpudevice.hpp index b1a781c998..e4134f674f 100644 --- a/rocclr/runtime/device/gpu/gpudevice.hpp +++ b/rocclr/runtime/device/gpu/gpudevice.hpp @@ -410,7 +410,8 @@ class Device : public NullDevice, public CALGSLDevice { //! Validates kernel before execution virtual bool validateKernel(const amd::Kernel& kernel, //!< AMD kernel object - const device::VirtualDevice* vdev); + const device::VirtualDevice* vdev, + bool coop_groups = false); virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, cl_set_device_clock_mode_output_amd* pSetClockModeOutput); diff --git a/rocclr/runtime/device/pal/palblit.cpp b/rocclr/runtime/device/pal/palblit.cpp index 1d960a286f..2c817144cf 100644 --- a/rocclr/runtime/device/pal/palblit.cpp +++ b/rocclr/runtime/device/pal/palblit.cpp @@ -2381,6 +2381,29 @@ void KernelBlitManager::writeRawData(device::Memory& memory, size_t size, const synchronize(); } +bool KernelBlitManager::RunGwsInit( + uint32_t value) const { + amd::ScopedLock k(lockXferOps_); + + size_t globalWorkOffset[1] = { 0 }; + size_t globalWorkSize[1] = { 1 }; + size_t localWorkSize[1] = { 1 }; + + // Program kernels arguments + setArgument(kernels_[GwsInit], 0, sizeof(uint32_t), &value); + + // Create ND range object for the kernel's execution + amd::NDRangeContainer ndrange(1, globalWorkOffset, globalWorkSize, localWorkSize); + + // Execute the blit + address parameters = kernels_[GwsInit]->parameters().values(); + bool result = gpu().submitKernelInternal(ndrange, *kernels_[GwsInit], parameters); + + synchronize(); + + return result; +} + amd::Memory* DmaBlitManager::pinHostMemory(const void* hostMem, size_t pinSize, size_t& partial) const { size_t pinAllocSize; diff --git a/rocclr/runtime/device/pal/palblit.hpp b/rocclr/runtime/device/pal/palblit.hpp index 4c9769d678..c8eec28db9 100644 --- a/rocclr/runtime/device/pal/palblit.hpp +++ b/rocclr/runtime/device/pal/palblit.hpp @@ -205,6 +205,7 @@ class KernelBlitManager : public DmaBlitManager { FillBuffer, FillImage, Scheduler, + GwsInit, BlitTotal }; @@ -352,6 +353,10 @@ class KernelBlitManager : public DmaBlitManager { const void* data //!< Raw data pointer ) const; + //! Runs a blit kernel for GWS init + bool RunGwsInit(uint32_t value //!< Initial value for GWS resource + ) const; + virtual amd::Monitor* lockXfer() const { return &lockXferOps_; } private: @@ -406,7 +411,7 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = { "copyImage", "copyImage1DA", "copyImageToBuffer", "copyBufferToImage", "copyBufferRect", "copyBufferRectAligned", "copyBuffer", "copyBufferAligned", "fillBuffer", - "fillImage", "scheduler", + "fillImage", "scheduler", "gwsInit" }; /*@}*/ // namespace pal diff --git a/rocclr/runtime/device/pal/paldevice.cpp b/rocclr/runtime/device/pal/paldevice.cpp index d746fbd3fb..e566ffd049 100644 --- a/rocclr/runtime/device/pal/paldevice.cpp +++ b/rocclr/runtime/device/pal/paldevice.cpp @@ -660,6 +660,8 @@ void NullDevice::fillDeviceInfo(const Pal::DeviceProperties& palProp, info_.pcieRevisionId_ = palProp.revisionId; info_.maxThreadsPerCU_ = info_.wavefrontWidth_ * info_.simdPerCU_ * palProp.gfxipProperties.shaderCore.numWavefrontsPerSimd; + info_.cooperativeGroups_ = GPU_ENABLE_COOP_GROUPS; + info_.cooperativeMultiDeviceGroups_ = GPU_ENABLE_COOP_GROUPS; } } @@ -852,6 +854,7 @@ Device::~Device() { } extern const char* SchedulerSourceCode; +extern const char* GwsInitSourceCode; Pal::IDevice* gDeviceList[Pal::MaxDevices] = {}; uint32_t gStartDevice = 0; uint32_t gNumDevices = 0; @@ -2059,7 +2062,8 @@ bool Device::allocScratch(uint regNum, const VirtualGPU* vgpu, uint vgprs) { return true; } -bool Device::validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev) { +bool Device::validateKernel( + const amd::Kernel& kernel, const device::VirtualDevice* vdev, bool coop_groups) { // Find the number of scratch registers used in the kernel const device::Kernel* devKernel = kernel.getDeviceKernel(*this); uint regNum = static_cast(devKernel->workGroupInfo()->scratchRegs_); @@ -2068,6 +2072,14 @@ bool Device::validateKernel(const amd::Kernel& kernel, const device::VirtualDevi if (!allocScratch(regNum, vgpu, devKernel->workGroupInfo()->usedVGPRs_)) { return false; } + // Runtime plans to launch cooperative groups on the device queue, thus + // validate the scratch buffer on that queue + if (coop_groups) { + vgpu = xferQueue(); + if (!allocScratch(regNum, vgpu, devKernel->workGroupInfo()->usedVGPRs_)) { + return false; + } + } if (devKernel->hsa()) { const HSAILKernel* hsaKernel = static_cast(devKernel); @@ -2224,6 +2236,35 @@ void Device::svmFree(void* ptr) const { } } +bool Device::AcquireExclusiveGpuAccess() { + // Lock the virtual GPU list + vgpusAccess().lock(); + + // Find all available virtual GPUs and lock them + // from the execution of commands + for (uint idx = 0; idx < vgpus().size(); ++idx) { + vgpus()[idx]->execution().lock(); + // Make sure a wait is done + vgpus()[idx]->WaitForIdleCompute(); + } +// if (!hsa_exclusive_gpu_access_) { + // @todo call rocr +// hsa_exclusive_gpu_access_ = true; +// } + return true; +} + +void Device::ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const { + vgpu.WaitForIdleCompute(); + // Find all available virtual GPUs and unlock them + // for the execution of commands + for (uint idx = 0; idx < vgpus().size(); ++idx) { + vgpus()[idx]->execution().unlock(); + } + + // Unock the virtual GPU list + vgpusAccess().unlock(); +} Device::SrdManager::~SrdManager() { for (uint i = 0; i < pool_.size(); ++i) { @@ -2333,7 +2374,7 @@ bool Device::createBlitProgram() { bool result = true; // Delayed compilation due to brig_loader memory allocation - const char* scheduler = nullptr; + const char* blits = nullptr; const char* ocl20 = nullptr; std::string sch = SchedulerSourceCode; @@ -2348,14 +2389,17 @@ bool Device::createBlitProgram() { sch.replace(loc, sizeof(AmdScheduler) - 1, AmdSchedulerPal); loc = sch.find(AmdScheduler, (loc + sizeof(AmdSchedulerPal) - 1)); sch.replace(loc, sizeof(AmdScheduler) - 1, AmdSchedulerPal); + if (info().cooperativeGroups_) { + sch.append(GwsInitSourceCode); + } } - scheduler = sch.c_str(); + blits = sch.c_str(); ocl20 = "-cl-std=CL2.0"; } blitProgram_ = new BlitProgram(context_); // Create blit programs - if (blitProgram_ == nullptr || !blitProgram_->create(this, scheduler, ocl20)) { + if (blitProgram_ == nullptr || !blitProgram_->create(this, blits, ocl20)) { delete blitProgram_; blitProgram_ = nullptr; LogError("Couldn't create blit kernels!"); diff --git a/rocclr/runtime/device/pal/paldevice.hpp b/rocclr/runtime/device/pal/paldevice.hpp index d4448d0ca2..1f51b560fb 100644 --- a/rocclr/runtime/device/pal/paldevice.hpp +++ b/rocclr/runtime/device/pal/paldevice.hpp @@ -342,7 +342,8 @@ class Device : public NullDevice { //! Validates kernel before execution virtual bool validateKernel(const amd::Kernel& kernel, //!< AMD kernel object - const device::VirtualDevice* vdev); + const device::VirtualDevice* vdev, + bool coop_group = false); virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, cl_set_device_clock_mode_output_amd* pSetClockModeOutput); @@ -525,6 +526,9 @@ class Device : public NullDevice { } } + bool AcquireExclusiveGpuAccess(); + void ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const; + private: static void PAL_STDCALL PalDeveloperCallback(void* pPrivateData, const Pal::uint32 deviceIndex, Pal::Developer::CallbackType type, void* pCbData); diff --git a/rocclr/runtime/device/pal/palschedcl.cpp b/rocclr/runtime/device/pal/palschedcl.cpp index f912c5fa7a..159c564b95 100644 --- a/rocclr/runtime/device/pal/palschedcl.cpp +++ b/rocclr/runtime/device/pal/palschedcl.cpp @@ -3,9 +3,9 @@ // namespace pal { -#define SCHEDULER_KERNEL(...) #__VA_ARGS__ +#define BLIT_KERNEL(...) #__VA_ARGS__ -const char* SchedulerSourceCode = SCHEDULER_KERNEL( +const char* SchedulerSourceCode = BLIT_KERNEL( %s \n extern void __amd_scheduler(__global void*, __global void*, uint); @@ -15,4 +15,24 @@ __kernel void scheduler(__global void* queue, __global void* params, uint paramI } \n); +const char* GwsInitSourceCode = BLIT_KERNEL( +\n +__kernel void gwsInit(uint value) { + unsigned int m0_backup, new_m0; + __asm__ __volatile__( + "s_mov_b32 %0 m0\n" + "v_readfirstlane_b32 %1 %2\n" + "s_nop 0\n" + "s_mov_b32 m0 %1\n" + "s_nop 0\n" + "ds_gws_init %3 offset:0 gds\n" + "s_waitcnt lgkmcnt(0) expcnt(0)\n" + "s_mov_b32 m0 %0\n" + "s_nop 0" + : "=s"(m0_backup), "=s"(new_m0) + : "v"(0 << 0x10), "{v0}"(value - 1) + : "memory"); +} +\n); + } // namespace pal diff --git a/rocclr/runtime/device/pal/palvirtual.cpp b/rocclr/runtime/device/pal/palvirtual.cpp index 5197810ddd..30b9b185a2 100644 --- a/rocclr/runtime/device/pal/palvirtual.cpp +++ b/rocclr/runtime/device/pal/palvirtual.cpp @@ -2166,24 +2166,62 @@ void VirtualGPU::PostDeviceEnqueue(const amd::Kernel& kernel, const HSAILKernel& // ================================================================================================ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) { - // Make sure VirtualGPU has an exclusive access to the resources - amd::ScopedLock lock(execution()); + if (vcmd.cooperativeGroups()) { + uint32_t workgroups = 0; + for (uint i = 0; i < vcmd.sizes().dimensions(); i++) { + if ((vcmd.sizes().local()[i] != 0) && (vcmd.sizes().global()[i] != 1)) { + workgroups += (vcmd.sizes().global()[i] / vcmd.sizes().local()[i]); + } + } + uint32_t counter = workgroups * + amd::alignUp(vcmd.sizes().local().product(), dev().info().wavefrontWidth_) / + dev().info().wavefrontWidth_; - profilingBegin(vcmd); + bool test = true; + VirtualGPU* queue = (test) ? this : dev().xferQueue(); - // Submit kernel to HW - if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), false, &vcmd.event(), - vcmd.sharedMemBytes())) { - vcmd.setStatus(CL_INVALID_OPERATION); + // Wait for the execution on the current queue, since the coop groups will use the device queue + waitAllEngines(); + + amd::ScopedLock lock(queue->blitMgr().lockXfer()); + + queue->profilingBegin(vcmd); + + static_cast(queue->blitMgr()).RunGwsInit(counter); + queue->addBarrier(RgpSqqtBarrierReason::PostDeviceEnqueue); + + // Submit kernel to HW + if (!queue->submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), + false, &vcmd.event(), vcmd.sharedMemBytes(), + vcmd.cooperativeGroups())) { + vcmd.setStatus(CL_INVALID_OPERATION); + } + + queue->profilingEnd(vcmd); + + // Wait for the execution on the device queue. Keep the current queue in-order + queue->waitAllEngines(); + } else { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + + profilingBegin(vcmd); + + // Submit kernel to HW + if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), false, &vcmd.event(), + vcmd.sharedMemBytes(), vcmd.cooperativeGroups())) { + vcmd.setStatus(CL_INVALID_OPERATION); + } + + profilingEnd(vcmd); } - - profilingEnd(vcmd); } // ================================================================================================ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel, const_address parameters, bool nativeMem, - amd::Event* enqueueEvent, uint32_t sharedMemBytes) { + amd::Event* enqueueEvent, uint32_t sharedMemBytes, + bool cooperativeGroup) { size_t newOffset[3] = {0, 0, 0}; size_t newGlobalSize[3] = {0, 0, 0}; @@ -2239,6 +2277,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const LogError("Wrong memory objects!"); return false; } + bool needFlush = false; // Avoid flushing when PerfCounter is enabled, to make sure PerfStart/dispatch/PerfEnd // are in the same cmdBuffer @@ -2358,6 +2397,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const return false; } } + // Perform post dispatch logic for RGP traces if (rgpCaptureEna()) { dev().rgpCaptureMgr()->PostDispatch(this); diff --git a/rocclr/runtime/device/pal/palvirtual.hpp b/rocclr/runtime/device/pal/palvirtual.hpp index 9e557e1f03..61242af0f2 100644 --- a/rocclr/runtime/device/pal/palvirtual.hpp +++ b/rocclr/runtime/device/pal/palvirtual.hpp @@ -298,7 +298,8 @@ class VirtualGPU : public device::VirtualDevice { const_address parameters, //!< Parameters for the kernel bool nativeMem = true, //!< Native memory objects amd::Event* enqueueEvent = nullptr, //!< Event provided in the enqueue kernel command - uint32_t sharedMemBytes = 0 //!< Shared memory size + uint32_t sharedMemBytes = 0, //!< Shared memory size + bool cooperativeGroups = false //!< TRUE if cooperative groups mode is required ); void submitNativeFn(amd::NativeFnCommand& vcmd); void submitFillMemory(amd::FillMemoryCommand& vcmd); @@ -512,6 +513,14 @@ class VirtualGPU : public device::VirtualDevice { //! Checks if RGP capture is enabled bool rgpCaptureEna() const { return state_.rgpCaptureEnabled_; } + //! Waits for idle on compute engine + void WaitForIdleCompute() { + if (events_[MainEngine].isValid()) { + queues_[events_[MainEngine].engineId_]->waitForEvent(events_[MainEngine].id_); + events_[MainEngine].invalidate(); + } + } + protected: void profileEvent(EngineType engine, bool type) const; diff --git a/rocclr/runtime/device/rocm/rocblit.cpp b/rocclr/runtime/device/rocm/rocblit.cpp index fa9b3b766c..11011aa550 100644 --- a/rocclr/runtime/device/rocm/rocblit.cpp +++ b/rocclr/runtime/device/rocm/rocblit.cpp @@ -2235,4 +2235,28 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, return true; } +bool KernelBlitManager::RunGwsInit( + uint32_t value) const { + amd::ScopedLock k(lockXferOps_); + + size_t globalWorkOffset[1] = { 0 }; + size_t globalWorkSize[1] = { 1 }; + size_t localWorkSize[1] = { 1 }; + + // Program kernels arguments + setArgument(kernels_[GwsInit], 0, sizeof(uint32_t), &value); + + // Create ND range object for the kernel's execution + amd::NDRangeContainer ndrange(1, globalWorkOffset, globalWorkSize, localWorkSize); + + // Execute the blit + address parameters = captureArguments(kernels_[GwsInit]); + + bool result = gpu().submitKernelInternal(ndrange, *kernels_[GwsInit], parameters, nullptr); + + releaseArguments(parameters); + + return result; +} + } // namespace pal diff --git a/rocclr/runtime/device/rocm/rocblit.hpp b/rocclr/runtime/device/rocm/rocblit.hpp index 0b556ea864..d2103005eb 100644 --- a/rocclr/runtime/device/rocm/rocblit.hpp +++ b/rocclr/runtime/device/rocm/rocblit.hpp @@ -234,6 +234,7 @@ class KernelBlitManager : public DmaBlitManager { FillBuffer, FillImage, Scheduler, + GwsInit, BlitTotal }; @@ -374,6 +375,10 @@ class KernelBlitManager : public DmaBlitManager { hsa_signal_t& schedulerSignal, uint threads); + //! Runs a blit kernel for GWS init + bool RunGwsInit(uint32_t value //!< Initial value for GWS resource + ) const; + private: static const size_t MaxXferBuffers = 2; static const uint TransferSplitSize = 1; @@ -434,7 +439,7 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = { "copyImage", "copyImage1DA", "copyImageToBuffer", "copyBufferToImage", "copyBufferRect", "copyBufferRectAligned", "copyBuffer", "copyBufferAligned", "fillBuffer", - "fillImage", "scheduler", + "fillImage", "scheduler", "gwsInit" }; inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index, size_t size, const void* value) const { diff --git a/rocclr/runtime/device/rocm/rocdevice.cpp b/rocclr/runtime/device/rocm/rocdevice.cpp index 4052daa821..6609d3d5ba 100644 --- a/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/rocclr/runtime/device/rocm/rocdevice.cpp @@ -148,6 +148,7 @@ Device::Device(hsa_agent_t bkendDevice) , pro_device_(nullptr) , pro_ena_(false) , freeMem_(0) + , hsa_exclusive_gpu_access_(false) , numOfVgpus_(0) { group_segment_.handle = 0; system_segment_.handle = 0; @@ -585,6 +586,7 @@ bool Device::init() { } extern const char* SchedulerSourceCode; +extern const char* GwsInitSourceCode; void Device::tearDown() { NullDevice::tearDown(); @@ -648,6 +650,9 @@ bool Device::create(bool sramEccEnabled) { #if defined(WITH_LIGHTNING_COMPILER) || defined(USE_COMGR_LIBRARY) std::string sch = SchedulerSourceCode; if (settings().useLightning_) { + if (info().cooperativeGroups_) { + sch.append(GwsInitSourceCode); + } scheduler = sch.c_str(); } #ifndef USE_COMGR_LIBRARY @@ -781,6 +786,38 @@ device::Program* NullDevice::createProgram(amd::option::Options* options) { return program; } +bool Device::AcquireExclusiveGpuAccess() { + // Lock the virtual GPU list + vgpusAccess().lock(); + + // Find all available virtual GPUs and lock them + // from the execution of commands + for (uint idx = 0; idx < vgpus().size(); ++idx) { + vgpus()[idx]->execution().lock(); + // Make sure a wait is done + vgpus()[idx]->releaseGpuMemoryFence(); + } + if (!hsa_exclusive_gpu_access_) { + // @todo call rocr + hsa_exclusive_gpu_access_ = true; + } + return true; +} + +void Device::ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const { + // Make sure the operation is done + vgpu.releaseGpuMemoryFence(); + + // Find all available virtual GPUs and unlock them + // for the execution of commands + for (uint idx = 0; idx < vgpus().size(); ++idx) { + vgpus()[idx]->execution().unlock(); + } + + // Unock the virtual GPU list + vgpusAccess().unlock(); +} + device::Program* Device::createProgram(amd::option::Options* options) { device::Program* program; if (settings().useLightning_) { @@ -1339,6 +1376,8 @@ bool Device::populateOCLDeviceConstants() { //TODO: set to true once thread trace support is available info_.threadTraceEnable_ = false; info_.pcieDeviceId_ = deviceInfo_.pciDeviceId_; + info_.cooperativeGroups_ = GPU_ENABLE_COOP_GROUPS; + info_.cooperativeMultiDeviceGroups_ = GPU_ENABLE_COOP_GROUPS; } info_.maxPipePacketSize_ = info_.maxMemAllocSize_; @@ -1356,8 +1395,12 @@ bool Device::populateOCLDeviceConstants() { } device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) { + amd::ScopedLock lock(vgpusAccess()); + bool profiling = (queue != nullptr) && queue->properties().test(CL_QUEUE_PROFILING_ENABLE); + profiling |= (queue == nullptr) ? true : false; + // Initialization of heap and other resources occur during the command // queue creation time. VirtualGPU* virtualDevice = new VirtualGPU(*this); diff --git a/rocclr/runtime/device/rocm/rocdevice.hpp b/rocclr/runtime/device/rocm/rocdevice.hpp index 4ffac0a454..7b7b8fb58c 100644 --- a/rocclr/runtime/device/rocm/rocdevice.hpp +++ b/rocclr/runtime/device/rocm/rocdevice.hpp @@ -393,6 +393,19 @@ class Device : public NullDevice { virtual amd::Memory* IpcAttach(const void* handle, size_t mem_size, unsigned int flags, void** dev_ptr) const; virtual void IpcDetach (amd::Memory& memory) const; + bool AcquireExclusiveGpuAccess(); + void ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const; + + //! Returns the lock object for the virtual gpus list + amd::Monitor& vgpusAccess() const { return vgpusAccess_; } + + typedef std::vector VirtualGPUs; + //! Returns the list of all virtual GPUs running on this device + const VirtualGPUs& vgpus() const { return vgpus_; } + VirtualGPUs vgpus_; //!< The list of all running virtual gpus (lock protected) + + VirtualGPU* xferQueue() const; + private: static hsa_ven_amd_loader_1_00_pfn_t amd_loader_ext_table; @@ -416,13 +429,13 @@ class Device : public NullDevice { amd::Context* context_; //!< A dummy context for internal data transfer VirtualGPU* xferQueue_; //!< Transfer queue, created on demand - VirtualGPU* xferQueue() const; - XferBuffers* xferRead_; //!< Transfer buffers read XferBuffers* xferWrite_; //!< Transfer buffers write const IProDevice* pro_device_; //!< AMDGPUPro device bool pro_ena_; //!< Extra functionality with AMDGPUPro device, beyond ROCr std::atomic freeMem_; //!< Total of free memory available + mutable amd::Monitor vgpusAccess_; //!< Lock to serialise virtual gpu list access + bool hsa_exclusive_gpu_access_; //!< TRUE if current device was moved into exclusive GPU access mode public: amd::Atomic numOfVgpus_; //!< Virtual gpu unique index diff --git a/rocclr/runtime/device/rocm/rocschedcl.cpp b/rocclr/runtime/device/rocm/rocschedcl.cpp index 2db69ca531..d0ca5651cd 100644 --- a/rocclr/runtime/device/rocm/rocschedcl.cpp +++ b/rocclr/runtime/device/rocm/rocschedcl.cpp @@ -3,9 +3,9 @@ // namespace roc { -#define SCHEDULER_KERNEL(...) #__VA_ARGS__ +#define BLIT_KERNEL(...) #__VA_ARGS__ -const char* SchedulerSourceCode = SCHEDULER_KERNEL( +const char* SchedulerSourceCode = BLIT_KERNEL( \n extern void __amd_scheduler_rocm(__global void*); \n @@ -13,5 +13,24 @@ __kernel void scheduler(__global void* params) { __amd_scheduler_rocm(params); } \n); +const char* GwsInitSourceCode = BLIT_KERNEL( +\n +__kernel void gwsInit(uint value) { + unsigned int m0_backup, new_m0; + __asm__ __volatile__( + "s_mov_b32 %0 m0\n" + "v_readfirstlane_b32 %1 %2\n" + "s_nop 0\n" + "s_mov_b32 m0 %1\n" + "s_nop 0\n" + "ds_gws_init %3 offset:0 gds\n" + "s_waitcnt lgkmcnt(0) expcnt(0)\n" + "s_mov_b32 m0 %0\n" + "s_nop 0" + : "=s"(m0_backup), "=s"(new_m0) + : "v"(0 << 0x10), "{v0}"(value - 1) + : "memory"); +} +\n); } // namespace roc diff --git a/rocclr/runtime/device/rocm/rocvirtual.cpp b/rocclr/runtime/device/rocm/rocvirtual.cpp index 405455d103..67d62e75d7 100644 --- a/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/rocclr/runtime/device/rocm/rocvirtual.cpp @@ -186,13 +186,16 @@ void VirtualGPU::MemoryDependency::clear(bool all) { } } -bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address params, size_t& ldsAddress) { +bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address params, + size_t& ldsAddress, bool cooperativeGroups) { Kernel& hsaKernel = const_cast(static_cast(*(kernel.getDeviceKernel(dev())))); const amd::KernelSignature& signature = kernel.signature(); const amd::KernelParameters& kernelParams = kernel.parameters(); - // AQL packets - setAqlHeader(kDispatchPacketHeaderNoSync); + if (!cooperativeGroups) { + // AQL packets + setAqlHeader(kDispatchPacketHeaderNoSync); + } // Mark the tracker with a new kernel, // so we can avoid checks of the aliased objects @@ -531,6 +534,7 @@ bool VirtualGPU::releaseGpuMemoryFence() { VirtualGPU::VirtualGPU(Device& device) : device::VirtualDevice(device), + execution_("Virtual GPU execution lock", true), roc_device_(device), virtualQueue_(nullptr), deviceQueueSize_(0), @@ -553,6 +557,10 @@ VirtualGPU::VirtualGPU(Device& device) kernarg_pool_cur_offset_ = 0; aqlHeader_ = kDispatchPacketHeaderNoSync; barrier_signal_.handle = 0; + + // Note: Virtual GPU device creation must be a thread safe operation + roc_device_.vgpus_.resize(roc_device_.numOfVgpus_); + roc_device_.vgpus_[index()] = this; } VirtualGPU::~VirtualGPU() { @@ -598,6 +606,10 @@ VirtualGPU::~VirtualGPU() { } --roc_device_.numOfVgpus_; // Virtual gpu unique index decrementing + roc_device_.vgpus_.erase(roc_device_.vgpus_.begin() + index()); + for (uint idx = index(); idx < roc_device_.vgpus().size(); ++idx) { + roc_device_.vgpus()[idx]->index_--; + } } bool VirtualGPU::create(bool profilingEna) { @@ -844,6 +856,9 @@ void VirtualGPU::updateCommandsState(amd::Command* list) { } void VirtualGPU::submitReadMemory(amd::ReadMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -932,6 +947,9 @@ void VirtualGPU::submitReadMemory(amd::ReadMemoryCommand& cmd) { } void VirtualGPU::submitWriteMemory(amd::WriteMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1026,6 +1044,9 @@ void VirtualGPU::submitWriteMemory(amd::WriteMemoryCommand& cmd) { } void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // in-order semantics: previous commands need to be done before we start releaseGpuMemoryFence(); @@ -1125,6 +1146,9 @@ bool VirtualGPU::copyMemory(cl_command_type type, amd::Memory& srcMem, amd::Memo } void VirtualGPU::submitCopyMemory(amd::CopyMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1142,6 +1166,9 @@ void VirtualGPU::submitCopyMemory(amd::CopyMemoryCommand& cmd) { } void VirtualGPU::submitSvmCopyMemory(amd::SvmCopyMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // in-order semantics: previous commands need to be done before we start releaseGpuMemoryFence(); @@ -1218,6 +1245,9 @@ void VirtualGPU::submitSvmCopyMemory(amd::SvmCopyMemoryCommand& cmd) { } void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1303,6 +1333,9 @@ void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) { } void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1339,6 +1372,9 @@ void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) { } void VirtualGPU::submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1376,6 +1412,9 @@ void VirtualGPU::submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd) { } void VirtualGPU::submitMapMemory(amd::MapMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1471,6 +1510,9 @@ void VirtualGPU::submitMapMemory(amd::MapMemoryCommand& cmd) { } void VirtualGPU::submitUnmapMemory(amd::UnmapMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + roc::Memory* devMemory = static_cast(cmd.memory().getDeviceMemory(dev(), false)); const device::Memory::WriteMapInfo* mapInfo = devMemory->writeMapInfo(cmd.mapPtr()); @@ -1560,6 +1602,9 @@ void VirtualGPU::submitUnmapMemory(amd::UnmapMemoryCommand& cmd) { bool VirtualGPU::fillMemory(cl_command_type type, amd::Memory* amdMemory, const void* pattern, size_t patternSize, const amd::Coord3D& origin, const amd::Coord3D& size) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + Memory* memory = dev().getRocMemory(amdMemory); bool entire = amdMemory->isEntirelyCovered(origin, size); @@ -1615,6 +1660,9 @@ bool VirtualGPU::fillMemory(cl_command_type type, amd::Memory* amdMemory, const } void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1628,6 +1676,9 @@ void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) { } void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // in-order semantics: previous commands need to be done before we start releaseGpuMemoryFence(); @@ -1669,6 +1720,9 @@ void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) { } void VirtualGPU::submitMigrateMemObjects(amd::MigrateMemObjectsCommand& vcmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + // Wait on a kernel if one is outstanding releaseGpuMemoryFence(); @@ -1875,13 +1929,13 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize) } bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel, - const_address parameters, void* eventHandle, uint32_t sharedMemBytes) { + const_address parameters, void* eventHandle, uint32_t sharedMemBytes, bool cooperativeGroups) { device::Kernel* devKernel = const_cast(kernel.getDeviceKernel(dev())); Kernel& gpuKernel = static_cast(*devKernel); size_t ldsUsage = gpuKernel.WorkgroupGroupSegmentByteSize(); // Check memory dependency and SVM objects - if (!processMemObjects(kernel, parameters, ldsUsage)) { + if (!processMemObjects(kernel, parameters, ldsUsage, cooperativeGroups)) { LogError("Wrong memory objects!"); return false; } @@ -2087,16 +2141,57 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const * the list of kernel parameters. */ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) { - profilingBegin(vcmd); + if (vcmd.cooperativeGroups()) { + uint32_t workgroups = 0; + for (uint i = 0; i < vcmd.sizes().dimensions(); i++) { + if ((vcmd.sizes().local()[i] != 0) && (vcmd.sizes().global()[i] != 1)) { + workgroups += (vcmd.sizes().global()[i] / vcmd.sizes().local()[i]); + } + } + uint32_t counter = workgroups * + amd::alignUp(vcmd.sizes().local().product(), dev().info().wavefrontWidth_) / + dev().info().wavefrontWidth_; - // Submit kernel to HW - if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), - static_cast(as_cl(&vcmd.event())), vcmd.sharedMemBytes())) { - LogError("AQL dispatch failed!"); - vcmd.setStatus(CL_INVALID_OPERATION); + // Get device queue for exclusive GPU access + VirtualGPU* queue = dev().xferQueue(); + + // Wait for the execution on the current queue, since the coop groups will use the device queue + releaseGpuMemoryFence(); + + // Lock the queue, using the blit manager lock + amd::ScopedLock lock(queue->blitMgr().lockXfer()); + queue->profilingBegin(vcmd); + + static_cast(queue->blitMgr()).RunGwsInit(counter); + + // Sync AQL packets + queue->setAqlHeader(kDispatchPacketHeader); + + // Submit kernel to HW + if (!queue->submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), + static_cast(as_cl(&vcmd.event())), vcmd.sharedMemBytes(), vcmd.cooperativeGroups())) { + LogError("AQL dispatch failed!"); + vcmd.setStatus(CL_INVALID_OPERATION); + } + // Wait for the execution on the device queue. Keep the current queue in-order + queue->releaseGpuMemoryFence(); + + queue->profilingEnd(vcmd); + } else { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + + profilingBegin(vcmd); + + // Submit kernel to HW + if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), + static_cast(as_cl(&vcmd.event())), vcmd.sharedMemBytes(), vcmd.cooperativeGroups())) { + LogError("AQL dispatch failed!"); + vcmd.setStatus(CL_INVALID_OPERATION); + } + + profilingEnd(vcmd); } - - profilingEnd(vcmd); } void VirtualGPU::submitNativeFn(amd::NativeFnCommand& cmd) { @@ -2108,6 +2203,9 @@ void VirtualGPU::submitMarker(amd::Marker& cmd) { } void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + profilingBegin(vcmd); auto fence = kBarrierAcquirePacket; dispatchAqlPacket(&fence, false); @@ -2115,6 +2213,8 @@ void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) { } void VirtualGPU::submitReleaseExtObjects(amd::ReleaseExtObjectsCommand& vcmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); profilingBegin(vcmd); auto fence = kBarrierReleasePacket; dispatchAqlPacket(&fence, false); @@ -2176,6 +2276,9 @@ amd::Memory* VirtualGPU::findPinnedMem(void* addr, size_t size) { void VirtualGPU::enableSyncBlit() const { blitMgr_->enableSynchronization(); } void VirtualGPU::submitTransferBufferFromFile(amd::TransferBufferFileCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + size_t copySize = cmd.size()[0]; size_t fileOffset = cmd.fileOffset(); Memory* mem = dev().getRocMemory(&cmd.memory()); @@ -2228,6 +2331,8 @@ void VirtualGPU::submitTransferBufferFromFile(amd::TransferBufferFileCommand& cm } void VirtualGPU::submitPerfCounter(amd::PerfCounterCommand& vcmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); const amd::PerfCounterCommand::PerfCounterList counters = vcmd.getCounters(); diff --git a/rocclr/runtime/device/rocm/rocvirtual.hpp b/rocclr/runtime/device/rocm/rocvirtual.hpp index cf33fe3b55..ea84bf1764 100644 --- a/rocclr/runtime/device/rocm/rocvirtual.hpp +++ b/rocclr/runtime/device/rocm/rocvirtual.hpp @@ -167,7 +167,8 @@ class VirtualGPU : public device::VirtualDevice { const amd::Kernel& kernel, //!< Kernel for execution const_address parameters, //!< Parameters for the kernel void* event_handle, //!< Handle to OCL event for debugging - uint32_t sharedMemBytes = 0 //!< Shared memory size + uint32_t sharedMemBytes = 0, //!< Shared memory size + bool cooperativeGroups = false //!< TRUE if cooperative groups mode is required ); void submitNativeFn(amd::NativeFnCommand& cmd); void submitMarker(amd::Marker& cmd); @@ -220,7 +221,8 @@ class VirtualGPU : public device::VirtualDevice { //! Detects memory dependency for HSAIL kernels and uses appropriate AQL header bool processMemObjects(const amd::Kernel& kernel, //!< AMD kernel object for execution const_address params, //!< Pointer to the param's store - size_t& ldsAddress //!< LDS usage + size_t& ldsAddress, //!< LDS usage + bool cooperativeGroups //!< Dispatch with cooperative groups ); // Retun the virtual gpu unique index uint index() const { return index_; } @@ -242,6 +244,9 @@ class VirtualGPU : public device::VirtualDevice { void enableSyncBlit() const; + //! Returns the monitor object for execution access by VirtualGPU + amd::Monitor& execution() { return execution_; } + // } roc OpenCL integration private: bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, bool blocking = true); @@ -296,6 +301,7 @@ class VirtualGPU : public device::VirtualDevice { * used to synchronized on kernel outputs. */ bool hasPendingDispatch_; + amd::Monitor execution_; //!< Lock to serialise access to all device objects Timestamp* timestamp_; hsa_agent_t gpu_device_; //!< Physical device hsa_queue_t* gpu_queue_; //!< Queue associated with a gpu @@ -321,7 +327,7 @@ class VirtualGPU : public device::VirtualDevice { uint kernarg_pool_cur_offset_; std::vector signal_pool_; //!< Pool of signals for profiling - const uint index_; //!< Virtual gpu unique index + uint index_; //!< Virtual gpu unique index friend class Timestamp; // PM4 packet for gfx8 performance counter diff --git a/rocclr/runtime/platform/command.cpp b/rocclr/runtime/platform/command.cpp index 034a3182e7..83bbb16fd1 100644 --- a/rocclr/runtime/platform/command.cpp +++ b/rocclr/runtime/platform/command.cpp @@ -226,8 +226,13 @@ void Command::enqueue() { const Context& Command::context() const { return queue_->context(); } NDRangeKernelCommand::NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList, - Kernel& kernel, const NDRangeContainer& sizes, uint32_t sharedMemBytes) - : Command(queue, CL_COMMAND_NDRANGE_KERNEL, eventWaitList), kernel_(kernel), sizes_(sizes), sharedMemBytes_(sharedMemBytes) { + Kernel& kernel, const NDRangeContainer& sizes, + uint32_t sharedMemBytes, uint32_t extraParam) + : Command(queue, CL_COMMAND_NDRANGE_KERNEL, eventWaitList) + , kernel_(kernel) + , sizes_(sizes) + , sharedMemBytes_(sharedMemBytes) + , extraParam_(extraParam) { auto& device = queue.device(); auto devKernel = const_cast(kernel.getDeviceKernel(device)); profilingInfo_.setCallback(devKernel->getProfilingCallback( @@ -393,7 +398,7 @@ bool MigrateMemObjectsCommand::validateMemory() { cl_int NDRangeKernelCommand::captureAndValidate() { const amd::Device& device = queue()->device(); // Validate the kernel before submission - if (!queue()->device().validateKernel(kernel(), queue()->vdev())) { + if (!queue()->device().validateKernel(kernel(), queue()->vdev(), cooperativeGroups())) { return CL_OUT_OF_RESOURCES; } diff --git a/rocclr/runtime/platform/command.hpp b/rocclr/runtime/platform/command.hpp index bad260289e..cf8b2ed507 100644 --- a/rocclr/runtime/platform/command.hpp +++ b/rocclr/runtime/platform/command.hpp @@ -753,11 +753,18 @@ class NDRangeKernelCommand : public Command { NDRangeContainer sizes_; address parameters_; uint32_t sharedMemBytes_; + uint32_t extraParam_; public: + enum { + CooperativeGroups = 0x01, + CooperativeMultiDeviceGroups = 0x02, + }; + //! Construct an ExecuteKernel command NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList, Kernel& kernel, - const NDRangeContainer& sizes, uint32_t sharedMemBytes = 0); + const NDRangeContainer& sizes, uint32_t sharedMemBytes = 0, + uint32_t extraParam = 0); virtual void submit(device::VirtualDevice& device) { device.submitKernel(*this); } @@ -776,6 +783,13 @@ class NDRangeKernelCommand : public Command { //! Return the shared memory size uint32_t sharedMemBytes() const { return sharedMemBytes_; } + //! Return the cooperative groups mode + bool cooperativeGroups() const { return (extraParam_ & CooperativeGroups) ? true : false; } + + //! Return the cooperative multi device groups mode + bool cooperativeMultiDeviceGroups() const { + return (extraParam_ & CooperativeMultiDeviceGroups) ? true : false; } + //! Set the local work size. void setLocalWorkSize(const NDRange& local) { sizes_.local() = local; } diff --git a/rocclr/runtime/utils/flags.hpp b/rocclr/runtime/utils/flags.hpp index 339e31845f..7e91958515 100644 --- a/rocclr/runtime/utils/flags.hpp +++ b/rocclr/runtime/utils/flags.hpp @@ -173,6 +173,8 @@ release(bool, GPU_ENABLE_LC, IS_LIGHTNING, \ "Enables LC path") \ release(bool, GPU_ENABLE_HW_P2P, false, \ "Enables HW P2P path") \ +release(bool, GPU_ENABLE_COOP_GROUPS, false, \ + "Enables cooperative group launch") \ release(uint, GPU_MAX_COMMAND_BUFFERS, 8, \ "The maximum number of command buffers allocated per queue") \ release(cstring, HIP_VISIBLE_DEVICES, "", \