From dbc47446b3d08c365fc0c7426ee1ecb6c2c78b46 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= Date: Mon, 2 Dec 2024 17:27:30 +0100 Subject: [PATCH] WIP WIP WIP --- .../sequential_async_event.drawio | 602 +++++++++--------- .../asynchronous/sequential_async_event.svg | 6 +- docs/how-to/hip_runtime_api/asynchronous.rst | 131 ++-- docs/how-to/performance_guidelines.rst | 2 + 4 files changed, 389 insertions(+), 352 deletions(-) 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 b2b30b858e..fee6cab1ed 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,301 +1,301 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - \ No newline at end of file + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + 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 0b62bbda0e..54b48bcd9a 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,4 @@ -
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
data1
D2H...
default stream
default stream
event1
event1
event2
event2
D2H
data2
D2H...
create events
create eventsText 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...
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
event1
event1
event2
event2
D2H
data1
D2H...
create events
create eventsText 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 bc6825ad4e..8596bf3d1c 100644 --- a/docs/how-to/hip_runtime_api/asynchronous.rst +++ b/docs/how-to/hip_runtime_api/asynchronous.rst @@ -21,30 +21,47 @@ data allocation/freeing all happen in the context of device streams. Streams are FIFO buffers of commands to execute in order on a given device. Commands which enqueue tasks on a stream all return promptly and the command is -executed asynchronously. Multiple streams may point to the same device and those -streams may be fed from multiple concurrent host-side threads. Execution on -multiple streams may be concurrent but isn't required to be. +executed asynchronously. Multiple streams may point to the same device and +those streams may be fed from multiple concurrent host-side threads. Execution +on multiple streams may be concurrent but isn't required to be. Managing streams ------------------------------------------------------------------------------- -Streams enable the overlap of computation and data transfer, ensuring continuous -GPU activity. To create a stream, the :cpp:func:`hipStreamCreate`, -:cpp:func:`hipStreamCreateWithFlags` and :cpp:func:`hipStreamCreateWithPriority` -functions are used, returning a handle to the newly created stream. Assigning -different operations to different streams allows multiple tasks to run -simultaneously, improving overall performance. +Streams enable the overlap of computation and data transfer, ensuring +continuous GPU activity. -The :cpp:func:`hipStreamSynchronize` function in the HIP API is used to block -the calling host thread until all previously submitted tasks in a specified HIP -stream have completed. It ensures that all operations in the given stream, such -as kernel executions or memory transfers, are finished before the host thread -proceeds. +To create a stream, the following functions are used, each returning a handle +to the newly created stream: + +- :cpp:func:`hipStreamCreate`: Creates a stream with default settings. +- :cpp:func:`hipStreamCreateWithFlags`: Allows creating a stream, with specific + flags, listed below, enabling more control over stream behavior: + + - ``hipStreamDefault``: creates a default stream suitable for most + operations. It ensures that the stream is not non-blocking. + - ``hipStreamNonBlocking``: creates a non-blocking stream, allowing + concurrent execution of operations. It ensures that tasks can run + simultaneously without waiting for each other to complete, thus improving + overall performance. + +- :cpp:func:`hipStreamCreateWithPriority``: Allows creating a stream with a + specified priority, enabling prioritization of certain tasks. + +Assigning different operations to different streams allows multiple tasks to +run simultaneously, improving overall +:ref:`performance`. + +The :cpp:func:`hipStreamSynchronize` function is used to block the calling host +thread until all previously submitted tasks in a specified HIP stream have +completed. It ensures that all operations in the given stream, such as kernel +executions or memory transfers, are finished before the host thread proceeds. .. note:: If the :cpp:func:`hipStreamSynchronize` function input stream is 0 (or the - default stream), it waits for all operations in the default stream to complete. + default stream), it waits for all operations in the default stream to + complete. Concurrent execution between host and device ------------------------------------------------------------------------------- @@ -61,17 +78,26 @@ Concurrent kernel execution ------------------------------------------------------------------------------- Concurrent execution of multiple kernels on the GPU allows different kernels to -run simultaneously to maximize GPU resource usage. Managing dependencies between -kernels is crucial for ensuring correct execution order. This can be achieved -using :cpp:func:`hipStreamWaitEvent`, which allows a kernel to wait for a -specific event before starting execution. +run simultaneously to maximize GPU resource usage. Managing dependencies +between kernels is crucial for ensuring correct execution order. This can be +achieved using :cpp:func:`hipStreamWaitEvent`, which allows a kernel to wait +for a specific event before starting execution. -Independent kernels can only run concurrently, if there are enough registers and -share memories for the kernels. To reach concurrent kernel executions, the +Independent kernels can only run concurrently, if there are enough registers +and share memories for the kernels. To reach concurrent kernel executions, the developer may have to reduce the block size of the kernels. The kernel runtimes -can be misleading at concurrent kernel runs, that's why during optimization it's -better to check the trace files, to see if a kernel is blocking another kernel, -while they are running parallel. +can be misleading at concurrent kernel runs, that's why during optimization +it's better to check the trace files, to see if a kernel is blocking another +kernel, while they are running parallel. + +When running kernels in parallel, the execution time can increase due to +contention for shared resources. This is because multiple kernels may attempt +to access the same GPU resources simultaneously, leading to delays. + +Asynchronous kernel execution is beneficial only under specific conditions It +is most effective when the kernels do not fully utilize the GPU's resources. In +such cases, overlapping kernel execution can improve overall throughput and +efficiency by keeping the GPU busy without exceeding its capacity. Overlap of data transfer and kernel execution =============================================================================== @@ -80,6 +106,11 @@ One of the primary benefits of asynchronous operations is the ability to overlap data transfer with kernel execution, leading to better resource utilization and improved performance. +Asynchronous execution is particularly advantageous in iterative processes. For +instance, if an iteration calculation is initiated, it can be efficient to +prepare the input data simultaneously, provided that this preparation does not +depend on the kernel's execution. + Querying device capabilities ------------------------------------------------------------------------------- @@ -112,9 +143,9 @@ without blocking other operations. :cpp:func:`hipMemcpyPeerAsync` enables data transfers between different GPUs, facilitating multi-GPU communication. Concurrent data transfers are important for applications that require frequent and large data movements. By overlapping data transfers with computation, -developers can minimize idle times and enhance performance. This is particularly -important for applications that need to handle large volumes of data -efficiently. +developers can minimize idle times and enhance performance. This is +particularly important for applications that need to handle large volumes of +data efficiently. Concurrent data transfers with intra-device copies ------------------------------------------------------------------------------- @@ -123,28 +154,28 @@ It is also possible to perform intra-device copies simultaneously with kernel execution on devices that support the ``concurrentKernels`` device property and/or with copies to or from the device (for devices that support the ``asyncEngineCount`` property). Intra-device copies can be initiated using -standard memory copy functions with destination and source addresses residing on -the same device. +standard memory copy functions with destination and source addresses residing +on the same device. Synchronization, event management and synchronous calls =============================================================================== Synchronization and event management are important for coordinating tasks and -ensuring correct execution order, and synchronous calls are -necessary for maintaining data consistency. +ensuring correct execution order, and synchronous calls are necessary for +maintaining data consistency. Synchronous calls ------------------------------------------------------------------------------- Synchronous calls ensure task completion before moving to the next operation. For example, :cpp:func:`hipMemcpy` for data transfers waits for completion -before returning control to the host. Similarly, synchronous kernel launches are -used when immediate completion is required. When a synchronous function is +before returning control to the host. Similarly, synchronous kernel launches +are used when immediate completion is required. When a synchronous function is called, control is not returned to the host thread before the device has completed the requested task. The behavior of the host thread—whether to yield, -block, or spin—can be specified using ``hipSetDeviceFlags`` with specific flags. -Understanding when to use synchronous calls is important for managing execution -flow and avoiding data races. +block, or spin—can be specified using ``hipSetDeviceFlags`` with specific +flags. Understanding when to use synchronous calls is important for managing +execution flow and avoiding data races. Events for synchronization ------------------------------------------------------------------------------- @@ -162,13 +193,13 @@ to start before the primary kernel finishes, HIP achieves similar functionality using streams and events. By employing :cpp:func:`hipStreamWaitEvent`, it is possible to manage the execution order without explicit hardware support. This mechanism allows a secondary kernel to launch as soon as the necessary -conditions are met, even if the primary kernel is still running. +conditions are met, even if the primary kernel is still running. Example ------------------------------------------------------------------------------- The examples shows the difference between sequential, asynchronous calls and -asynchronous calls with hipEvents. +asynchronous calls with ``hipEvents``. .. figure:: ../../data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg :alt: Compare the different calls @@ -365,28 +396,30 @@ The example codes // Stream 1: Host to Device 1 hipMemcpyAsync(d_data1, h_data1, N * sizeof(*d_data1), hipMemcpyHostToDevice, stream1); + // Stream 2: Host to Device 2 + hipMemcpyAsync(d_data2, h_data2, N * sizeof(*d_data2), hipMemcpyHostToDevice, stream2); + // Stream 1: Kernel 1 - hipLaunchKernelGGL(kernel, dim3(N/256), dim3(256), 0, stream1, d_data1, 1); + hipLaunchKernelGGL(kernel, dim3(N / 256), dim3(256), 0, stream1, d_data1, 1); - // Record event after the GPU kernel in stream1 + // Record event after the GPU kernel in Stream 1 hipEventRecord(event1, stream1); - // Stream 1: Device to Host 1 (after event) + // Stream 2: Wait for Event 1 before starting Kernel 2 hipStreamWaitEvent(stream2, event1, 0); - hipMemcpyAsync(h_data1, d_data1, N * sizeof(*h_data1), hipMemcpyDeviceToHost, stream2); - - // Stream 2: Host to Device 2 - hipMemcpyAsync(d_data2, h_data2, N * sizeof(*d_data2), hipMemcpyHostToDevice, stream2); // Stream 2: Kernel 2 - hipLaunchKernelGGL(kernel, dim3(N/256), dim3(256), 0, stream2, d_data2, 2); + hipLaunchKernelGGL(kernel, dim3(N / 256), dim3(256), 0, stream2, d_data2, 2); - // Record event after the GPU kernel in stream2 + // Record event after Kernel 2 in Stream 2 hipEventRecord(event2, stream2); - // Stream 2: Device to Host 2 (after event) + // Stream 1: Wait for Event 2 before Device to Host copy hipStreamWaitEvent(stream1, event2, 0); - hipMemcpyAsync(h_data2, d_data2, N * sizeof(*h_data2), hipMemcpyDeviceToHost, stream1); + hipMemcpyAsync(h_data1, d_data1, N * sizeof(*h_data1), hipMemcpyDeviceToHost, stream1); + + // Stream 2: Device to Host 2 (after Kernel 2) + hipMemcpyAsync(h_data2, d_data2, N * sizeof(*h_data2), hipMemcpyDeviceToHost, stream2); // Wait for all operations in both streams to complete hipStreamSynchronize(stream1); diff --git a/docs/how-to/performance_guidelines.rst b/docs/how-to/performance_guidelines.rst index bf74b63d16..9b62e4854e 100644 --- a/docs/how-to/performance_guidelines.rst +++ b/docs/how-to/performance_guidelines.rst @@ -3,6 +3,8 @@ developers optimize the performance of HIP-capable GPU architectures. :keywords: AMD, ROCm, HIP, CUDA, performance, guidelines +.. _how_to_performance_guidelines: + ******************************************************************************* Performance guidelines *******************************************************************************