HIP events synchronization implementation update #9328
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
wendyliang25
requested review from
chvamshi-xilinx,
sonals and
stsoe
as code owners
wendyliang25
force-pushed
the
hip-event
branch
from
20723cc to
752569a
Compare
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
wendyliang25
force-pushed
the
hip-event
branch
2 times, most recently
from
90b320b to
813a55d
Compare
|
clang-tidy review says "All clean, LGTM! 👍" |
Before this patch, event to record will be overwritten if hipEventRecord() is called again on the same event. However, this impelmentatino has a few issues: * it didn't erase the event from the stream which the event has been recorded before, instead, it tries to erase the event from the stream it is trying to record. * there is no lock to protect the event recording change. This the operation to change the event recording will impact two objects, the event itself and the stream. To simplify, this patch disable changing event recording if event is being recorded by a stream, and add locking to protect the event state change. Signed-off-by: Wendy Liang <wendy.liang@amd.com>
wendyliang25
force-pushed
the
hip-event
branch
from
813a55d to
9643953
Compare
…l start This patch update the hipStreamWaitEvent() implementation to chain the dummy event created for the stream to the actual event will be populated from other stream. Besides the actual event will be added as the dummy event's dependency, if there is wait event set before on the same stream, the previous wait event will also be added as dependency of the dummy event. When the actual event completes, it will try to submit the dummy event, the dummy event will only starts its chained commands if all of its dependencies got resolved. Signed-off-by: Wendy Liang <wendy.liang@amd.com>
wendyliang25
force-pushed
the
hip-event
branch
from
9643953 to
5df6c02
Compare
|
clang-tidy review says "All clean, LGTM! 👍" |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Problem solved by the commit
There are a couple findings when testing the HIP events synchronization APIs.
hipEventRecord()implementation, and it doesn't erase the event from the old hip stream, when trying to record the event with a new stream, instead, it erased the event from the stream it is trying to record.hipEventSynchronize(), when the event dependencies are complete, it only starts the commands if they are on the same stream as where the event is recorded. It the commands depends on the event is on a different stream, it will not start.hipStreamWaitEvent(), a dummy event is created as the waiting point on the stream, but only the event to wait is added as dependencies to this waiting point, if stream has another wait point before, it is it added to this wait point's dependencies. e.g.stream1: wait_on_ev1 -> do_cmd1 - wait_on_ev2 -> do_cmd2.do_cmd2should not start until ev1 and ev2 completes, however, today, the implementation waits on ev1, ev2 and cmd1 to complete before doingdo_cmd2.This patch set are to
Bug / issue (if any) fixed, which PR introduced the bug, how it was discovered
How problem was solved, alternative solutions (if any) and why they were rejected
Risks (if any) associated the changes in the commit
What has been tested and how, request additional testing if necessary
Tests to sync launches on the same stream and on different streams with hip events: https://gitenterprise.xilinx.com/XRT/testcases-v2/pull/6
Documentation impact (if any)