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

SWDEV-440866 - [hip-roclr] Adds support to batch memory operations APIs

Browse Source 
Change-Id: I449ffca44bbb04d13348d112e896d603c70fd485
This commit is contained in:
Sourabh Betigeri
2024-03-26 10:22:35 -07:00
committed by Sourabh Betigeri
parent c47f9dda58
commit bd5d8e9baf
18 changed files with 217 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/device/blit.hpp
+6
View File
@@ -233,6 +233,12 @@ class BlitManager : public amd::HeapObject {
uint64_t mask
) const = 0;
//! Stream batch memory operation
virtual bool batchMemOps(const void* paramArray,
size_t paramSize,
uint32_t count
) const = 0;
//! Enables synchronization on blit operations
void enableSynchronization() { syncOperation_ = true; }
rocclr/device/blitcl.cpp
+6 -1
View File
@@ -43,6 +43,8 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
extern void __amd_streamOpsWait(__global uint*, __global ulong*, ulong, ulong, ulong);
extern void __amd_batchMemOp(__global void*, uint count);
extern void __ockl_dm_init_v1(ulong, ulong, uint, uint);
extern void __ockl_gws_init(uint nwm1, uint rid);
@@ -162,6 +164,10 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
ulong4 srcRect, ulong4 dstRect, ulong4 size) {
__amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size);
}
__kernel void __amd_rocclr_batchMemOp(__global void* params, uint count) {
__amd_batchMemOp(params, count);
}
);
const char* HipExtraSourceCode = BLIT_KERNELS(
@@ -254,7 +260,6 @@ const char* BlitImageSourceCode = BLIT_KERNELS(
__amd_copyImageToBuffer(src, dstUInt, dstUShort, dstUChar, srcOrigin, dstOrigin, size, format,
pitch);
}
);
} // namespace amd::device
rocclr/device/device.hpp
+4
View File
@@ -92,6 +92,7 @@ class SvmMapMemoryCommand;
class SvmUnmapMemoryCommand;
class SvmPrefetchAsyncCommand;
class StreamOperationCommand;
class BatchMemoryOperationCommand;
class VirtualMapCommand;
class ExternalSemaphoreCmd;
class Isa;
@@ -1308,6 +1309,9 @@ class VirtualDevice : public amd::HeapObject {
ShouldNotReachHere();
}
virtual void submitStreamOperation(amd::StreamOperationCommand& cmd) { ShouldNotReachHere(); }
virtual void submitBatchMemoryOperation(amd::BatchMemoryOperationCommand& cmd) {
ShouldNotReachHere();
}
virtual void submitVirtualMap(amd::VirtualMapCommand& cmd) { ShouldNotReachHere(); }
virtual address allocKernelArguments(size_t size, size_t alignment) { return nullptr; }
rocclr/device/pal/palblit.hpp
+7
View File
@@ -475,6 +475,13 @@ class KernelBlitManager : public DmaBlitManager {
uint number_of_initial_blocks
) const;
//! Batch memory ops- Submits batch of streamWaits and streamWrite operations.
virtual bool batchMemOps(const void* paramArray,
size_t paramSize,
uint32_t count) const {
assert(!"Unimplemented");
return false;
}
private:
static constexpr size_t MaxXferBuffers = 2;
static constexpr uint TransferSplitSize = 3;
rocclr/device/rocm/rocblit.cpp
+32
View File
@@ -2583,6 +2583,38 @@ bool KernelBlitManager::streamOpsWait(device::Memory& memory, uint64_t value, si
return result;
}
// ================================================================================================
bool KernelBlitManager::batchMemOps(const void* paramArray, size_t paramSize,
uint32_t count) const {
amd::ScopedLock k(lockXferOps_);
bool result = false;
uint blitType = BatchMemOp;
size_t dim = 1;
size_t globalWorkOffset[1] = { 0 };
size_t globalWorkSize[1] = { count };
size_t localWorkSize[1] = { 1 };
// Get constant buffer and copy the array of parameters
constexpr bool kDirectVa = true;
auto constBuf = gpu().allocKernArg((count * paramSize), kCBAlignment);
memcpy(constBuf, paramArray, (count * paramSize));
setArgument(kernels_[blitType], 0, sizeof(cl_mem), constBuf, 0, nullptr, kDirectVa);
setArgument(kernels_[blitType], 1, sizeof(cl_mem), &count);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(dim, globalWorkOffset, globalWorkSize, localWorkSize);
// Execute the blit
address parameters = captureArguments(kernels_[blitType]);
result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr);
releaseArguments(parameters);
synchronize();
return result;
}
// ================================================================================================
bool KernelBlitManager::initHeap(device::Memory* heap_to_initialize, device::Memory* initial_blocks,
uint heap_size, uint number_of_initial_blocks) const {
rocclr/device/rocm/rocblit.hpp
+7 -3
View File
@@ -292,6 +292,7 @@ class KernelBlitManager : public DmaBlitManager {
BlitCopyBufferRectAligned,
StreamOpsWrite,
StreamOpsWait,
BatchMemOp,
Scheduler,
GwsInit,
InitHeap,
@@ -519,6 +520,9 @@ class KernelBlitManager : public DmaBlitManager {
uint64_t mask
) const;
//! Batch memory ops- Submits batch of streamWaits and streamWrite operations.
virtual bool batchMemOps(const void* paramArray, size_t paramSize, uint32_t count) const;
virtual amd::Monitor* lockXfer() const { return &lockXferOps_; }
virtual bool initHeap(device::Memory* heap_to_initialize,
@@ -599,9 +603,9 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = {
"__amd_rocclr_fillBufferAligned", "__amd_rocclr_fillBufferAligned2D", "__amd_rocclr_copyBuffer",
"__amd_rocclr_copyBufferAligned", "__amd_rocclr_copyBufferRect",
"__amd_rocclr_copyBufferRectAligned", "__amd_rocclr_streamOpsWrite", "__amd_rocclr_streamOpsWait",
"__amd_rocclr_scheduler", "__amd_rocclr_gwsInit", "__amd_rocclr_initHeap",
"__amd_rocclr_fillImage", "__amd_rocclr_copyImage", "__amd_rocclr_copyImage1DA",
"__amd_rocclr_copyImageToBuffer", "__amd_rocclr_copyBufferToImage"
"__amd_rocclr_batchMemOp", "__amd_rocclr_scheduler", "__amd_rocclr_gwsInit",
"__amd_rocclr_initHeap", "__amd_rocclr_fillImage", "__amd_rocclr_copyImage",
"__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", "__amd_rocclr_copyBufferToImage"
};
inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index,
rocclr/device/rocm/rocvirtual.cpp
+14 -4
View File
@@ -2715,8 +2715,7 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
else {
// mask is applied on value before performing
// the comparision defined by 'condition'
bool result = static_cast<KernelBlitManager&>(blitMgr()).streamOpsWait(*memory, value, offset,
sizeBytes, flags, mask);
bool result = blitMgr().streamOpsWait(*memory, value, offset, sizeBytes, flags, mask);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Waiting for value: 0x%lx."
" Flags: 0x%lx mask: 0x%lx", value, flags, mask);
if (!result) {
@@ -2729,8 +2728,7 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
// Ensure memory ordering preceding the write
dispatchBarrierPacket(kBarrierPacketReleaseHeader);
bool result = static_cast<KernelBlitManager&>(blitMgr()).streamOpsWrite(*memory, value,
offset, sizeBytes);
bool result = blitMgr().streamOpsWrite(*memory, value, offset, sizeBytes);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Writing value: 0x%lx", value);
if (!result) {
LogError("submitStreamOperation: Write failed!");
@@ -2741,6 +2739,18 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitBatchMemoryOperation(amd::BatchMemoryOperationCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
profilingBegin(cmd);
bool result = blitMgr().batchMemOps(cmd.getParamPtr(), cmd.paramSize(), cmd.count());
if (!result) {
LogError("submitBatchMemoryOperation failed!");
}
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
rocclr/device/rocm/rocvirtual.hpp
+1
View File
@@ -353,6 +353,7 @@ class VirtualGPU : public device::VirtualDevice {
void flush(amd::Command* list = nullptr, bool wait = false);
void submitFillMemory(amd::FillMemoryCommand& cmd);
void submitStreamOperation(amd::StreamOperationCommand& cmd);
void submitBatchMemoryOperation(amd::BatchMemoryOperationCommand& cmd);
void submitVirtualMap(amd::VirtualMapCommand& cmd);
void submitMigrateMemObjects(amd::MigrateMemObjectsCommand& cmd);