AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-513197 - Unify getBuffer implementation

Browse Source 
- Use getBuffer/releaseBuffer in BlitManager
- Cleanup XferBuffer as we use ManagedBuffer for both reads/writes

Change-Id: I2661b85dd012763b17a38a743fec1b1d79125f67


[ROCm/clr commit: 37d606d193]
This commit is contained in:
Saleel Kudchadker
2025-02-28 01:22:42 +00:00
parent ef505c7cd8
commit d0a7ae02cf
6 changed files with 71 additions and 191 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/rocm/rocblit.cpp
+44 -27
View File
@@ -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<const_address>(hostmem + copyOffset), outsize,
enablePin, firstTx, buffer);
size_t copysize = buffer.copySize_;
address stagingBuffer = buffer.buffer_;
BufferState outBuffer = {0};
getBuffer(static_cast<const_address>(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<const_address>(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<const_address>(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<const_address>(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);
}
}
}
projects/clr/rocclr/device/rocm/rocblit.hpp
+25 -26
View File
@@ -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
projects/clr/rocclr/device/rocm/rocdevice.cpp
-85
View File
@@ -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;
projects/clr/rocclr/device/rocm/rocdevice.hpp
+1 -48
View File
@@ -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<Memory*> 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<size_t> 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
projects/clr/rocclr/device/rocm/rocsettings.cpp
-2
View File
@@ -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;
projects/clr/rocclr/device/rocm/rocsettings.hpp
+1 -3
View File
@@ -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_;
};