diff --git a/projects/clr/hipamd/src/hip_graph_helper.hpp b/projects/clr/hipamd/src/hip_graph_helper.hpp index c146621185..7fdb378041 100644 --- a/projects/clr/hipamd/src/hip_graph_helper.hpp +++ b/projects/clr/hipamd/src/hip_graph_helper.hpp @@ -121,4 +121,6 @@ hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord size_t &start); hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams); + +hip::MemcpyType ihipGetMemcpyType(const void* src, void* dst, hipMemcpyKind kind); } // namespace hip \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 12fb04e29d..c1dd6fd8a2 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -1576,6 +1576,13 @@ class GraphMemcpyNode1D : public GraphMemcpyNode { WorkerThreadLock_.lock(); } status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *stream); + hip::MemcpyType type = ihipGetMemcpyType(src_, dst_, kind_); + if (type == hipCopyBuffer) { + amd::CopyMemoryCommand* cpycmd = reinterpret_cast(command); + amd::CopyMetadata copyMetadata = cpycmd->copyMetadata(); + copyMetadata.copyEnginePreference_ = amd::CopyMetadata::CopyEnginePreference::BLIT; + cpycmd->SetCopyMetadata(copyMetadata); + } if (!AMD_DIRECT_DISPATCH) { WorkerThreadLock_.unlock(); } @@ -1720,8 +1727,9 @@ class GraphMemcpyNode1D : public GraphMemcpyNode { virtual bool GraphCaptureEnabled() override { bool isGraphCapture = false; if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) { - switch (kind_) { - case hipMemcpyDeviceToDevice: + hip::MemcpyType type = ihipGetMemcpyType(src_, dst_, kind_); + switch (type) { + case hipCopyBuffer: isGraphCapture = true; break; default: diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index cec76d9bd2..3b8cedad76 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -69,6 +69,21 @@ typedef struct hipArray { }hipArray; namespace hip { +enum MemcpyType { + /// Memcpy from host to host + hipHostToHost, + /// Memcpy from host to device + hipWriteBuffer, + /// Memcpy from device to host + hipReadBuffer, + /// Memcpy from device A to device A + /// Memcpy from pinned host buffer to device/device to pinned host buffer + hipCopyBuffer, + /// Memcpy from device A to device A, user forced to SDMA + hipCopyBufferSDMA, + /// Memcpy from device A to device B + hipCopyBufferP2P, +}; struct Graph; struct GraphNode; struct GraphExec; diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index fc30c0f0db..31326ba766 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -463,6 +463,35 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, return hipSuccess; } +hip::MemcpyType ihipGetMemcpyType(const void* src, void* dst, hipMemcpyKind kind) { + size_t sOffset = 0; + amd::Memory* srcMemory = getMemoryObject(src, sOffset); + size_t dOffset = 0; + amd::Memory* dstMemory = getMemoryObject(dst, dOffset); + hip::MemcpyType type; + if(srcMemory == nullptr && dstMemory == nullptr) { + type = hipHostToHost; + } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { + type = hipWriteBuffer; + } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { + type = hipReadBuffer; + } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { + // Check if the queue device doesn't match the device on any memory object. + // And any of them are not host allocation. + // Hence it's a P2P transfer, because the app has requested access to another GPU + if ((srcMemory->GetDeviceById() != dstMemory->GetDeviceById()) && + ((srcMemory->getContext().devices().size() == 1) && + (dstMemory->getContext().devices().size() == 1))) { + type = hipCopyBufferP2P; + } else if( kind == hipMemcpyDeviceToDeviceNoCU) { + type = hipCopyBufferSDMA; + } else { + type = hipCopyBuffer; + } + } + return type; +} + // ================================================================================================ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hip::Stream& stream, bool isAsync) { @@ -473,37 +502,37 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, amd::Memory* dstMemory = getMemoryObject(dst, dOffset); amd::Device* queueDevice = &stream.device(); amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::NONE); - if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - hip::Stream* pStream = &stream; - if (queueDevice != dstMemory->GetDeviceById()) { - pStream = hip::getNullStream(dstMemory->GetDeviceById()->context()); - amd::Command* cmd = stream.getLastQueuedCommand(true); - if (cmd != nullptr) { - waitList.push_back(cmd); + hip::MemcpyType type = ihipGetMemcpyType(src, dst, kind); + hip::Stream* pStream = &stream; + switch (type) { + case hipWriteBuffer: + if (queueDevice != dstMemory->GetDeviceById()) { + pStream = hip::getNullStream(dstMemory->GetDeviceById()->context()); + amd::Command* cmd = stream.getLastQueuedCommand(true); + if (cmd != nullptr) { + waitList.push_back(cmd); + } } - } - command = new amd::WriteMemoryCommand(*pStream, CL_COMMAND_WRITE_BUFFER, waitList, - *dstMemory->asBuffer(), dOffset, sizeBytes, src, 0, 0, copyMetadata); - } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - hip::Stream* pStream = &stream; - if (queueDevice != srcMemory->GetDeviceById()) { - pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); - amd::Command* cmd = stream.getLastQueuedCommand(true); - if (cmd != nullptr) { - waitList.push_back(cmd); + command = new amd::WriteMemoryCommand(*pStream, CL_COMMAND_WRITE_BUFFER, waitList, + *dstMemory->asBuffer(), dOffset, sizeBytes, src, 0, 0, + copyMetadata); + break; + case hipReadBuffer: + if (queueDevice != srcMemory->GetDeviceById()) { + pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); + amd::Command* cmd = stream.getLastQueuedCommand(true); + if (cmd != nullptr) { + waitList.push_back(cmd); + } } - } - command = new amd::ReadMemoryCommand(*pStream, CL_COMMAND_READ_BUFFER, waitList, - *srcMemory->asBuffer(), sOffset, sizeBytes, dst, 0, 0, copyMetadata); - } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - // Check if the queue device doesn't match the device on any memory object. - // And any of them are not host allocation. - // Hence it's a P2P transfer, because the app has requested access to another GPU - if ((srcMemory->GetDeviceById() != dstMemory->GetDeviceById()) && - ((srcMemory->getContext().devices().size() == 1) && - (dstMemory->getContext().devices().size() == 1))) { + command = new amd::ReadMemoryCommand(*pStream, CL_COMMAND_READ_BUFFER, waitList, + *srcMemory->asBuffer(), sOffset, sizeBytes, dst, 0, 0, + copyMetadata); + break; + case hipCopyBufferP2P: command = new amd::CopyMemoryP2PCommand(stream, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes); + *srcMemory->asBuffer(), *dstMemory->asBuffer(), + sOffset, dOffset, sizeBytes); if (command == nullptr) { return hipErrorOutOfMemory; } @@ -513,10 +542,12 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, delete command; return hipErrorInvalidValue; } - } else { - hip::Stream* pStream = &stream; + break; + case hipCopyBufferSDMA: + copyMetadata.copyEnginePreference_ = amd::CopyMetadata::CopyEnginePreference::SDMA; + case hipCopyBuffer: if ((srcMemory->GetDeviceById() == dstMemory->GetDeviceById()) && - (queueDevice != srcMemory->GetDeviceById())) { + queueDevice != srcMemory->GetDeviceById()) { pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); amd::Command* cmd = stream.getLastQueuedCommand(true); if (cmd != nullptr) { @@ -531,7 +562,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, if (cmd != nullptr) { waitList.push_back(cmd); } - // Scenarios such as HtoD where src is pinned memory + // Scenarios such as HtoD where src is pinned memory } else if ((queueDevice != dstMemory->GetDeviceById()) && (srcMemory->getContext().devices().size() != 1)) { pStream = hip::getNullStream(dstMemory->GetDeviceById()->context()); @@ -541,14 +572,10 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, } } } - - copyMetadata.copyEnginePreference_ = (kind == hipMemcpyDeviceToDeviceNoCU) ? - amd::CopyMetadata::CopyEnginePreference::SDMA : - amd::CopyMetadata::CopyEnginePreference::NONE; command = new amd::CopyMemoryCommand(*pStream, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes, - copyMetadata); - } + *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, + dOffset, sizeBytes, copyMetadata); + break; } if (command == nullptr) { return hipErrorOutOfMemory; @@ -558,6 +585,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, } return hipSuccess; } + bool IsHtoHMemcpy(void* dst, const void* src) { size_t sOffset = 0; amd::Memory* srcMemory = getMemoryObject(src, sOffset); diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 84f659d44e..e05dd8f6f9 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -2263,12 +2263,12 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds #endif bool useShaderCopyPath = setup_.disableHwlCopyBuffer_ || - (sizeIn[0] <= dev().settings().sdmaCopyThreshold_) || - (!srcMemory.isHostMemDirectAccess() && - !dstMemory.isHostMemDirectAccess() && - !(p2p || asan) && !ipcShared && - !(copyMetadata.copyEnginePreference_ == - amd::CopyMetadata::CopyEnginePreference::SDMA)); + (sizeIn[0] <= dev().settings().sdmaCopyThreshold_) || + (!(p2p || asan || ipcShared) && + (!srcMemory.isHostMemDirectAccess() && !dstMemory.isHostMemDirectAccess() && + !(copyMetadata.copyEnginePreference_ == + amd::CopyMetadata::CopyEnginePreference::SDMA)) || + (copyMetadata.copyEnginePreference_ == amd::CopyMetadata::CopyEnginePreference::BLIT)); if (!useShaderCopyPath) { if (amd::IS_HIP) { diff --git a/projects/clr/rocclr/platform/command.hpp b/projects/clr/rocclr/platform/command.hpp index ef6e4b2d93..1001225a91 100644 --- a/projects/clr/rocclr/platform/command.hpp +++ b/projects/clr/rocclr/platform/command.hpp @@ -952,6 +952,8 @@ class CopyMemoryCommand : public TwoMemoryArgsCommand { const BufferRect& dstRect() const { return dstRect_; } //! Return the copy MetaData amd::CopyMetadata copyMetadata() const { return copyMetadata_; } + //! Updates copy MetaData + void SetCopyMetadata(amd::CopyMetadata copyMetadata) { copyMetadata_ = copyMetadata; } //! Updates the host memory to read from void setSource(Memory& srcMemory) { memory1_ = &srcMemory; } //! Updates the memory object to write to.