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

SWDEV-272496 - Delay enqueue until callback setup

Browse Source 
With direct dispatch enqueue occurs before callback update and
it can't be tracked in the device backend

Change-Id: Ie8793e3ddb68cc5bb36348f7a8dcdbdc87a2487c


[ROCm/hip commit: ffc89dff8a]
This commit is contained in:
German Andryeyev
2021-03-03 13:14:28 -05:00
parent 5fb266cfb9
commit 94afcd643c
1 changed files with 5 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/rocclr/hip_stream.cpp
+5 -2
View File
@@ -368,6 +368,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
} }
amd::HostQueue* hostQueue = hip::getQueue(stream); amd::HostQueue* hostQueue = hip::getQueue(stream);
amd::Command* command = hostQueue->getLastQueuedCommand(true); amd::Command* command = hostQueue->getLastQueuedCommand(true);
bool cmd_created = false;
if (command == nullptr || if (command == nullptr ||
// Add a marker if the last command is a notification, // Add a marker if the last command is a notification,
// since with multithreading AddCallback may occur before the previous // since with multithreading AddCallback may occur before the previous
@@ -375,7 +376,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
(command->type() == 0)) { (command->type() == 0)) {
amd::Command::EventWaitList eventWaitList; amd::Command::EventWaitList eventWaitList;
command = new amd::Marker(*hostQueue, kMarkerDisableFlush, eventWaitList); command = new amd::Marker(*hostQueue, kMarkerDisableFlush, eventWaitList);
command->enqueue(); cmd_created = true;
} }
amd::Event& event = command->event(); amd::Event& event = command->event();
StreamCallback* cbo = new StreamCallback(stream, callback, userData, command); StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);
@@ -384,7 +385,9 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
command->release(); command->release();
return hipErrorInvalidHandle; return hipErrorInvalidHandle;
} }
if (cmd_created) {
command->enqueue();
}
event.notifyCmdQueue(); event.notifyCmdQueue();
HIP_RETURN(hipSuccess); HIP_RETURN(hipSuccess);