SWDEV-514685 - Update documentation 2025-02-11 (#25)
Co-authored-by: Istvan Kiss <neon60@gmail.com>
Dieser Commit ist enthalten in:
Kiss, Istvan
GitHub
@@ -0,0 +1,534 @@
|
||||
.. meta::
|
||||
:description: This topic describes asynchronous concurrent execution in HIP
|
||||
:keywords: AMD, ROCm, HIP, asynchronous concurrent execution, asynchronous, async, concurrent, concurrency
|
||||
|
||||
.. _asynchronous_how-to:
|
||||
|
||||
*******************************************************************************
|
||||
Asynchronous concurrent execution
|
||||
*******************************************************************************
|
||||
|
||||
Asynchronous concurrent execution is important for efficient parallelism and
|
||||
resource utilization, with techniques such as overlapping computation and data
|
||||
transfer, managing concurrent kernel execution with streams on single or
|
||||
multiple devices, or using HIP graphs.
|
||||
|
||||
Streams and concurrent execution
|
||||
===============================================================================
|
||||
|
||||
All asynchronous APIs, such as kernel execution, data movement and potentially
|
||||
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 task is
|
||||
executed asynchronously. Multiple streams can point to the same device and
|
||||
those streams might be fed from multiple concurrent host-side threads. Multiple
|
||||
streams tied to the same device are not guaranteed to execute their commands in
|
||||
order.
|
||||
|
||||
Managing streams
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Streams enable the overlap of computation and data transfer, ensuring
|
||||
continuous GPU activity. By enabling tasks to run concurrently within the same
|
||||
GPU or across different GPUs, streams improve performance and throughput in
|
||||
high-performance computing (HPC).
|
||||
|
||||
To create a stream, the following functions are used, each defining a handle
|
||||
to the newly created stream:
|
||||
|
||||
- :cpp:func:`hipStreamCreate`: Creates a stream with default settings.
|
||||
- :cpp:func:`hipStreamCreateWithFlags`: Creates a stream, with specific
|
||||
flags, listed below, enabling more control over stream behavior:
|
||||
|
||||
- ``hipStreamDefault``: creates a default stream suitable for most
|
||||
operations. The default stream is a blocking operation.
|
||||
- ``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.
|
||||
|
||||
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.
|
||||
|
||||
Concurrent execution between host and device
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Concurrent execution between the host (CPU) and device (GPU) allows the CPU to
|
||||
perform other tasks while the GPU is executing kernels. Kernels are launched
|
||||
asynchronously using ``hipLaunchKernelGGL`` or using the triple chevron with a stream,
|
||||
enabling the CPU to continue executing other code while the GPU processes the
|
||||
kernel. Similarly, memory operations like :cpp:func:`hipMemcpyAsync` are
|
||||
performed asynchronously, allowing data transfers between the host and device
|
||||
without blocking the CPU.
|
||||
|
||||
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.
|
||||
|
||||
Independent kernels can only run concurrently if there are enough registers
|
||||
and shared memory for the kernels. To enable concurrent kernel executions, the
|
||||
developer may have to reduce the block size of the kernels. The kernel runtimes
|
||||
can be misleading for concurrent kernel runs, that is why during optimization
|
||||
it is a good practice to check the trace files, to see if one kernel is blocking
|
||||
another kernel, while they are running in parallel. For more information about
|
||||
the application tracing, check::doc:`rocprofiler:/how-to/using-rocprof`.
|
||||
|
||||
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.
|
||||
|
||||
Multiple kernels executing concurrently is only beneficial 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
|
||||
===============================================================================
|
||||
|
||||
One of the primary benefits of asynchronous operations and multiple streams 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 a kernel is initiated, it can be efficient to prepare the input
|
||||
data simultaneously, provided that this preparation does not depend on the
|
||||
kernel's execution. Such iterative data transfer and kernel execution overlap
|
||||
can be find in the :ref:`async_example`.
|
||||
|
||||
Querying device capabilities
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Some AMD HIP-enabled devices can perform asynchronous memory copy operations to
|
||||
or from the GPU concurrently with kernel execution. Applications can query this
|
||||
capability by checking the ``asyncEngineCount`` device property. Devices with
|
||||
an ``asyncEngineCount`` greater than zero support concurrent data transfers.
|
||||
Additionally, if host memory is involved in the copy, it should be page-locked
|
||||
to ensure optimal performance. Page-locking (or pinning) host memory increases
|
||||
the bandwidth between the host and the device, reducing the overhead associated
|
||||
with data transfers. For more details, visit :ref:`host_memory` page.
|
||||
|
||||
Asynchronous memory operations
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Asynchronous memory operations do not block the host while copying data and,
|
||||
when used with multiple streams, allow data to be transferred between the host
|
||||
and device while kernels are executed on the same GPU. Using operations like
|
||||
:cpp:func:`hipMemcpyAsync` or :cpp:func:`hipMemcpyPeerAsync`, developers can
|
||||
initiate data transfers without waiting for the previous operation to complete.
|
||||
This overlap of computation and data transfer ensures that the GPU is not idle
|
||||
while waiting for data. :cpp:func:`hipMemcpyPeerAsync` enables data transfers
|
||||
between different GPUs, facilitating multi-GPU communication.
|
||||
|
||||
:ref:`async_example`` include launching kernels in one stream while performing
|
||||
data transfers in another. This technique is especially useful in applications
|
||||
with large data sets that need to be processed quickly.
|
||||
|
||||
Concurrent data transfers with intra-device copies
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Devices that support the ``concurrentKernels`` property can perform
|
||||
intra-device copies concurrently with kernel execution. Additionally, devices
|
||||
that support the ``asyncEngineCount`` property can perform data transfers to
|
||||
or from the GPU simultaneously with kernel execution. Intra-device copies can
|
||||
be initiated using 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.
|
||||
|
||||
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
|
||||
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 :cpp:func:`hipSetDeviceFlags` with
|
||||
appropriate flags. Understanding when to use synchronous calls is important for
|
||||
managing execution flow and avoiding data races.
|
||||
|
||||
Events for synchronization
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
By creating an event with :cpp:func:`hipEventCreate` and recording it with
|
||||
:cpp:func:`hipEventRecord`, developers can synchronize operations across
|
||||
streams, ensuring correct task execution order. :cpp:func:`hipEventSynchronize`
|
||||
lets the application wait for an event to complete before proceeding with the next
|
||||
operation.
|
||||
|
||||
Programmatic dependent launch and synchronization
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
While CUDA supports programmatic dependent launches allowing a secondary kernel
|
||||
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.
|
||||
|
||||
.. _async_example:
|
||||
|
||||
Example
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
The examples shows the difference between sequential, asynchronous calls and
|
||||
asynchronous calls with ``hipEvents``.
|
||||
|
||||
.. figure:: ../../data/how-to/hip_runtime_api/asynchronous/sequential_async_event.svg
|
||||
:alt: Compare the different calls
|
||||
:align: center
|
||||
|
||||
The example codes
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: Sequential
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
// GPU Kernels
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += 1.0;}
|
||||
};
|
||||
__global__ void kernelB(double* arrayA, double* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] += arrayA[x] + 3.0;}
|
||||
};
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int numOfBlocks = 1 << 20;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr int numberOfIterations = 50;
|
||||
// The array size smaller to avoid the relatively short kernel launch compared to memory copies
|
||||
constexpr size_t arraySize = 1U << 25;
|
||||
double *d_dataA;
|
||||
double *d_dataB;
|
||||
|
||||
double initValueA = 0.0;
|
||||
double initValueB = 2.0;
|
||||
|
||||
std::vector<double> vectorA(arraySize, initValueA);
|
||||
std::vector<double> vectorB(arraySize, initValueB);
|
||||
// Allocate device memory
|
||||
HIP_CHECK(hipMalloc(&d_dataA, arraySize * sizeof(*d_dataA)));
|
||||
HIP_CHECK(hipMalloc(&d_dataB, arraySize * sizeof(*d_dataB)));
|
||||
for(int iteration = 0; iteration < numberOfIterations; iteration++)
|
||||
{
|
||||
// Host to Device copies
|
||||
HIP_CHECK(hipMemcpy(d_dataA, vectorA.data(), arraySize * sizeof(*d_dataA), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(d_dataB, vectorB.data(), arraySize * sizeof(*d_dataB), hipMemcpyHostToDevice));
|
||||
// Launch the GPU kernels
|
||||
hipLaunchKernelGGL(kernelA, dim3(numOfBlocks), dim3(threadsPerBlock), 0, 0, d_dataA, arraySize);
|
||||
hipLaunchKernelGGL(kernelB, dim3(numOfBlocks), dim3(threadsPerBlock), 0, 0, d_dataA, d_dataB, arraySize);
|
||||
// Device to Host copies
|
||||
HIP_CHECK(hipMemcpy(vectorA.data(), d_dataA, arraySize * sizeof(*vectorA.data()), hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipMemcpy(vectorB.data(), d_dataB, arraySize * sizeof(*vectorB.data()), hipMemcpyDeviceToHost));
|
||||
}
|
||||
// Wait for all operations to complete
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Verify results
|
||||
const double expectedA = (double)numberOfIterations;
|
||||
const double expectedB =
|
||||
initValueB + (3.0 * numberOfIterations) +
|
||||
(expectedA * (expectedA + 1.0)) / 2.0;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(vectorA[i] != expectedA){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedA << " got " << vectorA[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
if(vectorB[i] != expectedB){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedB << " got " << vectorB[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if(passed){
|
||||
std::cout << "Sequential execution completed successfully." << std::endl;
|
||||
}else{
|
||||
std::cerr << "Sequential execution failed." << std::endl;
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipFree(d_dataA));
|
||||
HIP_CHECK(hipFree(d_dataB));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: Asynchronous
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
// GPU Kernels
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += 1.0;}
|
||||
};
|
||||
__global__ void kernelB(double* arrayA, double* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] += arrayA[x] + 3.0;}
|
||||
};
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int numOfBlocks = 1 << 20;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr int numberOfIterations = 50;
|
||||
// The array size smaller to avoid the relatively short kernel launch compared to memory copies
|
||||
constexpr size_t arraySize = 1U << 25;
|
||||
double *d_dataA;
|
||||
double *d_dataB;
|
||||
|
||||
double initValueA = 0.0;
|
||||
double initValueB = 2.0;
|
||||
|
||||
std::vector<double> vectorA(arraySize, initValueA);
|
||||
std::vector<double> vectorB(arraySize, initValueB);
|
||||
// Allocate device memory
|
||||
HIP_CHECK(hipMalloc(&d_dataA, arraySize * sizeof(*d_dataA)));
|
||||
HIP_CHECK(hipMalloc(&d_dataB, arraySize * sizeof(*d_dataB)));
|
||||
// Create streams
|
||||
hipStream_t streamA, streamB;
|
||||
HIP_CHECK(hipStreamCreate(&streamA));
|
||||
HIP_CHECK(hipStreamCreate(&streamB));
|
||||
for(unsigned int iteration = 0; iteration < numberOfIterations; iteration++)
|
||||
{
|
||||
// Stream 1: Host to Device 1
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataA, vectorA.data(), arraySize * sizeof(*d_dataA), hipMemcpyHostToDevice, streamA));
|
||||
// Stream 2: Host to Device 2
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataB, vectorB.data(), arraySize * sizeof(*d_dataB), hipMemcpyHostToDevice, streamB));
|
||||
// Stream 1: Kernel 1
|
||||
hipLaunchKernelGGL(kernelA, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamA, d_dataA, arraySize);
|
||||
// Wait for streamA finish
|
||||
HIP_CHECK(hipStreamSynchronize(streamA));
|
||||
// Stream 2: Kernel 2
|
||||
hipLaunchKernelGGL(kernelB, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamB, d_dataA, d_dataB, arraySize);
|
||||
// Stream 1: Device to Host 2 (after Kernel 1)
|
||||
HIP_CHECK(hipMemcpyAsync(vectorA.data(), d_dataA, arraySize * sizeof(*vectorA.data()), hipMemcpyDeviceToHost, streamA));
|
||||
// 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(hipStreamSynchronize(streamA));
|
||||
HIP_CHECK(hipStreamSynchronize(streamB));
|
||||
// Verify results
|
||||
double expectedA = (double)numberOfIterations;
|
||||
double expectedB =
|
||||
initValueB + (3.0 * numberOfIterations) +
|
||||
(expectedA * (expectedA + 1.0)) / 2.0;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(vectorA[i] != expectedA){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedA << " got " << vectorA[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
if(vectorB[i] != expectedB){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedB << " got " << vectorB[i] << " at index: " << i << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if(passed){
|
||||
std::cout << "Asynchronous execution completed successfully." << std::endl;
|
||||
}else{
|
||||
std::cerr << "Asynchronous execution failed." << std::endl;
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipStreamDestroy(streamA));
|
||||
HIP_CHECK(hipStreamDestroy(streamB));
|
||||
HIP_CHECK(hipFree(d_dataA));
|
||||
HIP_CHECK(hipFree(d_dataB));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: hipStreamWaitEvent
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
// GPU Kernels
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += 1.0;}
|
||||
};
|
||||
__global__ void kernelB(double* arrayA, double* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] += arrayA[x] + 3.0;}
|
||||
};
|
||||
|
||||
int main()
|
||||
{
|
||||
constexpr int numOfBlocks = 1 << 20;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr int numberOfIterations = 50;
|
||||
// The array size smaller to avoid the relatively short kernel launch compared to memory copies
|
||||
constexpr size_t arraySize = 1U << 25;
|
||||
double *d_dataA;
|
||||
double *d_dataB;
|
||||
double initValueA = 0.0;
|
||||
double initValueB = 2.0;
|
||||
|
||||
std::vector<double> vectorA(arraySize, initValueA);
|
||||
std::vector<double> vectorB(arraySize, initValueB);
|
||||
// Allocate device memory
|
||||
HIP_CHECK(hipMalloc(&d_dataA, arraySize * sizeof(*d_dataA)));
|
||||
HIP_CHECK(hipMalloc(&d_dataB, arraySize * sizeof(*d_dataB)));
|
||||
// Create streams
|
||||
hipStream_t streamA, streamB;
|
||||
HIP_CHECK(hipStreamCreate(&streamA));
|
||||
HIP_CHECK(hipStreamCreate(&streamB));
|
||||
// Create events
|
||||
hipEvent_t event, eventA, eventB;
|
||||
HIP_CHECK(hipEventCreate(&event));
|
||||
HIP_CHECK(hipEventCreate(&eventA));
|
||||
HIP_CHECK(hipEventCreate(&eventB));
|
||||
for(unsigned int iteration = 0; iteration < numberOfIterations; iteration++)
|
||||
{
|
||||
// Stream 1: Host to Device 1
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataA, vectorA.data(), arraySize * sizeof(*d_dataA), hipMemcpyHostToDevice, streamA));
|
||||
// Stream 2: Host to Device 2
|
||||
HIP_CHECK(hipMemcpyAsync(d_dataB, vectorB.data(), arraySize * sizeof(*d_dataB), hipMemcpyHostToDevice, streamB));
|
||||
// Stream 1: Kernel 1
|
||||
hipLaunchKernelGGL(kernelA, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamA, d_dataA, arraySize);
|
||||
// Record event after the GPU kernel in Stream 1
|
||||
HIP_CHECK(hipEventRecord(event, streamA));
|
||||
// Stream 2: Wait for event before starting Kernel 2
|
||||
HIP_CHECK(hipStreamWaitEvent(streamB, event, 0));
|
||||
// Stream 2: Kernel 2
|
||||
hipLaunchKernelGGL(kernelB, dim3(numOfBlocks), dim3(threadsPerBlock), 0, streamB, d_dataA, d_dataB, arraySize);
|
||||
// Stream 1: Device to Host 2 (after Kernel 1)
|
||||
HIP_CHECK(hipMemcpyAsync(vectorA.data(), d_dataA, arraySize * sizeof(*vectorA.data()), hipMemcpyDeviceToHost, streamA));
|
||||
// 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));
|
||||
}
|
||||
// Verify results
|
||||
double expectedA = (double)numberOfIterations;
|
||||
double expectedB =
|
||||
initValueB + (3.0 * numberOfIterations) +
|
||||
(expectedA * (expectedA + 1.0)) / 2.0;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(vectorA[i] != expectedA){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedA << " got " << vectorA[i] << std::endl;
|
||||
break;
|
||||
}
|
||||
if(vectorB[i] != expectedB){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expectedB << " got " << vectorB[i] << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if(passed){
|
||||
std::cout << "Asynchronous execution with events completed successfully." << std::endl;
|
||||
}else{
|
||||
std::cerr << "Asynchronous execution with events failed." << std::endl;
|
||||
}
|
||||
|
||||
// Cleanup
|
||||
HIP_CHECK(hipEventDestroy(event));
|
||||
HIP_CHECK(hipEventDestroy(eventA));
|
||||
HIP_CHECK(hipEventDestroy(eventB));
|
||||
HIP_CHECK(hipStreamDestroy(streamA));
|
||||
HIP_CHECK(hipStreamDestroy(streamB));
|
||||
HIP_CHECK(hipFree(d_dataA));
|
||||
HIP_CHECK(hipFree(d_dataB));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
HIP Graphs
|
||||
===============================================================================
|
||||
|
||||
HIP graphs offer an efficient alternative to the standard method of launching
|
||||
GPU tasks via streams. Comprising nodes for operations and edges for
|
||||
dependencies, HIP graphs reduce kernel launch overhead and provide a high-level
|
||||
abstraction for managing dependencies and synchronization. By representing
|
||||
sequences of kernels and memory operations as a single graph, they simplify
|
||||
complex workflows and enhance performance, particularly for applications with
|
||||
intricate dependencies and multiple execution stages.
|
||||
For more details, see the :ref:`how_to_HIP_graph` documentation.
|
||||
@@ -1,52 +1,285 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the device memory of the HIP ecosystem
|
||||
ROCm software.
|
||||
:keywords: AMD, ROCm, HIP, device memory
|
||||
:keywords: AMD, ROCm, HIP, GPU, device memory, global, constant, texture, surface, shared
|
||||
|
||||
.. _device_memory:
|
||||
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
Device memory
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
|
||||
Device memory exists on the device, e.g. on GPUs in the video random access
|
||||
memory (VRAM), and is accessible by the kernels operating on the device. Recent
|
||||
architectures use graphics double data rate (GDDR) synchronous dynamic
|
||||
random-access memory (SDRAM) such as GDDR6, or high-bandwidth memory (HBM) such
|
||||
as HBM2e. Device memory can be allocated as global memory, constant, texture or
|
||||
surface memory.
|
||||
Device memory is random access memory that is physically located on a GPU. In
|
||||
general it is memory with a bandwidth that is an order of magnitude higher
|
||||
compared to RAM available to the host. That high bandwidth is only available to
|
||||
on-device accesses, accesses from the host or other devices have to go over a
|
||||
special interface which is considerably slower, usually the PCIe bus or the AMD
|
||||
Infinity Fabric.
|
||||
|
||||
On certain architectures like APUs, the GPU and CPU share the same physical
|
||||
memory.
|
||||
|
||||
There is also a special local data share on-chip directly accessible to the
|
||||
:ref:`compute units <hardware_implementation>`, that can be used for shared
|
||||
memory.
|
||||
|
||||
The physical device memory can be used to back up several different memory
|
||||
spaces in HIP, as described in the following.
|
||||
|
||||
Global memory
|
||||
================================================================================
|
||||
|
||||
Read-write storage visible to all threads on a given device. There are
|
||||
specialized versions of global memory with different usage semantics which are
|
||||
typically backed by the same hardware, but can use different caching paths.
|
||||
Global memory is the general read-write accessible memory visible to all threads
|
||||
on a given device. Since variables located in global memory have to be marked
|
||||
with the ``__device__`` qualifier, this memory space is also referred to as
|
||||
device memory.
|
||||
|
||||
Without explicitly copying it, it can only be accessed by the threads within a
|
||||
kernel operating on the device, however :ref:`unified_memory` can be used to
|
||||
let the runtime manage this, if desired.
|
||||
|
||||
Allocating global memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
This memory needs to be explicitly allocated.
|
||||
|
||||
It can be allocated from the host via the :ref:`HIP runtime memory management
|
||||
functions <memory_management_reference>` like :cpp:func:`hipMalloc`, or can be
|
||||
defined using the ``__device__`` qualifier on variables.
|
||||
|
||||
It can also be allocated within a kernel using ``malloc`` or ``new``.
|
||||
The specified amount of memory is allocated by each thread that executes the
|
||||
instructions. The recommended way to allocate the memory depends on the use
|
||||
case. If the memory is intended to be shared between the threads of a block, it
|
||||
is generally beneficial to allocate one large block of memory, due to the way
|
||||
the memory is accessed.
|
||||
|
||||
.. note::
|
||||
Memory allocated within a kernel can only be freed in kernels, not by the HIP
|
||||
runtime on the host, like :cpp:func:`hipFree`. It is also not possible to
|
||||
free device memory allocated on the host, with :cpp:func:`hipMalloc` for
|
||||
example, in a kernel.
|
||||
|
||||
|
||||
An example for how to share memory allocated within a kernel by only one thread
|
||||
is given in the following example. In case the device memory is only needed for
|
||||
communication between the threads in a single block, :ref:`shared_memory` is the
|
||||
better option, but is also limited in size.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void kernel_memory_allocation(TYPE* pointer){
|
||||
// The pointer is stored in shared memory, so that all
|
||||
// threads of the block can access the pointer
|
||||
__shared__ int *memory;
|
||||
|
||||
size_t blockSize = blockDim.x;
|
||||
constexpr size_t elementsPerThread = 1024;
|
||||
if(threadIdx.x == 0){
|
||||
// allocate memory in one contiguous block
|
||||
memory = new int[blockDim.x * elementsPerThread];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// load pointer into thread-local variable to avoid
|
||||
// unnecessary accesses to shared memory
|
||||
int *localPtr = memory;
|
||||
|
||||
// work with allocated memory, e.g. initialization
|
||||
for(int i = 0; i < elementsPerThread; ++i){
|
||||
// access in a contiguous way
|
||||
localPtr[i * blockSize + threadIdx.x] = i;
|
||||
}
|
||||
|
||||
// synchronize to make sure no thread is accessing the memory before freeing
|
||||
__syncthreads();
|
||||
if(threadIdx.x == 0){
|
||||
delete[] memory;
|
||||
}
|
||||
}
|
||||
|
||||
Copying between device and host
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
When not using :ref:`unified_memory`, memory has to be explicitly copied between
|
||||
the device and the host, using the HIP runtime API.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
size_t elements = 1 << 20;
|
||||
size_t size_bytes = elements * sizeof(int);
|
||||
|
||||
// allocate host and device memory
|
||||
int *host_pointer = new int[elements];
|
||||
int *device_input, *device_result;
|
||||
HIP_CHECK(hipMalloc(&device_input, size_bytes));
|
||||
HIP_CHECK(hipMalloc(&device_result, size_bytes));
|
||||
|
||||
// copy from host to the device
|
||||
HIP_CHECK(hipMemcpy(device_input, host_pointer, size_bytes, hipMemcpyHostToDevice));
|
||||
|
||||
// Use memory on the device, i.e. execute kernels
|
||||
|
||||
// copy from device to host, to e.g. get results from the kernel
|
||||
HIP_CHECK(hipMemcpy(host_pointer, device_result, size_bytes, hipMemcpyDeviceToHost));
|
||||
|
||||
// free memory when not needed any more
|
||||
HIP_CHECK(hipFree(device_result));
|
||||
HIP_CHECK(hipFree(device_input));
|
||||
delete[] host_pointer;
|
||||
|
||||
Constant memory
|
||||
================================================================================
|
||||
|
||||
Read-only storage visible to all threads on a given device. It is a limited
|
||||
segment backed by device memory with queryable size. It needs to be set by the
|
||||
host before kernel execution. Constant memory provides the best performance
|
||||
benefit when all threads within a warp access the same address.
|
||||
Constant memory is read-only storage visible to all threads on a given device.
|
||||
It is a limited segment backed by device memory, that takes a different caching
|
||||
route than normal device memory accesses. It needs to be set by the host before
|
||||
kernel execution.
|
||||
|
||||
In order to get the highest bandwidth from the constant memory, all threads of
|
||||
a warp have to access the same memory address. If they access different
|
||||
addresses, the accesses get serialized and the bandwidth is therefore reduced.
|
||||
|
||||
Using constant memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Constant memory can not be dynamically allocated, and the size has to be
|
||||
specified during compile time. If the values can not be specified during compile
|
||||
time, they have to be set by the host before the kernel, that accesses the
|
||||
constant memory, is called.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
constexpr size_t const_array_size = 32;
|
||||
__constant__ double const_array[const_array_size];
|
||||
|
||||
void set_constant_memory(double* values){
|
||||
hipMemcpyToSymbol(const_array, values, const_array_size * sizeof(double));
|
||||
}
|
||||
|
||||
__global__ void kernel_using_const_memory(double* array){
|
||||
|
||||
int warpIdx = threadIdx.x / warpSize;
|
||||
// uniform access of warps to const_array for best performance
|
||||
array[blockDim.x] *= const_array[warpIdx];
|
||||
}
|
||||
|
||||
Texture memory
|
||||
================================================================================
|
||||
|
||||
Read-only storage visible to all threads on a given device and accessible
|
||||
through additional APIs. Its origins come from graphics APIs, and provides
|
||||
performance benefits when accessing memory in a pattern where the
|
||||
addresses are close to each other in a 2D representation of the memory.
|
||||
Texture memory is special read-only memory visible to all threads on a given
|
||||
device and accessible through additional APIs. Its origins come from graphics
|
||||
APIs, and provides performance benefits when accessing memory in a pattern where
|
||||
the addresses are close to each other in a 2D or 3D representation of the
|
||||
memory. It also provides additional features like filtering and addressing for
|
||||
out-of-bounds accesses, which are further explained in :ref:`texture_fetching`.
|
||||
|
||||
The :ref:`texture management module <texture_management_reference>` of the HIP
|
||||
runtime API reference contains the functions of texture memory.
|
||||
The original use of the texture cache was also to take pressure off the global
|
||||
memory and other caches, however on modern GPUs, that support textures, the L1
|
||||
cache and texture cache are combined, so the main purpose is to make use of the
|
||||
texture specific features.
|
||||
|
||||
To find out whether textures are supported on a device, query
|
||||
:cpp:enumerator:`hipDeviceAttributeImageSupport`.
|
||||
|
||||
Using texture memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Textures are more complex than just a region of memory, so their layout has to
|
||||
be specified. They are represented by ``hipTextureObject_t`` and created using
|
||||
:cpp:func:`hipCreateTextureObject`.
|
||||
|
||||
The underlying memory is a 1D, 2D or 3D ``hipArray_t``, that needs to be
|
||||
allocated using :cpp:func:`hipMallocArray`.
|
||||
|
||||
On the device side, texture objects are accessed using the ``tex1D/2D/3D``
|
||||
functions.
|
||||
|
||||
The texture management functions can be found in the :ref:`Texture management
|
||||
API reference <texture_management_reference>`
|
||||
|
||||
A full example for how to use textures can be found in the `ROCm texture
|
||||
management example <https://github.com/ROCm/rocm-examples/blob/develop/HIP-Basic/texture_management/main.hip>`_
|
||||
|
||||
Surface memory
|
||||
================================================================================
|
||||
|
||||
A read-write version of texture memory, which can be useful for applications
|
||||
that require direct manipulation of 1D, 2D, or 3D hipArray_t.
|
||||
A read-write version of texture memory. It is created in the same way as a
|
||||
texture, but with :cpp:func:`hipCreateSurfaceObject`.
|
||||
|
||||
Since surfaces are also cached in the read-only texture cache, the changes
|
||||
written back to the surface can't be observed in the same kernel. A new kernel
|
||||
has to be launched in order to see the updated surface.
|
||||
|
||||
The corresponding functions are listed in the :ref:`Surface object API reference
|
||||
<surface_object_reference>`.
|
||||
|
||||
.. _shared_memory:
|
||||
|
||||
Shared memory
|
||||
================================================================================
|
||||
|
||||
Shared memory is read-write memory, that is only visible to the threads within a
|
||||
block. It is allocated per thread block, and needs to be either statically
|
||||
allocated at compile time, or can be dynamically allocated when launching the
|
||||
kernel, but not during kernel execution. Its general use-case is to share
|
||||
variables between the threads within a block, but can also be used as scratch
|
||||
pad memory.
|
||||
|
||||
Shared memory is not backed by the same physical memory as the other address
|
||||
spaces. It is on-chip memory local to the :ref:`compute units
|
||||
<hardware_implementation>`, providing low-latency, high-bandwidth access,
|
||||
comparable to the L1 cache. It is however limited in size, and as it is
|
||||
allocated per block, can restrict how many blocks can be scheduled to a compute
|
||||
unit concurrently, thereby potentially reducing occupancy.
|
||||
|
||||
An overview of the size of the local data share (LDS), that backs up shared
|
||||
memory, is given in the
|
||||
:doc:`GPU hardware specifications <rocm:reference/gpu-arch-specs>`.
|
||||
|
||||
Allocate shared memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Memory can be dynamically allocated by declaring an ``extern __shared__`` array,
|
||||
whose size can be set during kernel launch, which can then be accessed in the
|
||||
kernel.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
extern __shared__ int dynamic_shared[];
|
||||
__global__ void kernel(int array1SizeX, int array1SizeY, int array2Size){
|
||||
// at least (array1SizeX * array1SizeY + array2Size) * sizeof(int) bytes
|
||||
// dynamic shared memory need to be allocated when the kernel is launched
|
||||
int* array1 = dynamic_shared;
|
||||
// array1 is interpreted as 2D of size:
|
||||
int array1Size = array1SizeX * array1SizeY;
|
||||
|
||||
int* array2 = &(array1[array1Size]);
|
||||
|
||||
if(threadIdx.x < array1SizeX && threadIdx.y < array1SizeY){
|
||||
// access array1 with threadIdx.x + threadIdx.y * array1SizeX
|
||||
}
|
||||
if(threadIdx.x < array2Size){
|
||||
// access array2 threadIdx.x
|
||||
}
|
||||
}
|
||||
|
||||
A more in-depth example on dynamically allocated shared memory can be found in
|
||||
the `ROCm dynamic shared example
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/dynamic_shared>`_.
|
||||
|
||||
To statically allocate shared memory, just declare it in the kernel. The memory
|
||||
is allocated per block, not per thread. If the kernel requires more shared
|
||||
memory than is available to the architecture, the compilation fails.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void kernel(){
|
||||
__shared__ int array[128];
|
||||
__shared__ double result;
|
||||
}
|
||||
|
||||
A more in-depth example on statically allocated shared memory can be found in
|
||||
the `ROCm shared memory example
|
||||
<https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/shared_memory>`_.
|
||||
|
||||
The :ref:`surface objects module <surface_object_reference>` of HIP runtime API
|
||||
contains the functions for creating, destroying and reading surface memory.
|
||||
+66
-39
@@ -5,56 +5,67 @@
|
||||
|
||||
.. _texture_fetching:
|
||||
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
Texture fetching
|
||||
*******************************************************************************
|
||||
********************************************************************************
|
||||
|
||||
`Textures <../../../../doxygen/html/group___texture.html>`_ are more than just a buffer
|
||||
interpreted as a 1D, 2D, or 3D array.
|
||||
Textures give access to specialized hardware on GPUs that is usually used in
|
||||
graphics processing. In particular, textures use a different way of accessing
|
||||
their underlying device memory. Memory accesses to textures are routed through
|
||||
a special read-only texture cache, that is optimized for logical spatial
|
||||
locality, e.g. locality in 2D grids. This can also benefit certain algorithms
|
||||
used in GPGPU computing, when the access pattern is the same as used when
|
||||
accessing normal textures.
|
||||
|
||||
As textures are associated with graphics, they are indexed using floating-point
|
||||
values. The index can be in the range of [0 to size-1] or [0 to 1].
|
||||
Additionally, textures can be indexed using floating-point values. This is used
|
||||
in graphics applications to interpolate between neighboring values of a texture.
|
||||
Depending on the interpolation mode the index can be in the range of ``0`` to
|
||||
``size - 1`` or ``0`` to ``1``. Textures also have a way of handling
|
||||
out-of-bounds accesses.
|
||||
|
||||
Depending on the index, texture sampling or texture addressing is performed,
|
||||
which decides the return value.
|
||||
Depending on the value of the index, :ref:`texture filtering <texture_filtering>`
|
||||
or :ref:`texture addressing <texture_addressing>` is performed.
|
||||
|
||||
**Texture sampling**: When a texture is indexed with a fraction, the queried
|
||||
value is often between two or more texels (texture elements). The sampling
|
||||
method defines what value to return in such cases.
|
||||
|
||||
**Texture addressing**: Sometimes, the index is outside the bounds of the
|
||||
texture. This condition might look like a problem but helps to put a texture on
|
||||
a surface multiple times or to create a visible sign of out-of-bounds indexing,
|
||||
in computer graphics. The addressing mode defines what value to return when
|
||||
indexing a texture out of bounds.
|
||||
|
||||
The different sampling and addressing modes are described in the following
|
||||
sections.
|
||||
|
||||
Here is the sample texture used in this document for demonstration purposes. It
|
||||
Here is the example texture used in this document for demonstration purposes. It
|
||||
is 2x2 texels and indexed in the [0 to 1] range.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/original.png
|
||||
:width: 150
|
||||
:alt: Sample texture
|
||||
:alt: Example texture
|
||||
:align: center
|
||||
|
||||
Texture used as example
|
||||
|
||||
Texture sampling
|
||||
===============================================================================
|
||||
In HIP textures objects are of type :cpp:struct:`hipTextureObject_t` and created
|
||||
using :cpp:func:`hipCreateTextureObject`.
|
||||
|
||||
Texture sampling handles the usage of fractional indices. It is the method that
|
||||
describes, which nearby values will be used, and how they are combined into the
|
||||
resulting value.
|
||||
For a full list of available texture functions see the :ref:`HIP texture API
|
||||
reference <texture_management_reference>`.
|
||||
|
||||
The various texture sampling methods are discussed in the following sections.
|
||||
A code example for how to use textures can be found in the `ROCm texture
|
||||
management example <https://github.com/ROCm/rocm-examples/blob/develop/HIP-Basic/texture_management/main.hip>`_
|
||||
|
||||
.. _texture_filtering:
|
||||
|
||||
Texture filtering
|
||||
================================================================================
|
||||
|
||||
Texture filtering handles the usage of fractional indices. When the index is a
|
||||
fraction, the queried value lies between two or more texels (texture elements),
|
||||
depending on the dimensionality of the texture. The filtering method defines how
|
||||
to interpolate between these values.
|
||||
|
||||
The filter modes are specified in :cpp:enumerator:`hipTextureFilterMode`.
|
||||
|
||||
The various texture filtering methods are discussed in the following sections.
|
||||
|
||||
.. _texture_fetching_nearest:
|
||||
|
||||
Nearest point sampling
|
||||
Nearest point filtering
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This filter mode corresponds to ``hipFilterModePoint``.
|
||||
|
||||
In this method, the modulo of index is calculated as:
|
||||
|
||||
``tex(x) = T[floor(x)]``
|
||||
@@ -70,22 +81,24 @@ of the nearest texel.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/nearest.png
|
||||
:width: 300
|
||||
:alt: Texture upscaled with nearest point sampling
|
||||
:alt: Texture upscaled with nearest point filtering
|
||||
:align: center
|
||||
|
||||
Texture upscaled with nearest point sampling
|
||||
Texture upscaled with nearest point filtering
|
||||
|
||||
.. _texture_fetching_linear:
|
||||
|
||||
Linear filtering
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This filter mode corresponds to ``hipFilterModeLinear``.
|
||||
|
||||
The linear filtering method does a linear interpolation between values. Linear
|
||||
interpolation is used to create a linear transition between two values. The
|
||||
formula used is ``(1-t)P1 + tP2`` where ``P1`` and ``P2`` are the values and
|
||||
``t`` is within the [0 to 1] range.
|
||||
|
||||
In the case of texture sampling the following formulas are used:
|
||||
In the case of linear texture filtering the following formulas are used:
|
||||
|
||||
* For one dimensional textures: ``tex(x) = (1-α)T[i] + αT[i+1]``
|
||||
* For two dimensional textures: ``tex(x,y) = (1-α)(1-β)T[i,j] + α(1-β)T[i+1,j] + (1-α)βT[i,j+1] + αβT[i+1,j+1]``
|
||||
@@ -95,7 +108,7 @@ Where x, y, and, z are the floating-point indices. i, j, and, k are the integer
|
||||
indices and, α, β, and, γ values represent how far along the sampled point is on
|
||||
the three axes. These values are calculated by these formulas: ``i = floor(x')``, ``α = frac(x')``, ``x' = x - 0.5``, ``j = floor(y')``, ``β = frac(y')``, ``y' = y - 0.5``, ``k = floor(z')``, ``γ = frac(z')`` and ``z' = z - 0.5``
|
||||
|
||||
This following image shows a texture stretched out to a 4x4 pixel quad, but
|
||||
The following image shows a texture stretched out to a 4x4 pixel quad, but
|
||||
still indexed in the [0 to 1] range. The in-between values are interpolated
|
||||
between the neighboring texels.
|
||||
|
||||
@@ -106,12 +119,18 @@ between the neighboring texels.
|
||||
|
||||
Texture upscaled with linear filtering
|
||||
|
||||
.. _texture_addressing:
|
||||
|
||||
Texture addressing
|
||||
===============================================================================
|
||||
|
||||
Texture addressing mode handles the index that is out of bounds of the texture.
|
||||
This mode describes which values of the texture or a preset value to use when
|
||||
the index is out of bounds.
|
||||
The texture addressing modes are specified in
|
||||
:cpp:enumerator:`hipTextureAddressMode`.
|
||||
|
||||
The texture addressing mode handles out-of-bounds accesses to the texture. This
|
||||
can be used in graphics applications to e.g. repeat a texture on a surface
|
||||
multiple times in various ways or create visible signs of out-of-bounds
|
||||
indexing.
|
||||
|
||||
The following sections describe the various texture addressing methods.
|
||||
|
||||
@@ -120,8 +139,10 @@ The following sections describe the various texture addressing methods.
|
||||
Address mode border
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
In this method, the texture fetching returns a border value when indexing out of
|
||||
bounds. The border value must be set before texture fetching.
|
||||
This addressing mode is set using ``hipAddressModeBorder``.
|
||||
|
||||
This addressing mode returns a border value when indexing out of bounds. The
|
||||
border value must be set before texture fetching.
|
||||
|
||||
The following image shows the texture on a 4x4 pixel quad, indexed in the
|
||||
[0 to 3] range. The out-of-bounds values are the border color, which is yellow.
|
||||
@@ -141,6 +162,8 @@ the addressing begins.
|
||||
Address mode clamp
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This addressing mode is set using ``hipAddressModeClamp``.
|
||||
|
||||
This mode clamps the index between [0 to size-1]. Due to this, when indexing
|
||||
out-of-bounds, the values on the edge of the texture repeat. The clamp mode is
|
||||
the default addressing mode.
|
||||
@@ -164,6 +187,8 @@ the addressing begins.
|
||||
Address mode wrap
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This addressing mode is set using ``hipAddressModeWrap``.
|
||||
|
||||
Wrap mode addressing is only available for normalized texture coordinates. In
|
||||
this addressing mode, the fractional part of the index is used:
|
||||
|
||||
@@ -189,6 +214,8 @@ the addressing begins.
|
||||
Address mode mirror
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
This addressing mode is set using ``hipAddressModeMirror``.
|
||||
|
||||
Similar to the wrap mode the mirror mode is only available for normalized
|
||||
texture coordinates and also creates a repeating image, but mirroring the
|
||||
neighboring instances.
|
||||
|
||||
@@ -111,8 +111,7 @@ allocator can be used.
|
||||
❌: **Unsupported**
|
||||
|
||||
:sup:`1` Works only with ``HSA_XNACK=1`` and kernels with HMM support. First GPU
|
||||
access causes recoverable page-fault. For more details, visit `GPU memory
|
||||
<https://rocm.docs.amd.com/en/latest/conceptual/gpu-memory.html#xnack>`_.
|
||||
access causes recoverable page-fault.
|
||||
|
||||
.. _unified memory allocators:
|
||||
|
||||
@@ -144,8 +143,7 @@ GPUs, it is essential to set the environment variable ``HSA_XNACK=1`` and use
|
||||
a GPU kernel mode driver that supports HMM
|
||||
<https://www.kernel.org/doc/html/latest/mm/hmm.html>`_. Without this
|
||||
configuration, the behavior will be similar to that of systems without HMM
|
||||
support. For more details, visit
|
||||
`GPU memory <https://rocm.docs.amd.com/en/latest/conceptual/gpu-memory.html#xnack>`_.
|
||||
support.
|
||||
|
||||
The table below illustrates the expected behavior of managed and unified memory
|
||||
functions on ROCm and CUDA, both with and without HMM support.
|
||||
|
||||
@@ -28,7 +28,7 @@ reduce memory usage and unnecessary ``memcpy`` calls.
|
||||
.. _memory_allocation_virtual_memory:
|
||||
|
||||
Memory allocation
|
||||
================================================================================
|
||||
=================
|
||||
|
||||
Standard memory allocation uses the :cpp:func:`hipMalloc` function to allocate a
|
||||
block of memory on the device. However, when using virtual memory, this process
|
||||
@@ -37,10 +37,34 @@ is separated into multiple steps using the :cpp:func:`hipMemCreate`,
|
||||
:cpp:func:`hipMemSetAccess` functions. This guide explains what these functions
|
||||
do and how you can use them for virtual memory management.
|
||||
|
||||
Allocate physical memory
|
||||
--------------------------------------------------------------------------------
|
||||
.. _vmm_support:
|
||||
|
||||
The first step is to allocate the physical memory itself with the
|
||||
Virtual memory management support
|
||||
---------------------------------
|
||||
|
||||
The first step is to check if the targeted device or GPU supports virtual memory management.
|
||||
Use the :cpp:func:`hipDeviceGetAttribute` function to get the
|
||||
``hipDeviceAttributeVirtualMemoryManagementSupported`` attribute for a specific GPU, as shown in the following example.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
int vmm = 0, currentDev = 0;
|
||||
hipDeviceGetAttribute(
|
||||
&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, currentDev
|
||||
);
|
||||
|
||||
if (vmm == 0) {
|
||||
std::cout << "GPU " << currentDev << " doesn't support virtual memory management." << std::endl;
|
||||
} else {
|
||||
std::cout << "GPU " << currentDev << " support virtual memory management." << std::endl;
|
||||
}
|
||||
|
||||
.. _allocate_physical_memory:
|
||||
|
||||
Allocate physical memory
|
||||
------------------------
|
||||
|
||||
The next step is to allocate the physical memory using the
|
||||
:cpp:func:`hipMemCreate` function. This function accepts the size of the buffer,
|
||||
an ``unsigned long long`` variable for the flags, and a
|
||||
:cpp:struct:`hipMemAllocationProp` variable. :cpp:struct:`hipMemAllocationProp`
|
||||
@@ -48,42 +72,54 @@ contains the properties of the memory to be allocated, such as where the memory
|
||||
is physically located and what kind of shareable handles are available. If the
|
||||
allocation is successful, the function returns a value of
|
||||
:cpp:enumerator:`hipSuccess`, with :cpp:type:`hipMemGenericAllocationHandle_t`
|
||||
representing a valid physical memory allocation. The allocated memory size must
|
||||
be aligned with the granularity appropriate for the properties of the
|
||||
allocation. You can use the :cpp:func:`hipMemGetAllocationGranularity` function
|
||||
to determine the correct granularity.
|
||||
representing a valid physical memory allocation.
|
||||
|
||||
The allocated memory must be aligned with the appropriate granularity. The
|
||||
granularity value can be queried with :cpp:func:`hipMemGetAllocationGranularity`,
|
||||
and its value depends on the target device hardware and the type of memory
|
||||
allocation. If the allocation size is not aligned, meaning it is not cleanly
|
||||
divisible by the minimum granularity value, :cpp:func:`hipMemCreate` will return
|
||||
an out-of-memory error.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
size_t granularity = 0;
|
||||
hipMemGenericAllocationHandle_t allocHandle;
|
||||
hipMemAllocationProp prop = {};
|
||||
prop.type = HIP_MEM_ALLOCATION_TYPE_PINNED;
|
||||
prop.location.type = HIP_MEM_LOCATION_TYPE_DEVICE;
|
||||
// The pinned allocation type cannot be migrated from its current location
|
||||
// while the application is actively using it.
|
||||
prop.type = hipMemAllocationTypePinned;
|
||||
// Set the location type to device, currently there are no other valid option.
|
||||
prop.location.type = hipMemLocationTypeDevice;
|
||||
// Set the device id, where the memory will be allocated.
|
||||
prop.location.id = currentDev;
|
||||
hipMemGetAllocationGranularity(&granularity, &prop, HIP_MEM_ALLOC_GRANULARITY_MINIMUM);
|
||||
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum);
|
||||
padded_size = ROUND_UP(size, granularity);
|
||||
hipMemCreate(&allocHandle, padded_size, &prop, 0);
|
||||
|
||||
Reserve virtual address range
|
||||
--------------------------------------------------------------------------------
|
||||
.. _reserve_virtual_address:
|
||||
|
||||
After you have acquired an allocation of physical memory, you must map it before
|
||||
you can use it. To do so, you need a virtual address to map it to. Mapping
|
||||
means the physical memory allocation is available from the virtual address range
|
||||
it is mapped to. To reserve a virtual memory range, use the
|
||||
:cpp:func:`hipMemAddressReserve` function. The size of the virtual memory must
|
||||
match the amount of physical memory previously allocated. You can then map the
|
||||
physical memory allocation to the newly-acquired virtual memory address range
|
||||
using the :cpp:func:`hipMemMap` function.
|
||||
Reserve virtual address range
|
||||
-----------------------------
|
||||
|
||||
After you have acquired an allocation of physical memory, you must map it to a
|
||||
virtual address before you can use it. Mapping means the physical memory
|
||||
allocation is available from the virtual address range it is mapped to. To
|
||||
reserve a virtual memory range, use the :cpp:func:`hipMemAddressReserve`
|
||||
function. The size of the virtual memory must match the amount of physical
|
||||
memory previously allocated. You can then map the physical memory allocation to
|
||||
the newly-acquired virtual memory address range using the :cpp:func:`hipMemMap`
|
||||
function.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&ptr, padded_size, 0, 0, 0);
|
||||
hipMemMap(ptr, padded_size, 0, allocHandle, 0);
|
||||
|
||||
.. _set_memory_access:
|
||||
|
||||
Set memory access
|
||||
--------------------------------------------------------------------------------
|
||||
-----------------
|
||||
|
||||
Finally, use the :cpp:func:`hipMemSetAccess` function to enable memory access.
|
||||
It accepts the pointer to the virtual memory, the size, and a
|
||||
@@ -103,16 +139,39 @@ devices.
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAccessDesc accessDesc = {};
|
||||
accessDesc.location.type = HIP_MEM_LOCATION_TYPE_DEVICE;
|
||||
accessDesc.location.type = hipMemLocationTypeDevice;
|
||||
accessDesc.location.id = currentDev;
|
||||
accessDesc.flags = HIP_MEM_ACCESS_FLAGS_PROT_READWRITE;
|
||||
accessDesc.flags = hipMemAccessFlagsProtReadwrite;
|
||||
hipMemSetAccess(ptr, padded_size, &accessDesc, 1);
|
||||
|
||||
At this point the memory is allocated, mapped, and ready for use. You can read
|
||||
and write to it, just like you would a C style memory allocation.
|
||||
|
||||
.. _usage_virtual_memory:
|
||||
|
||||
Dynamically increase allocation size
|
||||
------------------------------------
|
||||
|
||||
To increase the amount of pre-allocated memory, use
|
||||
:cpp:func:`hipMemAddressReserve`, which accepts the starting address, and the
|
||||
size of the reservation in bytes. This allows you to have a continuous virtual
|
||||
address space without worrying about the underlying physical allocation.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&new_ptr, (new_size - padded_size), 0, ptr + padded_size, 0);
|
||||
hipMemMap(new_ptr, (new_size - padded_size), 0, newAllocHandle, 0);
|
||||
hipMemSetAccess(new_ptr, (new_size - padded_size), &accessDesc, 1);
|
||||
|
||||
The code sample above assumes that :cpp:func:`hipMemAddressReserve` was able to
|
||||
reserve the memory address at the specified location. However, this isn't
|
||||
guaranteed to be true, so you should validate that ``new_ptr`` points to a
|
||||
specific virtual address before using it.
|
||||
|
||||
.. _free_virtual_memory:
|
||||
|
||||
Free virtual memory
|
||||
--------------------------------------------------------------------------------
|
||||
-------------------
|
||||
|
||||
To free the memory allocated in this manner, use the corresponding free
|
||||
functions. To unmap the memory, use :cpp:func:`hipMemUnmap`. To release the
|
||||
@@ -128,27 +187,197 @@ synchronizes the device. This causes worse resource usage and performance.
|
||||
hipMemRelease(allocHandle);
|
||||
hipMemAddressFree(ptr, size);
|
||||
|
||||
.. _usage_virtual_memory:
|
||||
Example code
|
||||
============
|
||||
|
||||
Memory usage
|
||||
================================================================================
|
||||
The virtual memory management example follows these steps:
|
||||
|
||||
Dynamically increase allocation size
|
||||
--------------------------------------------------------------------------------
|
||||
1. Check virtual memory management :ref:`support <vmm_support>`:
|
||||
The :cpp:func:`hipDeviceGetAttribute` function is used to check the virtual
|
||||
memory management support of the GPU with ID 0.
|
||||
|
||||
The :cpp:func:`hipMemAddressReserve` function allows you to increase the amount
|
||||
of pre-allocated memory. This function accepts a parameter representing the
|
||||
requested starting address of the virtual memory. This allows you to have a
|
||||
continuous virtual address space without worrying about the underlying physical
|
||||
allocation.
|
||||
2. Physical memory :ref:`allocation <allocate_physical_memory>`: Physical memory
|
||||
is allocated using :cpp:func:`hipMemCreate` with pinned memory on the
|
||||
device.
|
||||
|
||||
3. Virtual memory :ref:`reservation <reserve_virtual_address>`: Virtual address
|
||||
range is reserved using :cpp:func:`hipMemAddressReserve`.
|
||||
|
||||
4. Mapping virtual address to physical memory: The physical memory is mapped
|
||||
to a virtual address (``virtualPointer``) using :cpp:func:`hipMemMap`.
|
||||
|
||||
5. Memory :ref:`access permissions<set_memory_access>`: Permission is set for
|
||||
pointer to allow read and write access using :cpp:func:`hipMemSetAccess`.
|
||||
|
||||
6. Memory operation: Data is written to the memory via ``virtualPointer``.
|
||||
|
||||
7. Launch kernels: The ``zeroAddr`` and ``fillAddr`` kernels are
|
||||
launched using the virtual memory pointer.
|
||||
|
||||
8. :ref:`Cleanup <free_virtual_memory>`: The mappings, physical memory, and
|
||||
virtual address are released at the end to avoid memory leaks.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&new_ptr, (new_size - padded_size), 0, ptr + padded_size, 0);
|
||||
hipMemMap(new_ptr, (new_size - padded_size), 0, newAllocHandle, 0);
|
||||
hipMemSetAccess(new_ptr, (new_size - padded_size), &accessDesc, 1);
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define ROUND_UP(SIZE,GRANULARITY) ((1 + SIZE / GRANULARITY) * GRANULARITY)
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void zeroAddr(int* pointer) {
|
||||
*pointer = 0;
|
||||
}
|
||||
|
||||
__global__ void fillAddr(int* pointer) {
|
||||
*pointer = 42;
|
||||
}
|
||||
|
||||
|
||||
int main() {
|
||||
|
||||
int currentDev = 0;
|
||||
|
||||
// Step 1: Check virtual memory management support on device 0
|
||||
int vmm = 0;
|
||||
HIP_CHECK(
|
||||
hipDeviceGetAttribute(
|
||||
&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, currentDev
|
||||
)
|
||||
);
|
||||
|
||||
std::cout << "Virtual memory management support value: " << vmm << std::endl;
|
||||
|
||||
if (vmm == 0) {
|
||||
std::cout << "GPU 0 doesn't support virtual memory management.";
|
||||
return 0;
|
||||
}
|
||||
|
||||
// Size of memory to allocate
|
||||
size_t size = 4 * 1024;
|
||||
|
||||
// Step 2: Allocate physical memory
|
||||
hipMemGenericAllocationHandle_t allocHandle;
|
||||
hipMemAllocationProp prop = {};
|
||||
prop.type = hipMemAllocationTypePinned;
|
||||
prop.location.type = hipMemLocationTypeDevice;
|
||||
prop.location.id = currentDev;
|
||||
size_t granularity = 0;
|
||||
HIP_CHECK(
|
||||
hipMemGetAllocationGranularity(
|
||||
&granularity,
|
||||
&prop,
|
||||
hipMemAllocationGranularityMinimum));
|
||||
size_t padded_size = ROUND_UP(size, granularity);
|
||||
HIP_CHECK(hipMemCreate(&allocHandle, padded_size * 2, &prop, 0));
|
||||
|
||||
// Step 3: Reserve a virtual memory address range
|
||||
void* virtualPointer = nullptr;
|
||||
HIP_CHECK(hipMemAddressReserve(&virtualPointer, padded_size, granularity, nullptr, 0));
|
||||
|
||||
// Step 4: Map the physical memory to the virtual address range
|
||||
HIP_CHECK(hipMemMap(virtualPointer, padded_size, 0, allocHandle, 0));
|
||||
|
||||
// Step 5: Set memory access permission for pointer
|
||||
hipMemAccessDesc accessDesc = {};
|
||||
accessDesc.location.type = hipMemLocationTypeDevice;
|
||||
accessDesc.location.id = currentDev;
|
||||
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
|
||||
|
||||
HIP_CHECK(hipMemSetAccess(virtualPointer, padded_size, &accessDesc, 1));
|
||||
|
||||
// Step 6: Perform memory operation
|
||||
int value = 42;
|
||||
HIP_CHECK(hipMemcpy(virtualPointer, &value, sizeof(int), hipMemcpyHostToDevice));
|
||||
|
||||
int result = 1;
|
||||
HIP_CHECK(hipMemcpy(&result, virtualPointer, sizeof(int), hipMemcpyDeviceToHost));
|
||||
if( result == 42) {
|
||||
std::cout << "Success. Value: " << result << std::endl;
|
||||
} else {
|
||||
std::cout << "Failure. Value: " << result << std::endl;
|
||||
}
|
||||
|
||||
// Step 7: Launch kernels
|
||||
// Launch zeroAddr kernel
|
||||
zeroAddr<<<1, 1>>>((int*)virtualPointer);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Check zeroAddr kernel result
|
||||
result = 1;
|
||||
HIP_CHECK(hipMemcpy(&result, virtualPointer, sizeof(int), hipMemcpyDeviceToHost));
|
||||
if( result == 0) {
|
||||
std::cout << "Success. zeroAddr kernel: " << result << std::endl;
|
||||
} else {
|
||||
std::cout << "Failure. zeroAddr kernel: " << result << std::endl;
|
||||
}
|
||||
|
||||
// Launch fillAddr kernel
|
||||
fillAddr<<<1, 1>>>((int*)virtualPointer);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Check fillAddr kernel result
|
||||
result = 1;
|
||||
HIP_CHECK(hipMemcpy(&result, virtualPointer, sizeof(int), hipMemcpyDeviceToHost));
|
||||
if( result == 42) {
|
||||
std::cout << "Success. fillAddr kernel: " << result << std::endl;
|
||||
} else {
|
||||
std::cout << "Failure. fillAddr kernel: " << result << std::endl;
|
||||
}
|
||||
|
||||
// Step 8: Cleanup
|
||||
HIP_CHECK(hipMemUnmap(virtualPointer, padded_size));
|
||||
HIP_CHECK(hipMemRelease(allocHandle));
|
||||
HIP_CHECK(hipMemAddressFree(virtualPointer, padded_size));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Virtual aliases
|
||||
================================================================================
|
||||
|
||||
Virtual aliases are multiple virtual memory addresses mapping to the same
|
||||
physical memory on the GPU. When this occurs, different threads, processes, or memory
|
||||
allocations to access shared physical memory through different virtual
|
||||
addresses on different devices.
|
||||
|
||||
Multiple virtual memory mappings can be created using multiple calls to
|
||||
:cpp:func:`hipMemMap` on the same memory allocation.
|
||||
|
||||
.. note::
|
||||
|
||||
RDNA cards may not produce correct results, if users access two different
|
||||
virtual addresses that map to the same physical address. In this case, the
|
||||
L1 data caches will be incoherent due to the virtual-to-physical aliasing.
|
||||
These GPUs will produce correct results if users access virtual-to-physical
|
||||
aliases using volatile pointers.
|
||||
|
||||
NVIDIA GPUs require special fences to produce correct results when
|
||||
using virtual aliases.
|
||||
|
||||
In the following code block, the kernels input device pointers are virtual
|
||||
aliases of the same memory allocation:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
__global__ void updateBoth(int* pointerA, int* pointerB) {
|
||||
// May produce incorrect results on RDNA and NVIDIA cards.
|
||||
*pointerA = 0;
|
||||
*pointerB = 42;
|
||||
}
|
||||
|
||||
__global__ void updateBoth_v2(volatile int* pointerA, volatile int* pointerB) {
|
||||
// May produce incorrect results on NVIDIA cards.
|
||||
*pointerA = 0;
|
||||
*pointerB = 42;
|
||||
}
|
||||
|
||||
The code sample above assumes that :cpp:func:`hipMemAddressReserve` was able to
|
||||
reserve the memory address at the specified location. However, this isn't
|
||||
guaranteed to be true, so you should validate that ``new_ptr`` points to a
|
||||
specific virtual address before using it.
|
||||
|
||||
In neuem Issue referenzieren
Einen Benutzer sperren