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: I5ac63a6626af8c2b4ac382c52dfe1aaf0b3716b8
This commit is contained in:
Sourabh Betigeri
2024-03-26 10:22:35 -07:00
committed by Ajay GunaShekar
parent 3ad8f1b811
commit 03dbcd8ca7
18 changed files with 220 additions and 20 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
@@ -2511,6 +2511,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
+13 -7
View File
@@ -285,6 +285,7 @@ class KernelBlitManager : public DmaBlitManager {
BlitCopyImage1DA,
BlitCopyImageToBuffer,
BlitCopyBufferToImage,
BatchMemOp,
BlitTotal
};
@@ -503,6 +504,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,
@@ -586,13 +590,15 @@ class KernelBlitManager : public DmaBlitManager {
};
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_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"};
inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index,
size_t size, const void* value, size_t offset,
rocclr/device/rocm/rocvirtual.cpp
+14 -4
View File
@@ -2719,8 +2719,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) {
@@ -2733,8 +2732,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!");
@@ -2745,6 +2743,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
@@ -368,6 +368,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);
rocclr/platform/command.hpp
+42 -1
View File
@@ -845,7 +845,7 @@ class FillMemoryCommand : public OneMemoryArgCommand {
class StreamOperationCommand : public OneMemoryArgCommand {
private:
uint64_t value_; // !< Value to Wait on or to Write.
uint64_t value_; // !< Value to Wait on or to Write.
uint64_t mask_; // !< Mask to be applied on signal value for Wait operation.
unsigned int flags_; // !< Flags defining the Wait condition.
size_t offset_; // !< Offset into memory for Write
@@ -888,6 +888,46 @@ class StreamOperationCommand : public OneMemoryArgCommand {
const size_t sizeBytes() const { return sizeBytes_; }
};
/*! \brief A batch memory operation command.
*
* \details Batch operations to synchronize the stream via memory operations
* Operations are either 32-bit stream wait or write.
* Wait: All the commands issued after stream wait are not executed
* until the wait condition is true.
* Write: Writes a 32 or 64 bit vaue to the memory using a GPU Blit.
* The operations are enqueued in the order they appear in the array.
*/
class BatchMemoryOperationCommand : public Command {
public:
BatchMemoryOperationCommand(HostQueue& queue, cl_command_type cmdType, uint32_t count,
uint32_t flags, EventWaitList& eventWaitList, const void* paramArray,
size_t paramSize)
: Command(queue, cmdType, eventWaitList),
count_(count),
paramArray_(paramArray),
flags_(flags),
paramSize_(paramSize) {
// Sanity check
assert(((cmdType == ROCCLR_COMMAND_BATCH_STREAM)) && "Invalid batch memory operation");
}
virtual void submit(device::VirtualDevice& device) { device.submitBatchMemoryOperation(*this); }
//! Returns the value
const uint64_t count() const { return count_; }
//! Return the pointer to the paramList
const void* getParamPtr() { return paramArray_; }
//! Return the size of a single mem op param in bytes
const size_t paramSize() const { return paramSize_; }
private:
uint32_t count_; // !< The number of operations in the array.
uint32_t flags_; // !< Reserved for future expansion. Must be 0.
const void* paramArray_; // !< Pointer to the array of individual operations
size_t paramSize_; // !< size in bytes of the param array passed
};
/*! \brief A generic copy memory command
*
* \details Used for both buffers and images. Backends are expected
@@ -1858,6 +1898,7 @@ union ComputeCommand {
CopyMemoryP2PCommand cmd25;
SvmPrefetchAsyncCommand cmd26;
VirtualMapCommand cmd27;
BatchMemoryOperationCommand cmd28;
ComputeCommand() {}
~ComputeCommand() {}
};
rocclr/platform/command_utils.hpp
+1
View File
@@ -24,6 +24,7 @@
// Dummy command types for Stream Wait and Write commands.
#define ROCCLR_COMMAND_STREAM_WAIT_VALUE 0x4501
#define ROCCLR_COMMAND_STREAM_WRITE_VALUE 0x4502
#define ROCCLR_COMMAND_BATCH_STREAM 0x4503
// Stream Wait Value Conidtions
#define ROCCLR_STREAM_WAIT_VALUE_GTE 0x0