diff --git a/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.drawio b/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.drawio index 3d548182d7..2ea9376cf3 100644 --- a/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.drawio +++ b/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.drawio @@ -1,21 +1,21 @@ - + - + - + - - + + - + @@ -92,10 +92,10 @@ - + - + @@ -104,7 +104,7 @@ - + @@ -113,10 +113,10 @@ - + - + @@ -125,10 +125,10 @@ - + - + @@ -137,7 +137,7 @@ - + @@ -146,10 +146,10 @@ - + - + @@ -157,22 +157,16 @@ - - - - - - - - + + - - + + @@ -182,13 +176,13 @@ - + - + - + @@ -197,7 +191,7 @@ - + @@ -206,13 +200,13 @@ - + - + - + @@ -221,7 +215,7 @@ - + @@ -230,10 +224,10 @@ - + - + @@ -241,23 +235,17 @@ - - - - - - - - + + - + - + @@ -265,32 +253,20 @@ - - - - - - - - + + - - - - - - - + - - + + - - + + diff --git a/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg b/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg index 43d146db2d..fe52799858 100644 --- a/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg +++ b/docs/data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg @@ -1,2 +1,2 @@ -
time
time
default stream
default stream
H2D
data1
H2D...
H2D
data2
H2D...
kernel
data1
kernel...
kernel
data2
kernel...
D2H
data1
D2H...
D2H
data2
D2H...
H2D
data2
H2D...
kernel
data2
kernel...
stream2
stream2
D2H
data2
D2H...
H2D
data1
H2D...
kernel
data1
kernel...
stream1
stream1
D2H
data1
D2H...
default stream
default stream
Seqeuntial calls:
Seqeuntial calls:
Asynchronous calls:
Asynchronous calls:
Asynchronous calls with hipEvent:
Asynchronous calls with hipEvent: -
H2D
data2
H2D...
kernel
data2
kernel...
stream2
stream2
H2D
data1
H2D...
kernel
data1
kernel...
stream1
stream1
D2H
data2
D2H...
default stream
default stream
event
event
D2H
data1
D2H...
create events
create events
Text is not SVG - cannot display
\ No newline at end of file +
time
time
default stream
default stream
H2D
data1
H2D...
H2D
data2
H2D...
kernel
data1
kernel...
kernel
data2
kernel...
D2H
data1
D2H...
D2H
data2
D2H...
H2D
data2
H2D...
kernel
data2
kernel...
stream2
stream2
D2H
data2
D2H...
H2D
data1
H2D...
kernel
data1
kernel...
stream1
stream1
D2H
data1
D2H...
Seqeuntial calls:
Seqeuntial calls:
Asynchronous calls:
Asynchronous calls:
Asynchronous calls with hipEvent:
Asynchronous calls with hipEvent: +
H2D
data2
H2D...
kernel
data2
kernel...
stream2
stream2
H2D
data1
H2D...
kernel
data1
kernel...
stream1
stream1
D2H
data2
D2H...
event
event
D2H
data1
D2H...
eventA
eventB
eventA...
Text is not SVG - cannot display
\ No newline at end of file diff --git a/docs/how-to/hip_runtime_api/asynchronous.rst b/docs/how-to/hip_runtime_api/asynchronous.rst index 2015d31c39..588553ab94 100644 --- a/docs/how-to/hip_runtime_api/asynchronous.rst +++ b/docs/how-to/hip_runtime_api/asynchronous.rst @@ -442,7 +442,7 @@ The example codes // Record event after the GPU kernel in Stream 1 HIP_CHECK(hipEventRecord(event, streamA)); - // Stream 2: Wait for Event 1 before starting Kernel 2 + // Stream 2: Wait for event before starting Kernel 2 HIP_CHECK(hipStreamWaitEvent(streamB, event, 0)); // Stream 2: Kernel 2 @@ -453,13 +453,13 @@ The example codes // Stream 2: Device to Host 2 (after Kernel 2) HIP_CHECK(hipMemcpyAsync(vectorB.data(), d_dataB, arraySize * sizeof(*vectorB.data()), hipMemcpyDeviceToHost, streamB)); - } - // Wait for all operations in both streams to complete - HIP_CHECK(hipEventRecord(eventA, streamA)); - HIP_CHECK(hipEventRecord(eventB, streamB)); - HIP_CHECK(hipStreamWaitEvent(streamA, eventA, 0)); - HIP_CHECK(hipStreamWaitEvent(streamB, eventB, 0)); + // Wait for all operations in both streams to complete + HIP_CHECK(hipEventRecord(eventA, streamA)); + HIP_CHECK(hipEventRecord(eventB, streamB)); + HIP_CHECK(hipStreamWaitEvent(streamA, eventA, 0)); + HIP_CHECK(hipStreamWaitEvent(streamB, eventB, 0)); + } // Cleanup HIP_CHECK(hipEventDestroy(event));