SWDEV-301667 - Missing marker for hipExtLaunch* optimization

- Use isFenceDirty() check in hipStreamQuery.
In case if there is a stopEvent for hipExtLaunch* API we avoid system
scope if there are consecutive hipExtLaunch* and replace them with agent
scope, but then if one uses hipStreamQuery it should detect we need a
barrier with system scope.

Change-Id: I539810f659d4e1c49104e3a0da6ff2803c6b3dee


[ROCm/clr commit: 4b3c4709ff]
Este commit está contenido en:
Saleel Kudchadker
2024-01-11 07:03:40 +00:00
padre 898e647f84
commit 9864d6e56c
+9
Ver fichero
@@ -592,6 +592,14 @@ hipError_t hipStreamQuery_common(hipStream_t stream) {
bool wait = (stream == nullptr) ? true : false;
hip::Stream* hip_stream = hip::getStream(stream, wait);
if (hip_stream->vdev()->isFenceDirty()) {
amd::Command* command = new amd::Marker(*hip_stream, kMarkerDisableFlush);
if (command != nullptr) {
command->enqueue();
command->release();
}
}
amd::Command* command = hip_stream->getLastQueuedCommand(true);
if (command == nullptr) {
// Nothing was submitted to the queue
@@ -602,6 +610,7 @@ hipError_t hipStreamQuery_common(hipStream_t stream) {
if (command->type() != 0) {
event.notifyCmdQueue();
}
// Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status
bool ready = command->queue()->device().IsHwEventReady(event);
if (!ready) {