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

SWDEV-508435 - Use the stream of the src/dst image memory object in A2H and H2A commands

Browse Source 
Change-Id: I9b776a54760a4633d5f84cf7b467d2d3ba8cbdde
This commit is contained in:
Ioannis Assiouras
2025-02-06 14:59:57 +00:00
parent de83d7a6ae
commit a8edb8d467
1 changed files with 26 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_memory.cpp
+26 -2
View File
@@ -1973,8 +1973,20 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, amd::Image* dstImage,
}
command = copyCommand;
} else {
hip::Stream* pStream = stream;
amd::Device* queueDevice = &stream->device();
amd::Command::EventWaitList waitList;
if (queueDevice != dstImage->GetDeviceById()) {
pStream = hip::getNullStream(dstImage->GetDeviceById()->context());
amd::Command* cmd = stream->getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
}
amd::WriteMemoryCommand* writeMemCmd = new amd::WriteMemoryCommand(
*stream, CL_COMMAND_WRITE_IMAGE, amd::Command::EventWaitList{}, *dstImage, dstOrigin,
*pStream, CL_COMMAND_WRITE_IMAGE, waitList, *dstImage, dstOrigin,
copyRegion, static_cast<const char*>(srcHost) + start, srcRowPitch, srcSlicePitch,
copyMetadata);
if (writeMemCmd == nullptr) {
@@ -2010,8 +2022,20 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, void* dstHost, amd::Coo
}
command = copyCommand;
} else {
hip::Stream* pStream = stream;
amd::Device* queueDevice = &stream->device();
amd::Command::EventWaitList waitList;
if (queueDevice != srcImage->GetDeviceById()) {
pStream = hip::getNullStream(srcImage->GetDeviceById()->context());
amd::Command* cmd = stream->getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
}
amd::ReadMemoryCommand* readMemCmd = new amd::ReadMemoryCommand(
*stream, CL_COMMAND_READ_IMAGE, amd::Command::EventWaitList{}, *srcImage, srcOrigin,
*pStream, CL_COMMAND_READ_IMAGE, waitList, *srcImage, srcOrigin,
copyRegion, static_cast<char*>(dstHost) + start, dstRowPitch, dstSlicePitch,
copyMetadata);