From d0a7ae02cf15c50bd73f40f7784071a606a5166c Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Fri, 28 Feb 2025 01:22:42 +0000 Subject: [PATCH] SWDEV-513197 - Unify getBuffer implementation - Use getBuffer/releaseBuffer in BlitManager - Cleanup XferBuffer as we use ManagedBuffer for both reads/writes Change-Id: I2661b85dd012763b17a38a743fec1b1d79125f67 [ROCm/clr commit: 37d606d193b7a05b3b27c9952c8482c678858a16] --- projects/clr/rocclr/device/rocm/rocblit.cpp | 71 ++++++++++------ projects/clr/rocclr/device/rocm/rocblit.hpp | 51 ++++++----- projects/clr/rocclr/device/rocm/rocdevice.cpp | 85 ------------------- projects/clr/rocclr/device/rocm/rocdevice.hpp | 49 +---------- .../clr/rocclr/device/rocm/rocsettings.cpp | 2 - .../clr/rocclr/device/rocm/rocsettings.hpp | 4 +- 6 files changed, 71 insertions(+), 191 deletions(-) diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 765790c378..a150bf260e 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -650,21 +650,20 @@ bool DmaBlitManager::hsaCopyStagedOrPinned(const_address hostSrc, address hostDs hsa_agent_t dstAgent = hostToDev ? dev().getBackendDevice() : dev().getCpuAgent(); bool firstTx = true; while(totalSize > 0) { - size_t outsize = totalSize; const_address hostmem = hostToDev ? hostSrc : hostDst; // Get Pinned Host Memory or Staging buffer based on copy size - BufferState buffer{0}; - getBuffer(static_cast(hostmem + copyOffset), outsize, - enablePin, firstTx, buffer); - size_t copysize = buffer.copySize_; - address stagingBuffer = buffer.buffer_; + BufferState outBuffer = {0}; + getBuffer(static_cast(hostmem + copyOffset), totalSize, + enablePin, firstTx, outBuffer); + size_t copysize = outBuffer.copySize_; + address stagingBuffer = outBuffer.buffer_; if (stagingBuffer == 0) { LogWarning("DmaBlitManager::hsaCopyStagedOrPinned Buffer creation failed!"); status = false; break; } if (hostToDev) { // H2D Path - if (buffer.pinnedMem_ == nullptr) { // Copy to Staging Buffer + if (outBuffer.pinnedMem_ == nullptr) { // Copy to Staging Buffer memcpy(stagingBuffer, hostSrc + copyOffset, copysize); } ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "HSA Async Copy staged H2D"); @@ -678,7 +677,8 @@ bool DmaBlitManager::hsaCopyStagedOrPinned(const_address hostSrc, address hostDs const_address src = static_cast(hostSrc) + copyOffset; status = rocrCopyBuffer(stagingBuffer, dstAgent, src , srcAgent, copysize, copyMetadata); if (status ) { - if (buffer.pinnedMem_ == nullptr) { // Blocking copy from Staging Buffer + // Wait for current signal of previous rocr copy if its not pinned mem + if (outBuffer.pinnedMem_ == nullptr) { gpu().Barriers().WaitCurrent(); memcpy(hostDst + copyOffset, stagingBuffer, copysize); } @@ -686,8 +686,8 @@ bool DmaBlitManager::hsaCopyStagedOrPinned(const_address hostSrc, address hostDs break; } } - // Release Pinned Memory back to pool - releaseBuffer(buffer); + // Release Pinned Memory back to pool if any + releaseBuffer(outBuffer); // Update Offset and Transfer Size copyOffset += copysize; totalSize -= copysize; @@ -1683,13 +1683,16 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, amd::Coord3D dstOrigin(0, 0, 0); size_t copySize = 0; size_t stagedCopyOffset = 0; - size_t maxStagedXferSize = dev().settings().stagedXferSize_; - Memory& xferBuf = dev().xferRead().acquire(); - address xferBufAddr = xferBuf.getDeviceMemory(); - constexpr bool kAttachSignal = true; + while (totalSize > 0) { - copySize = std::min(totalSize, maxStagedXferSize); + BufferState outBuffer = {0}; + constexpr bool kEnablePin = true; + constexpr bool kFirstTx = false; + getBuffer(static_cast(dstAddr + stagedCopyOffset), totalSize, + kEnablePin, kFirstTx, outBuffer); + copySize = outBuffer.copySize_; + address xferBufAddr = outBuffer.buffer_; srcAddr += stagedCopyOffset; ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Blit staging D2H copy stg buf=%p, src=%p, " "dstOrigin=0x%x, size=%zu", xferBufAddr, srcAddr, dstOrigin[0], copySize); @@ -1703,16 +1706,20 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, if (!result) { break; } - // Wait on current signal of previous blit copy - gpu().Barriers().WaitCurrent(); - ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "memcpy host dst=%p, stg buf=%p, size=%zu", - (void*)(dstAddr + stagedCopyOffset), xferBufAddr, copySize); - memcpy(dstAddr + stagedCopyOffset, xferBufAddr, copySize); + // Wait for current signal of previous blit copy if its not pinned mem + if (outBuffer.pinnedMem_ == nullptr) { + gpu().Barriers().WaitCurrent(); + ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "memcpy host dst=%p, stg buf=%p, size=%zu", + (void*)(dstAddr + stagedCopyOffset), xferBufAddr, copySize); + memcpy(dstAddr + stagedCopyOffset, xferBufAddr, copySize); + } totalSize -= copySize; stagedCopyOffset += copySize; + // Release Pinned Memory back to pool + releaseBuffer(outBuffer); } - dev().xferRead().release(gpu(), xferBuf); + } } @@ -1814,16 +1821,24 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo amd::Coord3D srcOrigin(0, 0, 0); size_t copySize = 0; size_t stagedCopyOffset = 0; - size_t maxStagedXferSize = dev().settings().stagedXferSize_; while (totalSize > 0) { - copySize = std::min(totalSize, maxStagedXferSize); + BufferState outBuffer = {0}; + // Disable pinned writes + constexpr bool kEnablePin = false; + constexpr bool kFirstTx = false; + // Do not enable pinning for uploads. Always use staging buffer + getBuffer(static_cast(srcAddr + stagedCopyOffset), totalSize, + kEnablePin, kFirstTx, outBuffer); // Get an address from managed staging buffer - address stagingBuffer = gpu().Staging().Acquire(std::min(copySize, maxStagedXferSize)); + address stagingBuffer = outBuffer.buffer_; + copySize = outBuffer.copySize_; dstAddr += stagedCopyOffset; - ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "memcpy stg buf=%p, host src=%p, size=%zu", - stagingBuffer, (void*)(srcAddr + stagedCopyOffset), copySize); - memcpy(stagingBuffer, srcAddr + stagedCopyOffset, copySize); + if (outBuffer.pinnedMem_ == nullptr) { + ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "memcpy stg buf=%p, host src=%p, size=%zu", + stagingBuffer, (void*)(srcAddr + stagedCopyOffset), copySize); + memcpy(stagingBuffer, srcAddr + stagedCopyOffset, copySize); + } ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Blit staging H2D copy dst=%p, stg buf=%p, " "dstOrigin=0x%x, size=%zu", dstAddr, stagingBuffer, origin[0], copySize); // No cache flush is needed here as we use a staging buffer, and the acquire logic @@ -1836,6 +1851,8 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo } totalSize -= copySize; stagedCopyOffset += copySize; + // Release pinned memory if any + releaseBuffer(outBuffer); } } } diff --git a/projects/clr/rocclr/device/rocm/rocblit.hpp b/projects/clr/rocclr/device/rocm/rocblit.hpp index d2bad22d20..8a7f03f454 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.hpp +++ b/projects/clr/rocclr/device/rocm/rocblit.hpp @@ -208,7 +208,11 @@ class DmaBlitManager : public device::HostBlitManager { } protected: - + struct BufferState{ + address buffer_; //!< Staging Buffer or Pinned Host Mem Address + amd::Memory* pinnedMem_; //!< Pinned Memory + size_t copySize_; //!< last copy size + }; //! Synchronizes the blit operations if necessary inline void synchronize() const; @@ -236,6 +240,26 @@ class DmaBlitManager : public device::HostBlitManager { const_address src, hsa_agent_t& srcAgent, size_t size, amd::CopyMetadata& copyMetadata) const; + // Get Pinned Host Memory or Staging Buffer + void getBuffer(const_address hostMem, //!< Host Mem Address + size_t size, //!< Transfer Size + bool enablePin, //!< True when Pinning is enabled + bool first_tx, //!< True for first copy + BufferState &buffer //!< State of Buffer + ) const; + + // Release Pinned host memory + void releaseBuffer(BufferState &buff //!< True if last copy used Pinned resource + ) const; + + bool hsaCopyStagedOrPinned(const_address hostSrc, //!< Src buffer address + address hostDst, //!< Dst Buffer address + size_t size, //!< Size in bytes + bool hostToDev, //!< True for H2D copy + amd::CopyMetadata& copyMetadata, //!< copy MetaData + bool enPinning = false //!< True if pinning required + ) const; + const size_t PinXferSize; //!< Copy size for Pinned Copy const size_t MinSizeForPinnedXfer; //!< Mininum copy size for Pinned Copy const size_t StagingXferSize; //!< Copy size for Staging Buffer Copy @@ -251,31 +275,6 @@ class DmaBlitManager : public device::HostBlitManager { //! Disable operator= DmaBlitManager& operator=(const DmaBlitManager&); - - bool hsaCopyStagedOrPinned(const_address hostSrc, //!< Src buffer address - address hostDst, //!< Dst Buffer address - size_t size, //!< Size of copy data in bytes - bool hostToDev, //!< True for H2D copy - amd::CopyMetadata& copyMetadata, //!< copy MetaData - bool enPinning = false //!< True if pinning required - ) const; - struct BufferState{ - address buffer_; //!< Staging Buffer or Pinned Host Mem Address - amd::Memory* pinnedMem_; //!< Pinned Memory - size_t copySize_; //!< last copy size - }; - - // Get Pinned Host Memory or Staging Buffer - void getBuffer(const_address hostMem, //!< Host Mem Address - size_t size, //!< Transfer Size - bool enablePin, //!< True when Pinning is enabled - bool first_tx, //!< True for first copy - BufferState &buffer //!< State of Buffer - ) const; - - // Release Pinned host memory - void releaseBuffer(BufferState &buff //!< True if last copy used Pinned resource - ) const; }; //! Kernel Blit Manager diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index 17a6a223c9..6d44d2f080 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -180,7 +180,6 @@ Device::Device(hsa_agent_t bkendDevice) , gpuvm_segment_max_alloc_(0) , alloc_granularity_(0) , xferQueue_(nullptr) - , xferRead_(nullptr) , freeMem_(0) , vgpusAccess_(true) /* Virtual GPU List Ops Lock */ , hsa_exclusive_gpu_access_(false) @@ -289,9 +288,6 @@ Device::~Device() { } queuePool_.clear(); - // Destroy temporary buffers for read/write - delete xferRead_; - // Destroy transfer queue delete xferQueue_; @@ -391,76 +387,6 @@ hsa_status_t Device::loaderQueryHostAddress(const void* device, const void** hos : HSA_STATUS_ERROR; } -Device::XferBuffers::~XferBuffers() { - // Destroy temporary buffer for reads - for (const auto& buf : freeBuffers_) { - delete buf; - } - freeBuffers_.clear(); -} - -bool Device::XferBuffers::create() { - Memory* xferBuf = nullptr; - bool result = false; - - // Create a buffer object - xferBuf = new Buffer(dev(), bufSize_); - - // Try to allocate memory for the transfer buffer - if ((nullptr == xferBuf) || !xferBuf->create()) { - delete xferBuf; - xferBuf = nullptr; - LogError("Couldn't allocate a transfer buffer!"); - } else { - result = true; - freeBuffers_.push_back(xferBuf); - } - - return result; -} - -Memory& Device::XferBuffers::acquire() { - Memory* xferBuf = nullptr; - size_t listSize; - - // Lock the operations with the staged buffer list - amd::ScopedLock l(lock_); - listSize = freeBuffers_.size(); - - // If the list is empty, then attempt to allocate a staged buffer - if (listSize == 0) { - // Allocate memory - xferBuf = new Buffer(dev(), bufSize_); - - // Allocate memory for the transfer buffer - if ((nullptr == xferBuf) || !xferBuf->create()) { - delete xferBuf; - xferBuf = nullptr; - LogError("Couldn't allocate a transfer buffer!"); - } else { - ++acquiredCnt_; - } - } - - if (xferBuf == nullptr) { - xferBuf = *(freeBuffers_.begin()); - freeBuffers_.erase(freeBuffers_.begin()); - ++acquiredCnt_; - } - - return *xferBuf; -} - -void Device::XferBuffers::release(VirtualGPU& gpu, Memory& buffer) { - // Make sure buffer isn't busy on the current VirtualGPU, because - // the next aquire can come from different queue - // buffer.wait(gpu); - // Lock the operations with the staged buffer list - amd::ScopedLock l(lock_); - freeBuffers_.push_back(&buffer); - --acquiredCnt_; -} - // ================================================================================================ bool Device::init() { ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Initializing HSA stack."); @@ -817,17 +743,6 @@ bool Device::create() { // Use just 1 entry by default for the map cache mapCache_->push_back(nullptr); - if (settings().stagedXferSize_ != 0) { - // Initialize staged read buffers - if (settings().stagedXferRead_) { - xferRead_ = new XferBuffers(*this, amd::alignUp(settings().stagedXferSize_, 4 * Ki)); - if ((xferRead_ == nullptr) || !xferRead_->create()) { - LogError("Couldn't allocate transfer buffer objects for write"); - return false; - } - } - } - // Create signal for HMM prefetch operation on device if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &prefetch_signal_)) { return false; diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp index 1b53e6227a..b1307486e8 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp @@ -329,50 +329,7 @@ struct AgentInfo { //! A HSA device ordinal (physical HSA device) class Device : public NullDevice { public: - //! Transfer buffers - class XferBuffers : public amd::HeapObject { - public: - static const size_t MaxXferBufListSize = 8; - - //! Default constructor - XferBuffers(const Device& device, size_t bufSize) - : bufSize_(bufSize), acquiredCnt_(0), gpuDevice_(device) {} - - //! Default destructor - ~XferBuffers(); - - //! Creates the xfer buffers object - bool create(); - - //! Acquires an instance of the transfer buffers - Memory& acquire(); - - //! Releases transfer buffer - void release(VirtualGPU& gpu, //!< Virual GPU object used with the buffer - Memory& buffer //!< Transfer buffer for release - ); - - //! Returns the buffer's size for transfer - size_t bufSize() const { return bufSize_; } - - private: - //! Disable copy constructor - XferBuffers(const XferBuffers&); - - //! Disable assignment operator - XferBuffers& operator=(const XferBuffers&); - - //! Get device object - const Device& dev() const { return gpuDevice_; } - - size_t bufSize_; //!< Staged buffer size - std::list freeBuffers_; //!< The list of free buffers - std::atomic_uint acquiredCnt_; //!< The total number of acquired buffers - amd::Monitor lock_; //!< Stgaed buffer acquire/release lock - const Device& gpuDevice_; //!< GPU device object - }; - - //! Initialise the whole HSA device subsystem (CAL init, device enumeration, etc). + //! Initialise the whole HSA device subsystem (init, device enumeration, etc). static bool init(); static void tearDown(); @@ -517,9 +474,6 @@ class Device : public NullDevice { //! Adds a map target to the cache bool addMapTarget(amd::Memory* memory) const; - //! Returns transfer buffer object - XferBuffers& xferRead() const { return *xferRead_; } - //! Returns a ROC memory object from AMD memory object roc::Memory* getRocMemory(amd::Memory* mem //!< Pointer to AMD memory object ) const; @@ -648,7 +602,6 @@ class Device : public NullDevice { static constexpr bool offlineDevice_ = false; VirtualGPU* xferQueue_; //!< Transfer queue, created on demand - XferBuffers* xferRead_; //!< Transfer buffers read 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 diff --git a/projects/clr/rocclr/device/rocm/rocsettings.cpp b/projects/clr/rocclr/device/rocm/rocsettings.cpp index fe3fab3ed2..c68065f9f5 100644 --- a/projects/clr/rocclr/device/rocm/rocsettings.cpp +++ b/projects/clr/rocclr/device/rocm/rocsettings.cpp @@ -53,8 +53,6 @@ Settings::Settings() { // Disable image DMA by default (ROCM runtime doesn't support it) imageDMA_ = false; - stagedXferRead_ = true; - stagedXferWrite_ = true; stagedXferSize_ = flagIsDefault(GPU_STAGING_BUFFER_SIZE) ? 1 * Mi : GPU_STAGING_BUFFER_SIZE * Mi; diff --git a/projects/clr/rocclr/device/rocm/rocsettings.hpp b/projects/clr/rocclr/device/rocm/rocsettings.hpp index b15b8cbe64..b4da5df0d5 100644 --- a/projects/clr/rocclr/device/rocm/rocsettings.hpp +++ b/projects/clr/rocclr/device/rocm/rocsettings.hpp @@ -45,14 +45,12 @@ class Settings : public device::Settings { uint enableLocalMemory_ : 1; //!< Enable GPUVM memory uint enableNCMode_ : 1; //!< Enable Non Coherent mode for system memory uint imageDMA_ : 1; //!< Enable direct image DMA transfers - uint stagedXferRead_ : 1; //!< Uses a staged buffer read - uint stagedXferWrite_ : 1; //!< Uses a staged buffer write uint imageBufferWar_ : 1; //!< Image buffer workaround for Gfx10 uint cpu_wait_for_signal_ : 1; //!< Wait for HSA signal on CPU uint system_scope_signal_ : 1; //!< HSA signal is visibile to the entire system uint fgs_kernel_arg_ : 1; //!< Use fine grain kernel arg segment uint barrier_value_packet_ : 1; //!< Barrier value packet functionality - uint reserved_ : 21; + uint reserved_ : 23; }; uint value_; };