SWDEV-502480 - Update documentation from GitHub 2024-12-05
Change-Id: I179814351b77935aff55e8ae47dd322a3e15a868
This commit is contained in:
Istvan Kiss
@@ -0,0 +1,129 @@
|
||||
.. meta::
|
||||
:description: This page describes call stack concept in HIP
|
||||
:keywords: AMD, ROCm, HIP, call stack
|
||||
|
||||
*******************************************************************************
|
||||
Call stack
|
||||
*******************************************************************************
|
||||
|
||||
The call stack is a data structure for managing function calls, by saving the
|
||||
state of the current function. Each time a function is called, a new call frame
|
||||
is added to the top of the stack, containing information such as local
|
||||
variables, return addresses and function parameters. When the function
|
||||
execution completes, the frame is removed from the stack and loaded back into
|
||||
the corresponding registers. This concept allows the program to return to the
|
||||
calling function and continue execution from where it left off.
|
||||
|
||||
The call stack for each thread must track its function calls, local variables,
|
||||
and return addresses. However, in GPU programming, the memory required to store
|
||||
the call stack increases due to the parallelism inherent to the GPUs. NVIDIA
|
||||
and AMD GPUs use different approaches. NVIDIA GPUs have the independent thread
|
||||
scheduling feature where each thread has its own call stack and effective
|
||||
program counter. On AMD GPUs threads are grouped; each warp has its own call
|
||||
stack and program counter. Warps are described and explained in the
|
||||
:ref:`inherent_thread_hierarchy`
|
||||
|
||||
If a thread or warp exceeds its stack size, a stack overflow occurs, causing
|
||||
kernel failure. This can be detected using debuggers.
|
||||
|
||||
Call stack management with HIP
|
||||
===============================================================================
|
||||
|
||||
You can adjust the call stack size as shown in the following example, allowing
|
||||
fine-tuning based on specific kernel requirements. This helps prevent stack
|
||||
overflow errors by ensuring sufficient stack memory is allocated.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
} \
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
size_t stackSize;
|
||||
HIP_CHECK(hipDeviceGetLimit(&stackSize, hipLimitStackSize));
|
||||
std::cout << "Default stack size: " << stackSize << " bytes" << std::endl;
|
||||
|
||||
// Set a new stack size
|
||||
size_t newStackSize = 1024 * 8; // 8 KiB
|
||||
HIP_CHECK(hipDeviceSetLimit(hipLimitStackSize, newStackSize));
|
||||
|
||||
HIP_CHECK(hipDeviceGetLimit(&stackSize, hipLimitStackSize));
|
||||
std::cout << "Updated stack size: " << stackSize << " bytes" << std::endl;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Depending on the GPU model, at full occupancy, it can consume a significant
|
||||
amount of memory. For instance, an MI300X with 304 compute units (CU) and up to
|
||||
2048 threads per CU could use 304 · 2048 · 1024 bytes = 608 MiB for the call
|
||||
stack by default.
|
||||
|
||||
Handling recursion and deep function calls
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Similar to CPU programming, recursive functions and deeply nested function
|
||||
calls are supported. However, developers must ensure that these functions do
|
||||
not exceed the available stack memory, considering the huge amount of memory
|
||||
needed for the call stack due to the GPUs inherent parallelism. This can be
|
||||
achieved by increasing stack size or optimizing code to reduce stack usage. To
|
||||
detect stack overflow add proper error handling or use debugging tools.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
} \
|
||||
}
|
||||
|
||||
__device__ unsigned long long fibonacci(unsigned long long n)
|
||||
{
|
||||
if (n == 0 || n == 1)
|
||||
{
|
||||
return n;
|
||||
}
|
||||
return fibonacci(n - 1) + fibonacci(n - 2);
|
||||
}
|
||||
|
||||
__global__ void kernel(unsigned long long n)
|
||||
{
|
||||
unsigned long long result = fibonacci(n);
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
|
||||
if (x == 0)
|
||||
printf("%llu! = %llu \n", n, result);
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
kernel<<<1, 1>>>(10);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// With -O0 optimization option hit the stack limit
|
||||
// kernel<<<1, 256>>>(2048);
|
||||
// HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,502 @@
|
||||
.. meta::
|
||||
:description: This topic describes how to use cooperative groups in HIP
|
||||
:keywords: AMD, ROCm, HIP, cooperative groups
|
||||
|
||||
.. _cooperative_groups_how-to:
|
||||
|
||||
*******************************************************************************
|
||||
Cooperative groups
|
||||
*******************************************************************************
|
||||
|
||||
The cooperative groups API is an extension to the HIP programming model, which
|
||||
provides developers with a flexible, dynamic grouping mechanism for the
|
||||
communicating threads. Cooperative groups let you define your own set of thread
|
||||
groups which may fit your use-cases better than those defined by the hardware.
|
||||
This lets you specify the level of granularity for thread communication which
|
||||
can lead to more efficient parallel decompositions.
|
||||
|
||||
The API is accessible in the ``cooperative_groups`` namespace after the
|
||||
``hip_cooperative_groups.h`` header is included. The header contains the following
|
||||
elements:
|
||||
|
||||
* Static functions to create groups and subgroups.
|
||||
* Hardware-accelerated operations over the whole group, like shuffles.
|
||||
* Data types of cooperative groups.
|
||||
* Synchronize member function of the groups.
|
||||
* Get group properties member functions.
|
||||
|
||||
Cooperative groups thread model
|
||||
================================================================================
|
||||
|
||||
The thread hierarchy abstractions of cooperative groups are depicted in the following figures: :ref:`grid hierarchy <coop_thread_top_hierarchy>` and :ref:`block hierarchy <coop_thread_bottom_hierarchy>`.
|
||||
|
||||
.. _coop_thread_top_hierarchy:
|
||||
|
||||
.. figure:: ../../data/how-to/hip_runtime_api/cooperative_groups/thread_hierarchy_coop_top.svg
|
||||
:alt: Diagram depicting nested rectangles of varying color. The outermost one
|
||||
titled "Grid", inside sets of different sized rectangles layered on
|
||||
one another titled "Block". Each "Block" containing sets of uniform
|
||||
rectangles layered on one another titled "Warp". Each of the "Warp"
|
||||
titled rectangles filled with downward pointing arrows inside.
|
||||
|
||||
Cooperative group thread hierarchy in grids.
|
||||
|
||||
The **multi grid** is an abstraction of potentially multiple simultaneous
|
||||
launches of the same kernel over multiple devices. The **grid** in cooperative
|
||||
groups is a single dispatch of kernels for execution like the original grid.
|
||||
|
||||
.. note::
|
||||
|
||||
* The ability to synchronize over a grid or multi grid requires the kernel to
|
||||
be launched using the specific cooperative groups API.
|
||||
|
||||
* Multi grid deprecated since ROCm 5.0.
|
||||
|
||||
The **block** is the same as the :ref:`inherent_thread_model` block entity.
|
||||
|
||||
.. note::
|
||||
|
||||
Explicit warp-level thread handling is absent from the Cooperative groups API. In order to exploit the known hardware SIMD width on which built-in functionality translates to simpler logic, you can use the group partitioning part of the API, such as ``tiled_partition``.
|
||||
|
||||
.. _coop_thread_bottom_hierarchy:
|
||||
|
||||
.. figure:: ../../data/how-to/hip_runtime_api/cooperative_groups/thread_hierarchy_coop_bottom.svg
|
||||
:alt: The new level between block thread and threads.
|
||||
|
||||
Cooperative group thread hierarchy in blocks.
|
||||
|
||||
The cooperative groups API introduce a new level between block thread and threads. The :ref:`thread-block tile <coop_thread_block_tile>` give the opportunity to have tiles in the thread block, while the :ref:`coalesced group <coop_coalesced_groups>` holds the active threads of the parent group. These groups further discussed in the :ref:`groups types <coop_group_types>` section.
|
||||
|
||||
For details on memory model, check the :ref:`memory model description <memory_hierarchy>`.
|
||||
|
||||
.. _coop_group_types:
|
||||
|
||||
Group types
|
||||
===========
|
||||
|
||||
Group types are based on the levels of synchronization and data sharing among threads.
|
||||
|
||||
Thread-block group
|
||||
------------------
|
||||
|
||||
Represents an intra-block cooperative groups type where the participating threads within the group are the same threads that participated in the currently executing ``block``.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
class thread_block;
|
||||
|
||||
Constructed via:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
thread_block g = this_thread_block();
|
||||
|
||||
The ``group_index()`` , ``thread_index()`` , ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()`` , ``sync()`` and ``group_dim()`` member functions are public of the thread_block class. For further details, check the :ref:`thread_block references <thread_block_ref>` .
|
||||
|
||||
Grid group
|
||||
------------
|
||||
|
||||
Represents an inter-block cooperative groups type where the group's participating threads span multiple blocks running the same kernel on the same device. Use the cooperative launch API to enable synchronization across the grid group.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
class grid_group;
|
||||
|
||||
Constructed via:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
grid_group g = this_grid();
|
||||
|
||||
The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()`` and ``sync()`` member functions
|
||||
are public of the ``grid_group`` class. For further details, check the :ref:`grid_group references <grid_group_ref>`.
|
||||
|
||||
Multi-grid group
|
||||
------------------
|
||||
|
||||
Represents an inter-device cooperative groups type where the participating threads within the group span multiple devices that run the same kernel on the devices. Use the cooperative launch API to enable synchronization across the multi-grid group.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
class multi_grid_group;
|
||||
|
||||
Constructed via:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
// Kernel must be launched with the cooperative multi-device API
|
||||
multi_grid_group g = this_multi_grid();
|
||||
|
||||
The ``num_grids()`` , ``grid_rank()`` , ``thread_rank()``, ``size()``, ``cg_type()``, ``is_valid()`` ,
|
||||
and ``sync()`` member functions are public of the ``multi_grid_group`` class. For
|
||||
further details check the :ref:`multi_grid_group references <multi_grid_group_ref>` .
|
||||
|
||||
.. _coop_thread_block_tile:
|
||||
|
||||
Thread-block tile
|
||||
------------------
|
||||
|
||||
This constructs a templated class derived from ``thread_group``. The template defines the tile
|
||||
size of the new thread group at compile time. This group type also supports sub-wave level intrinsics.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
template <unsigned int Size, typename ParentT = void>
|
||||
class thread_block_tile;
|
||||
|
||||
Constructed via:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
template <unsigned int Size, typename ParentT>
|
||||
_CG_QUALIFIER thread_block_tile<Size, ParentT> tiled_partition(const ParentT& g)
|
||||
|
||||
|
||||
.. note::
|
||||
|
||||
* Size must be a power of 2 and not larger than warp (wavefront) size.
|
||||
* ``shfl()`` functions support integer or float type.
|
||||
|
||||
The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()``, ``sync()``, ``meta_group_rank()``, ``meta_group_size()``, ``shfl()``, ``shfl_down()``, ``shfl_up()``, ``shfl_xor()``, ``ballot()``, ``any()``, ``all()``, ``match_any()`` and ``match_all()`` member functions are public of the ``thread_block_tile`` class. For further details, check the :ref:`thread_block_tile references <thread_block_tile_ref>` .
|
||||
|
||||
.. _coop_coalesced_groups:
|
||||
|
||||
Coalesced groups
|
||||
------------------
|
||||
|
||||
Threads (64 threads on CDNA and 32 threads on RDNA) in a warp cannot execute different instructions simultaneously, so conditional branches are executed serially within the warp. When threads encounter a conditional branch, they can diverge, resulting in some threads being disabled, if they do not meet the condition to execute that branch. The active threads referred as coalesced, and coalesced group represents an active thread group within a warp.
|
||||
|
||||
.. note::
|
||||
|
||||
The NVIDIA GPU's independent thread scheduling presents the appearance that threads on different branches execute concurrently.
|
||||
|
||||
.. warning::
|
||||
|
||||
AMD GPUs do not support independent thread scheduling. Some CUDA application can rely on this feature and the ported HIP version on AMD GPUs can deadlock, when they try to make use of independent thread scheduling.
|
||||
|
||||
This group type also supports sub-wave level intrinsics.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
class coalesced_group;
|
||||
|
||||
Constructed via:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
coalesced_group active = coalesced_threads();
|
||||
|
||||
.. note::
|
||||
|
||||
``shfl()`` functions support integer or float type.
|
||||
|
||||
The ``thread_rank()`` , ``size()``, ``cg_type()``, ``is_valid()``, ``sync()``, ``meta_group_rank()``, ``meta_group_size()``, ``shfl()``, ``shfl_down()``, ``shfl_up()``, ``ballot()``, ``any()``, ``all()``, ``match_any()`` and ``match_all()`` member functions are public of the ``coalesced_group`` class. For more information, see :ref:`coalesced_group references <coalesced_group_ref>` .
|
||||
|
||||
Cooperative groups simple example
|
||||
=================================
|
||||
|
||||
The difference to the original block model in the ``reduce_sum`` device function is the following.
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: Original Block
|
||||
:sync: original-block
|
||||
|
||||
.. code-block:: cuda
|
||||
|
||||
__device__ int reduce_sum(int *shared, int val) {
|
||||
|
||||
// Thread ID
|
||||
const unsigned int thread_id = threadIdx.x;
|
||||
|
||||
// Every iteration the number of active threads
|
||||
// halves, until we processed all values
|
||||
for(unsigned int i = blockDim.x / 2; i > 0; i /= 2) {
|
||||
// Store value in shared memory with thread ID
|
||||
shared[thread_id] = val;
|
||||
|
||||
// Synchronize all threads
|
||||
__syncthreads();
|
||||
|
||||
// Active thread sum up
|
||||
if(thread_id < i)
|
||||
val += shared[thread_id + i];
|
||||
|
||||
// Synchronize all threads in the group
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
// ...
|
||||
}
|
||||
|
||||
.. tab-item:: Cooperative groups
|
||||
:sync: cooperative-groups
|
||||
|
||||
.. code-block:: cuda
|
||||
|
||||
__device__ int reduce_sum(thread_group g,
|
||||
int *shared,
|
||||
int val) {
|
||||
|
||||
// Thread ID
|
||||
const unsigned int group_thread_id = g.thread_rank();
|
||||
|
||||
// Every iteration the number of active threads
|
||||
// halves, until we processed all values
|
||||
for(unsigned int i = g.size() / 2; i > 0; i /= 2) {
|
||||
// Store value in shared memroy with thread ID
|
||||
shared[group_thread_id] = val;
|
||||
|
||||
// Synchronize all threads in the group
|
||||
g.sync();
|
||||
|
||||
// Active thread sum up
|
||||
if(group_thread_id < i)
|
||||
val += shared[group_thread_id + i];
|
||||
|
||||
// Synchronize all threads in the group
|
||||
g.sync();
|
||||
}
|
||||
|
||||
// ...
|
||||
}
|
||||
|
||||
The ``reduce_sum()`` function call and input data initialization difference to the original block model is the following.
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: Original Block
|
||||
:sync: original-block
|
||||
|
||||
.. code-block:: cuda
|
||||
|
||||
__global__ void sum_kernel(...) {
|
||||
|
||||
// ...
|
||||
|
||||
// Workspace array in shared memory
|
||||
__shared__ unsigned int workspace[2048];
|
||||
|
||||
// ...
|
||||
|
||||
// Perform reduction
|
||||
output = reduce_sum(workspace, input);
|
||||
|
||||
// ...
|
||||
}
|
||||
|
||||
.. tab-item:: Cooperative groups
|
||||
:sync: cooperative-groups
|
||||
|
||||
.. code-block:: cuda
|
||||
|
||||
__global__ void sum_kernel(...) {
|
||||
|
||||
// ...
|
||||
|
||||
// Workspace array in shared memory
|
||||
__shared__ unsigned int workspace[2048];
|
||||
|
||||
// ...
|
||||
|
||||
// Initialize the thread_block
|
||||
thread_block thread_block_group = this_thread_block();
|
||||
// Perform reduction
|
||||
output = reduce_sum(thread_block_group, workspace, input);
|
||||
|
||||
// ...
|
||||
}
|
||||
|
||||
At the device function, the input group type is the ``thread_group``, which is the parent class of all the cooperative groups type. With this, you can write generic functions, which can work with any type of cooperative groups.
|
||||
|
||||
.. _coop_synchronization:
|
||||
|
||||
Synchronization
|
||||
===============
|
||||
|
||||
With each group type, the synchronization requires using the correct cooperative groups launch API.
|
||||
|
||||
**Check the kernel launch capability**
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: Thread-block
|
||||
:sync: thread-block
|
||||
|
||||
Do not need kernel launch validation.
|
||||
|
||||
.. tab-item:: Grid
|
||||
:sync: grid
|
||||
|
||||
Confirm the cooperative launch capability on the single AMD GPU:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
int device = 0;
|
||||
int supports_coop_launch = 0;
|
||||
// Check support
|
||||
// Use hipDeviceAttributeCooperativeMultiDeviceLaunch when launching across multiple devices
|
||||
HIP_CHECK(hipGetDevice(&device));
|
||||
HIP_CHECK(
|
||||
hipDeviceGetAttribute(&supports_coop_launch, hipDeviceAttributeCooperativeLaunch, device));
|
||||
if(!supports_coop_launch)
|
||||
{
|
||||
std::cout << "Skipping, device " << device << " does not support cooperative groups"
|
||||
<< std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: Multi-grid
|
||||
:sync: multi-grid
|
||||
|
||||
Confirm the cooperative launch capability over multiple GPUs:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
// Check support of cooperative groups
|
||||
std::vector<int> deviceIDs;
|
||||
for(int deviceID = 0; deviceID < device_count; deviceID++) {
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
int supports_coop_launch = 0;
|
||||
HIP_CHECK(
|
||||
hipDeviceGetAttribute(
|
||||
&supports_coop_launch,
|
||||
hipDeviceAttributeCooperativeMultiDeviceLaunch,
|
||||
deviceID));
|
||||
if(!supports_coop_launch) {
|
||||
std::cout << "Skipping, device " << deviceID << " does not support cooperative groups"
|
||||
<< std::endl;
|
||||
}
|
||||
else
|
||||
#endif
|
||||
{
|
||||
std::cout << deviceID << std::endl;
|
||||
// Collect valid deviceIDs.
|
||||
deviceIDs.push_back(deviceID);
|
||||
}
|
||||
}
|
||||
|
||||
**Kernel launch**
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: Thread-block
|
||||
:sync: thread-block
|
||||
|
||||
You can access the new block representation using the original kernel launch methods.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
void* params[] = {&d_vector, &d_block_reduced, &d_partition_reduced};
|
||||
// Launching kernel from host.
|
||||
HIP_CHECK(hipLaunchKernelGGL(vector_reduce_kernel<partition_size>,
|
||||
dim3(num_blocks),
|
||||
dim3(threads_per_block),
|
||||
0,
|
||||
hipStreamDefault,
|
||||
&d_vector,
|
||||
&d_block_reduced,
|
||||
&d_partition_reduced));
|
||||
|
||||
.. tab-item:: Grid
|
||||
:sync: grid
|
||||
|
||||
Launch the cooperative kernel on a single GPU:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
void* params[] = {};
|
||||
// Launching kernel from host.
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(vector_reduce_kernel<partition_size>,
|
||||
dim3(num_blocks),
|
||||
dim3(threads_per_block),
|
||||
0,
|
||||
0,
|
||||
hipStreamDefault));
|
||||
|
||||
.. tab-item:: Multi-grid
|
||||
:sync: multi-grid
|
||||
|
||||
Launch the cooperative kernel over multiple GPUs:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipLaunchParams *launchParamsList = (hipLaunchParams*)malloc(sizeof(hipLaunchParams) * deviceIDs.size());
|
||||
for(int deviceID : deviceIDs) {
|
||||
|
||||
// Set device
|
||||
HIP_CHECK(hipSetDevice(deviceID));
|
||||
|
||||
// Create stream
|
||||
hipStream_t stream;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
|
||||
// Parameters
|
||||
void* params[] = {&(d_vector[deviceID]), &(d_block_reduced[deviceID]), &(d_partition_reduced[deviceID])};
|
||||
|
||||
// Set launchParams
|
||||
launchParamsList[deviceID].func = (void*)vector_reduce_kernel<partition_size>;
|
||||
launchParamsList[deviceID].gridDim = dim3(1);
|
||||
launchParamsList[deviceID].blockDim = dim3(threads_per_block);
|
||||
launchParamsList[deviceID].sharedMem = 0;
|
||||
launchParamsList[deviceID].stream = stream;
|
||||
launchParamsList[deviceID].args = params;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList,
|
||||
(int)deviceIDs.size(),
|
||||
hipCooperativeLaunchMultiDeviceNoPreSync));
|
||||
|
||||
**Device side synchronization**
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: Thread-block
|
||||
:sync: thread-block
|
||||
|
||||
The device side code of the thread_block synchronization over single GPUs:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
thread_block g = this_thread_block();
|
||||
g.sync();
|
||||
|
||||
.. tab-item:: Grid
|
||||
:sync: grid
|
||||
|
||||
The device side code of the grid synchronization over single GPUs:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
grid_group grid = this_grid();
|
||||
grid.sync();
|
||||
|
||||
.. tab-item:: Multi-grid
|
||||
:sync: multi-grid
|
||||
|
||||
The device side code of the multi-grid synchronization over multiple GPUs:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
multi_grid_group multi_grid = this_multi_grid();
|
||||
multi_grid.sync();
|
||||
|
||||
Unsupported NVIDIA CUDA features
|
||||
================================
|
||||
|
||||
HIP doesn't support the following NVIDIA CUDA optional headers:
|
||||
|
||||
* ``cooperative_groups/memcpy_async.h``
|
||||
* ``cooperative_groups/reduce.h``
|
||||
* ``cooperative_groups/scan.h``
|
||||
|
||||
HIP doesn't support the following CUDA class in ``cooperative_groups`` namespace:
|
||||
|
||||
* ``cluster_group``
|
||||
|
||||
HIP doesn't support the following CUDA functions/operators in ``cooperative_groups`` namespace:
|
||||
|
||||
* ``synchronize``
|
||||
* ``memcpy_async``
|
||||
* ``wait`` and ``wait_prior``
|
||||
* ``barrier_arrive`` and ``barrier_wait``
|
||||
* ``invoke_one`` and ``invoke_one_broadcast``
|
||||
* ``reduce``
|
||||
* ``reduce_update_async`` and ``reduce_store_async``
|
||||
* Reduce operators ``plus`` , ``less`` , ``greater`` , ``bit_and`` , ``bit_xor`` and ``bit_or``
|
||||
* ``inclusive_scan`` and ``exclusive_scan``
|
||||
@@ -0,0 +1,136 @@
|
||||
.. meta::
|
||||
:description: Error Handling
|
||||
:keywords: AMD, ROCm, HIP, error handling, error
|
||||
|
||||
.. _error_handling:
|
||||
|
||||
********************************************************************************
|
||||
Error handling
|
||||
********************************************************************************
|
||||
|
||||
HIP provides functionality to detect, report, and manage errors that occur
|
||||
during the execution of HIP runtime functions or when launching kernels. Every
|
||||
HIP runtime function, apart from launching kernels, has :cpp:type:`hipError_t`
|
||||
as return type. :cpp:func:`hipGetLastError` and :cpp:func:`hipPeekAtLastError`
|
||||
can be used for catching errors from kernel launches, as kernel launches don't
|
||||
return an error directly. HIP maintains an internal state, that includes the
|
||||
last error code. :cpp:func:`hipGetLastError` returns and resets that error to
|
||||
``hipSuccess``, while :cpp:func:`hipPeekAtLastError` just returns the error
|
||||
without changing it. To get a human readable version of the errors,
|
||||
:cpp:func:`hipGetErrorString` and :cpp:func:`hipGetErrorName` can be used.
|
||||
|
||||
.. note::
|
||||
|
||||
:cpp:func:`hipGetLastError` returns the returned error code of the last HIP
|
||||
runtime API call even if it's ``hipSuccess``, while ``cudaGetLastError``
|
||||
returns the error returned by any of the preceding CUDA APIs in the same
|
||||
host thread. :cpp:func:`hipGetLastError` behavior will be matched with
|
||||
``cudaGetLastError`` in ROCm release 7.0.
|
||||
|
||||
Best practices of HIP error handling:
|
||||
|
||||
1. Check errors after each API call - Avoid error propagation.
|
||||
2. Use macros for error checking - Check :ref:`hip_check_macros`.
|
||||
3. Handle errors gracefully - Free resources and provide meaningful error
|
||||
messages to the user.
|
||||
|
||||
For more details on the error handling functions, see :ref:`error handling
|
||||
functions reference page <error_handling_reference>`.
|
||||
|
||||
.. _hip_check_macros:
|
||||
|
||||
HIP check macros
|
||||
================================================================================
|
||||
|
||||
HIP uses check macros to simplify error checking and reduce code duplication.
|
||||
The ``HIP_CHECK`` macros are mainly used to detect and report errors. It can
|
||||
also exit from application with ``exit(1);`` function call after the error
|
||||
print. The ``HIP_CHECK`` macro example:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t status = expression; \
|
||||
if(status != hipSuccess){ \
|
||||
std::cerr << "HIP error " \
|
||||
<< status << ": " \
|
||||
<< hipGetErrorString(status) \
|
||||
<< " at " << __FILE__ << ":" \
|
||||
<< __LINE__ << std::endl; \
|
||||
} \
|
||||
}
|
||||
|
||||
Complete example
|
||||
================================================================================
|
||||
|
||||
A complete example to demonstrate the error handling with a simple addition of
|
||||
two values kernel:
|
||||
|
||||
.. 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; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c, size_t size) {
|
||||
const size_t index = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(index < size) {
|
||||
c[index] += a[index] + b[index];
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
constexpr int numOfBlocks = 256;
|
||||
constexpr int threadsPerBlock = 256;
|
||||
constexpr size_t arraySize = 1U << 16;
|
||||
|
||||
std::vector<int> a(arraySize), b(arraySize), c(arraySize);
|
||||
int *d_a, *d_b, *d_c;
|
||||
|
||||
// Setup input values.
|
||||
std::fill(a.begin(), a.end(), 1);
|
||||
std::fill(b.begin(), b.end(), 2);
|
||||
|
||||
// Allocate device copies of a, b and c.
|
||||
HIP_CHECK(hipMalloc(&d_a, arraySize * sizeof(*d_a)));
|
||||
HIP_CHECK(hipMalloc(&d_b, arraySize * sizeof(*d_b)));
|
||||
HIP_CHECK(hipMalloc(&d_c, arraySize * sizeof(*d_c)));
|
||||
|
||||
// Copy input values to device.
|
||||
HIP_CHECK(hipMemcpy(d_a, &a, arraySize * sizeof(*d_a), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(d_b, &b, arraySize * sizeof(*d_b), hipMemcpyHostToDevice));
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(numOfBlocks), dim3(threadsPerBlock), 0, 0, d_a, d_b, d_c, arraySize);
|
||||
// Check the kernel launch
|
||||
HIP_CHECK(hipGetLastError());
|
||||
// Check for kernel execution error
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Copy the result back to the host.
|
||||
HIP_CHECK(hipMemcpy(&c, d_c, arraySize * sizeof(*d_c), hipMemcpyDeviceToHost));
|
||||
|
||||
// Cleanup allocated memory.
|
||||
HIP_CHECK(hipFree(d_a));
|
||||
HIP_CHECK(hipFree(d_b));
|
||||
HIP_CHECK(hipFree(d_c));
|
||||
|
||||
// Print the result.
|
||||
std::cout << a[0] << " + " << b[0] << " = " << c[0] << std::endl;
|
||||
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,140 @@
|
||||
.. meta::
|
||||
:description: HIP provides an external resource interoperability API that
|
||||
allows efficient data sharing between HIP's computing power and
|
||||
OpenGL's graphics rendering.
|
||||
:keywords: AMD, ROCm, HIP, external, interop, interoperability
|
||||
|
||||
*******************************************************************************
|
||||
External resource interoperability
|
||||
*******************************************************************************
|
||||
|
||||
This feature allows HIP to work with resources -- like memory and semaphores --
|
||||
created by other APIs. This means resources can be used from APIs like CUDA,
|
||||
OpenCL and Vulkan within HIP, making it easier to integrate HIP into existing
|
||||
projects.
|
||||
|
||||
To use external resources in HIP, you typically follow these steps:
|
||||
|
||||
- Import resources from other APIs using HIP provided functions
|
||||
- Use external resources as if they were created in HIP
|
||||
- Destroy the HIP resource object to clean up
|
||||
|
||||
Semaphore Functions
|
||||
===============================================================================
|
||||
|
||||
Semaphore functions are essential for synchronization in parallel computing.
|
||||
These functions facilitate communication and coordination between different
|
||||
parts of a program or between different programs. By managing semaphores, tasks
|
||||
are executed in the correct order, and resources are utilized effectively.
|
||||
Semaphore functions ensure smooth operation, preventing conflicts and
|
||||
maintaining the integrity of processes; upholding the integrity and performance
|
||||
of concurrent processes.
|
||||
|
||||
External semaphore functions can be used in HIP as described in :ref:`external_resource_interoperability_reference`.
|
||||
|
||||
Memory Functions
|
||||
===============================================================================
|
||||
|
||||
HIP external memory functions focus on the efficient sharing and management of
|
||||
memory resources. These functions enable importing memory created by external
|
||||
systems, enabling the HIP program to use this memory seamlessly. Memory
|
||||
functions include mapping memory for effective use and ensuring proper cleanup
|
||||
to prevent resource leaks. This is critical for performance, particularly in
|
||||
applications handling large datasets or complex structures such as textures in
|
||||
graphics. Proper memory management ensures stability and efficient resource
|
||||
utilization.
|
||||
|
||||
Example
|
||||
===============================================================================
|
||||
|
||||
ROCm examples include a
|
||||
`HIP--Vulkan interoperation example <https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/vulkan_interop>`_
|
||||
demonstrates how to perform interoperation between HIP and Vulkan.
|
||||
|
||||
In this example, a simple HIP kernel is used to compute a sine wave, which is
|
||||
then rendered to a window as a graphical output using Vulkan. The process
|
||||
requires several initialization steps, such as setting up a HIP context,
|
||||
creating a Vulkan instance, and configuring the GPU device and queue. After
|
||||
these initial steps, the kernel executes the sine wave computation, and Vulkan
|
||||
continuously updates the window framebuffer to display the computed data until
|
||||
the window is closed.
|
||||
|
||||
The following code converts a Vulkan memory handle to its equivalent HIP
|
||||
handle. The input ``VkDeviceMemory`` and the created HIP memory represents the
|
||||
same physical area of GPU memory, through the handles of each respective API.
|
||||
Writing to the buffer in one API will allow us to read the results through the
|
||||
other. Note that access to the buffer should be synchronized between the APIs,
|
||||
for example using queue syncs or semaphores.
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/external_interop.hip
|
||||
:start-after: // [Sphinx vulkan memory to hip start]
|
||||
:end-before: // [Sphinx vulkan memory to hip end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
The Vulkan semaphore is converted to HIP semaphore shown in the following
|
||||
example. Signaling on the semaphore in one API will allow the other API to wait
|
||||
on it, which is how we can guarantee synchronized access to resources in a
|
||||
cross-API manner.
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/external_interop.hip
|
||||
:start-after: // [Sphinx semaphore import start]
|
||||
:end-before: // [Sphinx semaphore import end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
When the HIP external memory is exported from Vulkan and imported to HIP, it is
|
||||
not yet ready for use. The Vulkan handle is shared, allowing for memory sharing
|
||||
rather than copying during the export process. To actually use the memory, we
|
||||
need to map it to a pointer so that we may pass it to the kernel so that it can
|
||||
be read from and written to. The external memory map to HIP in the following
|
||||
example:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/external_interop.hip
|
||||
:start-after: // [Sphinx map external memory start]
|
||||
:end-before: // [Sphinx map external memory end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
Wait for buffer is ready and not under modification at Vulkan side:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/external_interop.hip
|
||||
:start-after: // [Sphinx wait semaphore start]
|
||||
:end-before: // [Sphinx wait semaphore end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
The sinewave kernel implementation:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/external_interop.hip
|
||||
:start-after: [Sphinx sinewave kernel start]
|
||||
:end-before: // [Sphinx sinewave kernel end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
Signal to Vulkan that we are done with the buffer and that it can proceed with
|
||||
rendering:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/external_interop.hip
|
||||
:start-after: // [Sphinx signal semaphore start]
|
||||
:end-before: // [Sphinx signal semaphore end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
@@ -0,0 +1,516 @@
|
||||
.. meta::
|
||||
:description: This chapter describes how to use HIP graphs and highlights their use cases.
|
||||
:keywords: ROCm, HIP, graph, stream
|
||||
|
||||
.. _how_to_HIP_graph:
|
||||
|
||||
********************************************************************************
|
||||
HIP graphs
|
||||
********************************************************************************
|
||||
|
||||
.. note::
|
||||
The HIP graph API is currently in Beta. Some features can change and might
|
||||
have outstanding issues. Not all features supported by CUDA graphs are yet
|
||||
supported. For a list of all currently supported functions see the
|
||||
:ref:`HIP graph API documentation<graph_management_reference>`.
|
||||
|
||||
HIP graphs are an alternative way of executing tasks on a GPU that can provide
|
||||
performance benefits over launching kernels using the standard
|
||||
method via streams. A HIP graph is made up of nodes and edges. The nodes of a
|
||||
HIP graph represent the operations performed, while the edges mark dependencies
|
||||
between those operations.
|
||||
|
||||
The nodes can be one of the following:
|
||||
|
||||
- empty nodes
|
||||
- nested graphs
|
||||
- kernel launches
|
||||
- host-side function calls
|
||||
- HIP memory functions (copy, memset, ...)
|
||||
- HIP events
|
||||
- signalling or waiting on external semaphores
|
||||
|
||||
.. note::
|
||||
The available node types are specified by :cpp:enum:`hipGraphNodeType`.
|
||||
|
||||
The following figure visualizes the concept of graphs, compared to using streams.
|
||||
|
||||
.. figure:: ../../data/how-to/hip_runtime_api/hipgraph/hip_graph.svg
|
||||
:alt: Diagram depicting the difference between using streams to execute
|
||||
kernels with dependencies, resolved by explicitly synchronizing,
|
||||
or using graphs, where the edges denote the dependencies.
|
||||
|
||||
The standard method of launching kernels incurs a small overhead for each
|
||||
iteration of the operation involved. That overhead is negligible, when the
|
||||
kernel is launched directly with the HIP C/C++ API, but depending on the
|
||||
framework used, there can be several levels of redirection, until the actual
|
||||
kernel is launched by the HIP runtime, leading to significant overhead.
|
||||
Especially for some AI frameworks, a GPU kernel might run faster than the time
|
||||
it takes for the framework to set up and launch the kernel, and so the overhead
|
||||
of repeatedly launching kernels can have a significant impact on performance.
|
||||
|
||||
HIP graphs are designed to address this issue, by predefining the HIP API calls
|
||||
and their dependencies with a graph, and performing most of the initialization
|
||||
beforehand. Launching a graph only requires a single call, after which the
|
||||
HIP runtime takes care of executing the operations within the graph.
|
||||
Graphs can provide additional performance benefits, by enabling optimizations
|
||||
that are only possible when knowing the dependencies between the operations.
|
||||
|
||||
.. figure:: ../../data/how-to/hip_runtime_api/hipgraph/hip_graph_speedup.svg
|
||||
:alt: Diagram depicting the speed up achievable with HIP graphs compared to
|
||||
HIP streams when launching many short-running kernels.
|
||||
|
||||
Qualitative presentation of the execution time of many short-running kernels
|
||||
when launched using HIP stream versus HIP graph. This does not include the
|
||||
time needed to set up the graph.
|
||||
|
||||
Using HIP graphs
|
||||
================================================================================
|
||||
|
||||
There are two different ways of creating graphs: Capturing kernel launches from
|
||||
a stream, or explicitly creating graphs. The difference between the two
|
||||
approaches is explained later in this chapter.
|
||||
|
||||
The general flow for using HIP graphs includes the following steps.
|
||||
|
||||
#. Create a :cpp:type:`hipGraph_t` graph template using one of the two approaches described in this chapter
|
||||
#. Create a :cpp:type:`hipGraphExec_t` executable instance of the graph template using :cpp:func:`hipGraphInstantiate`
|
||||
#. Use :cpp:func:`hipGraphLaunch` to launch the executable graph to a stream
|
||||
#. After execution completes free and destroy graph resources
|
||||
|
||||
The first two steps are the initial setup and only need to be executed once. First
|
||||
step is the definition of the operations (nodes) and the dependencies (edges)
|
||||
between them. The second step is the instantiation of the graph. This takes care
|
||||
of validating and initializing the graph, to reduce the overhead when executing
|
||||
the graph. The third step is the execution of the graph, which takes care of
|
||||
launching all the kernels and executing the operations while respecting their
|
||||
dependencies and necessary synchronizations as specified.
|
||||
|
||||
Because HIP graphs require some setup and initialization overhead before their
|
||||
first execution, graphs only provide a benefit for workloads that require
|
||||
many iterations to complete.
|
||||
|
||||
In both methods the :cpp:type:`hipGraph_t` template for a graph is used to define the graph.
|
||||
In order to actually launch a graph, the template needs to be instantiated using
|
||||
:cpp:func:`hipGraphInstantiate`, which results in an executable graph of type :cpp:type:`hipGraphExec_t`.
|
||||
This executable graph can then be launched with :cpp:func:`hipGraphLaunch`, replaying the
|
||||
operations within the graph. Note, that launching graphs is fundamentally no
|
||||
different to executing other HIP functions on a stream, except for the fact,
|
||||
that scheduling the operations within the graph encompasses less overhead and
|
||||
can enable some optimizations, but they still need to be associated with a stream for execution.
|
||||
|
||||
Memory management
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Memory that is used by operations in graphs can either be pre-allocated or
|
||||
managed within the graph. Graphs can contain nodes that take care of allocating
|
||||
memory on the device or copying memory between the host and the device.
|
||||
Whether you want to pre-allocate the memory or manage it within the graph
|
||||
depends on the use-case. If the graph is executed in a tight loop the
|
||||
performance is usually better when the memory is preallocated, so that it
|
||||
does not need to be reallocated in every iteration.
|
||||
|
||||
The same rules as for normal memory allocations apply for memory allocated and
|
||||
freed by nodes, meaning that the nodes that access memory allocated in a graph
|
||||
must be ordered after allocation and before freeing.
|
||||
|
||||
Memory management within the graph enables the runtime to take care of memory reuse and optimizations.
|
||||
The lifetime of memory managed in a graph begins when the execution reaches the
|
||||
node allocating the memory, and ends when either reaching the corresponding
|
||||
free node within the graph, or after graph execution when a corresponding
|
||||
:cpp:func:`hipFreeAsync` or :cpp:func:`hipFree` call is reached.
|
||||
The memory can also be freed with a free node in a different graph that is
|
||||
associated with the same memory address.
|
||||
|
||||
Unlike device memory that is not associated with a graph, this does not necessarily
|
||||
mean that the freed memory is returned back to the operating system immediately.
|
||||
Graphs can retain a memory pool for quickly reusing memory within the graph.
|
||||
This can be especially useful when memory is freed and reallocated later on
|
||||
within a graph, as that memory doesn't have to be requested from the operating system.
|
||||
It also potentially reduces the total memory footprint of the graph, by reusing the same memory.
|
||||
|
||||
The amount of memory allocated for graph memory pools on a specific device can
|
||||
be queried using :cpp:func:`hipDeviceGetGraphMemAttribute`.
|
||||
In order to return the freed memory :cpp:func:`hipDeviceGraphMemTrim` can be used.
|
||||
This will return any memory that is not in active use by graphs.
|
||||
|
||||
These memory allocations can also be set up to allow access from multiple GPUs,
|
||||
just like normal allocations. HIP then takes care of allocating and mapping the
|
||||
memory to the GPUs. When capturing a graph from a stream, the node sets the
|
||||
accessibility according to :cpp:func:`hipMemPoolSetAccess` at the time of capturing.
|
||||
|
||||
|
||||
Capture graphs from a stream
|
||||
================================================================================
|
||||
|
||||
The easy way to integrate HIP graphs into already existing code is to use
|
||||
:cpp:func:`hipStreamBeginCapture` and :cpp:func:`hipStreamEndCapture` to obtain a :cpp:type:`hipGraph_t`
|
||||
graph template that includes the captured operations.
|
||||
|
||||
When starting to capture operations for a graph using :cpp:func:`hipStreamBeginCapture`,
|
||||
the operations assigned to the stream are captured into a graph instead of being
|
||||
executed. The associated graph is returned when calling :cpp:func:`hipStreamEndCapture`, which
|
||||
also stops capturing operations.
|
||||
In order to capture to an already existing graph use :cpp:func:`hipStreamBeginCaptureToGraph`.
|
||||
|
||||
The functions assigned to the capturing stream are not executed, but instead are
|
||||
captured and defined as nodes in the graph, to be run when the instantiated
|
||||
graph is launched.
|
||||
|
||||
Functions must be associated with a stream in order to be captured.
|
||||
This means that non-HIP API functions are not captured by default, but are
|
||||
executed as standard functions when encountered and not added to the graph.
|
||||
In order to assign host functions to a stream use
|
||||
:cpp:func:`hipLaunchHostFunc`, as shown in the following code example.
|
||||
They will then be captured and defined as a host node in the resulting graph,
|
||||
and won't be executed when encountered.
|
||||
|
||||
Synchronous HIP API calls that are implicitly assigned to the default stream are
|
||||
not permitted while capturing a stream and will return an error. This is
|
||||
because they implicitly synchronize and cause a dependency that can not be
|
||||
captured within the stream. This includes functions like :cpp:func:`hipMalloc`,
|
||||
:cpp:func:`hipMemcpy` and :cpp:func:`hipFree`. In order to capture these to the stream, replace
|
||||
them with the corresponding asynchronous calls like :cpp:func:`hipMallocAsync`, :cpp:func:`hipMemcpyAsync` or :cpp:func:`hipFreeAsync`.
|
||||
|
||||
The general flow for using stream capture to create a graph template is:
|
||||
|
||||
#. Create a stream from which to capture the operations
|
||||
|
||||
#. Call :cpp:func:`hipStreamBeginCapture` before the first operation to be captured
|
||||
|
||||
#. Call :cpp:func:`hipStreamEndCapture` after the last operation to be captured
|
||||
|
||||
#. Define a :cpp:type:`hipGraph_t` graph template to which :cpp:func:`hipStreamEndCapture`
|
||||
passes the captured graph
|
||||
|
||||
The following code is an example of how to use the HIP graph API to capture a
|
||||
graph from a stream.
|
||||
|
||||
.. 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; \
|
||||
} \
|
||||
}
|
||||
|
||||
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] *= 2.0;}
|
||||
};
|
||||
__global__ void kernelB(int* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] = 3;}
|
||||
};
|
||||
__global__ void kernelC(double* arrayA, const int* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += arrayB[x];}
|
||||
};
|
||||
|
||||
struct set_vector_args{
|
||||
std::vector<double>& h_array;
|
||||
double value;
|
||||
};
|
||||
|
||||
void set_vector(void* args){
|
||||
set_vector_args h_args{*(reinterpret_cast<set_vector_args*>(args))};
|
||||
|
||||
std::vector<double>& vec{h_args.h_array};
|
||||
vec.assign(vec.size(), h_args.value);
|
||||
}
|
||||
|
||||
int main(){
|
||||
constexpr int numOfBlocks = 1024;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
constexpr size_t arraySize = 1U << 20;
|
||||
|
||||
// This example assumes that kernelA operates on data that needs to be initialized on
|
||||
// and copied from the host, while kernelB initializes the array that is passed to it.
|
||||
// Both arrays are then used as input to kernelC, where arrayA is also used as
|
||||
// output, that is copied back to the host, while arrayB is only read from and not modified.
|
||||
|
||||
double* d_arrayA;
|
||||
int* d_arrayB;
|
||||
std::vector<double> h_array(arraySize);
|
||||
constexpr double initValue = 2.0;
|
||||
|
||||
hipStream_t captureStream;
|
||||
HIP_CHECK(hipStreamCreate(&captureStream));
|
||||
|
||||
// Start capturing the operations assigned to the stream
|
||||
HIP_CHECK(hipStreamBeginCapture(captureStream, hipStreamCaptureModeGlobal));
|
||||
|
||||
// hipMallocAsync and hipMemcpyAsync are needed, to be able to assign it to a stream
|
||||
HIP_CHECK(hipMallocAsync(&d_arrayA, arraySize*sizeof(double), captureStream));
|
||||
HIP_CHECK(hipMallocAsync(&d_arrayB, arraySize*sizeof(int), captureStream));
|
||||
|
||||
// Assign host function to the stream
|
||||
// Needs a custom struct to pass the arguments
|
||||
set_vector_args args{h_array, initValue};
|
||||
HIP_CHECK(hipLaunchHostFunc(captureStream, set_vector, &args));
|
||||
|
||||
HIP_CHECK(hipMemcpyAsync(d_arrayA, h_array.data(), arraySize*sizeof(double), hipMemcpyHostToDevice, captureStream));
|
||||
|
||||
kernelA<<<numOfBlocks, threadsPerBlock, 0, captureStream>>>(d_arrayA, arraySize);
|
||||
kernelB<<<numOfBlocks, threadsPerBlock, 0, captureStream>>>(d_arrayB, arraySize);
|
||||
kernelC<<<numOfBlocks, threadsPerBlock, 0, captureStream>>>(d_arrayA, d_arrayB, arraySize);
|
||||
|
||||
HIP_CHECK(hipMemcpyAsync(h_array.data(), d_arrayA, arraySize*sizeof(*d_arrayA), hipMemcpyDeviceToHost, captureStream));
|
||||
|
||||
HIP_CHECK(hipFreeAsync(d_arrayA, captureStream));
|
||||
HIP_CHECK(hipFreeAsync(d_arrayB, captureStream));
|
||||
|
||||
// Stop capturing
|
||||
hipGraph_t graph;
|
||||
HIP_CHECK(hipStreamEndCapture(captureStream, &graph));
|
||||
|
||||
// Create an executable graph from the captured graph
|
||||
hipGraphExec_t graphExec;
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
|
||||
// The graph template can be deleted after the instantiation if it's not needed for later use
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
|
||||
// Actually launch the graph. The stream does not have
|
||||
// to be the same as the one used for capturing.
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, captureStream));
|
||||
|
||||
// Verify results
|
||||
constexpr double expected = initValue * 2.0 + 3;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(h_array[i] != expected){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expected << " got " << h_array[0] << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if(passed){
|
||||
std::cerr << "Validation passed." << std::endl;
|
||||
}
|
||||
|
||||
// Free graph and stream resources after usage
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipStreamDestroy(captureStream));
|
||||
}
|
||||
|
||||
Explicit graph creation
|
||||
================================================================================
|
||||
|
||||
Graphs can also be created directly using the HIP graph API, giving more
|
||||
fine-grained control over the graph. In this case, the graph nodes are created
|
||||
explicitly, together with their parameters and dependencies, which specify the
|
||||
edges of the graph, thereby forming the graph structure.
|
||||
|
||||
The nodes are represented by the generic :cpp:type:`hipGraphNode_t` type. The actual
|
||||
node type is implicitly defined by the specific function used to add the node to
|
||||
the graph, for example :cpp:func:`hipGraphAddKernelNode` See the
|
||||
:ref:`HIP graph API documentation<graph_management_reference>` for the
|
||||
available functions, they are of type ``hipGraphAdd{Type}Node``. Each type of
|
||||
node also has a predefined set of parameters depending on the operation, for
|
||||
example :cpp:class:`hipKernelNodeParams` for a kernel launch. See the
|
||||
:doc:`documentation for the general hipGraphNodeParams type<../../doxygen/html/structhip_graph_node_params>`
|
||||
for a list of available parameter types and their members.
|
||||
|
||||
The general flow for explicitly creating a graph is usually:
|
||||
|
||||
#. Create a graph :cpp:type:`hipGraph_t`
|
||||
|
||||
#. Create the nodes and their parameters and add them to the graph
|
||||
|
||||
#. Define a :cpp:type:`hipGraphNode_t`
|
||||
|
||||
#. Define the parameter struct for the desired operation, by explicitly setting the appropriate struct's members.
|
||||
|
||||
#. Use the appropriate ``hipGraphAdd{Type}Node`` function to add the node to the graph.
|
||||
|
||||
#. The dependencies can be defined when adding the node to the graph, or afterwards by using :cpp:func:`hipGraphAddDependencies`
|
||||
|
||||
The following code example demonstrates how to explicitly create nodes in order to create a graph.
|
||||
|
||||
.. 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; \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void kernelA(double* arrayA, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] *= 2.0;}
|
||||
};
|
||||
__global__ void kernelB(int* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayB[x] = 3;}
|
||||
};
|
||||
__global__ void kernelC(double* arrayA, const int* arrayB, size_t size){
|
||||
const size_t x = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
if(x < size){arrayA[x] += arrayB[x];}
|
||||
};
|
||||
|
||||
struct set_vector_args{
|
||||
std::vector<double>& h_array;
|
||||
double value;
|
||||
};
|
||||
|
||||
void set_vector(void* args){
|
||||
set_vector_args h_args{*(reinterpret_cast<set_vector_args*>(args))};
|
||||
|
||||
std::vector<double>& vec{h_args.h_array};
|
||||
vec.assign(vec.size(), h_args.value);
|
||||
}
|
||||
|
||||
int main(){
|
||||
constexpr int numOfBlocks = 1024;
|
||||
constexpr int threadsPerBlock = 1024;
|
||||
size_t arraySize = 1U << 20;
|
||||
|
||||
// The pointers to the device memory don't need to be declared here,
|
||||
// they are contained within the hipMemAllocNodeParams as the dptr member
|
||||
std::vector<double> h_array(arraySize);
|
||||
constexpr double initValue = 2.0;
|
||||
|
||||
// Create graph an empty graph
|
||||
hipGraph_t graph;
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
|
||||
// Parameters to allocate arrays
|
||||
hipMemAllocNodeParams allocArrayAParams{};
|
||||
allocArrayAParams.poolProps.allocType = hipMemAllocationTypePinned;
|
||||
allocArrayAParams.poolProps.location.type = hipMemLocationTypeDevice;
|
||||
allocArrayAParams.poolProps.location.id = 0; // GPU on which memory resides
|
||||
allocArrayAParams.bytesize = arraySize * sizeof(double);
|
||||
|
||||
hipMemAllocNodeParams allocArrayBParams{};
|
||||
allocArrayBParams.poolProps.allocType = hipMemAllocationTypePinned;
|
||||
allocArrayBParams.poolProps.location.type = hipMemLocationTypeDevice;
|
||||
allocArrayBParams.poolProps.location.id = 0; // GPU on which memory resides
|
||||
allocArrayBParams.bytesize = arraySize * sizeof(int);
|
||||
|
||||
// Add the allocation nodes to the graph. They don't have any dependencies
|
||||
hipGraphNode_t allocNodeA, allocNodeB;
|
||||
HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocArrayAParams));
|
||||
HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeB, graph, nullptr, 0, &allocArrayBParams));
|
||||
|
||||
// Parameters for the host function
|
||||
// Needs custom struct to pass the arguments
|
||||
set_vector_args args{h_array, initValue};
|
||||
hipHostNodeParams hostParams{};
|
||||
hostParams.fn = set_vector;
|
||||
hostParams.userData = static_cast<void*>(&args);
|
||||
|
||||
// Add the host node that initializes the host array. It also doesn't have any dependencies
|
||||
hipGraphNode_t hostNode;
|
||||
HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams));
|
||||
|
||||
// Add memory copy node, that copies the initialized host array to the device.
|
||||
// It has to wait for the host array to be initialized and the device memory to be allocated
|
||||
hipGraphNode_t cpyNodeDependencies[] = {allocNodeA, hostNode};
|
||||
hipGraphNode_t cpyToDevNode;
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&cpyToDevNode, graph, cpyNodeDependencies, 1, allocArrayAParams.dptr, h_array.data(), arraySize * sizeof(double), hipMemcpyHostToDevice));
|
||||
|
||||
// Parameters for kernelA
|
||||
hipKernelNodeParams kernelAParams;
|
||||
void* kernelAArgs[] = {&allocArrayAParams.dptr, static_cast<void*>(&arraySize)};
|
||||
kernelAParams.func = reinterpret_cast<void*>(kernelA);
|
||||
kernelAParams.gridDim = numOfBlocks;
|
||||
kernelAParams.blockDim = threadsPerBlock;
|
||||
kernelAParams.sharedMemBytes = 0;
|
||||
kernelAParams.kernelParams = kernelAArgs;
|
||||
kernelAParams.extra = nullptr;
|
||||
|
||||
// Add the node for kernelA. It has to wait for the memory copy to finish, as it depends on the values from the host array.
|
||||
hipGraphNode_t kernelANode;
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelANode, graph, &cpyToDevNode, 1, &kernelAParams));
|
||||
|
||||
// Parameters for kernelB
|
||||
hipKernelNodeParams kernelBParams;
|
||||
void* kernelBArgs[] = {&allocArrayBParams.dptr, static_cast<void*>(&arraySize)};
|
||||
kernelBParams.func = reinterpret_cast<void*>(kernelB);
|
||||
kernelBParams.gridDim = numOfBlocks;
|
||||
kernelBParams.blockDim = threadsPerBlock;
|
||||
kernelBParams.sharedMemBytes = 0;
|
||||
kernelBParams.kernelParams = kernelBArgs;
|
||||
kernelBParams.extra = nullptr;
|
||||
|
||||
// Add the node for kernelB. It only has to wait for the memory to be allocated, as it initializes the array.
|
||||
hipGraphNode_t kernelBNode;
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelBNode, graph, &allocNodeB, 1, &kernelBParams));
|
||||
|
||||
// Parameters for kernelC
|
||||
hipKernelNodeParams kernelCParams;
|
||||
void* kernelCArgs[] = {&allocArrayAParams.dptr, &allocArrayBParams.dptr, static_cast<void*>(&arraySize)};
|
||||
kernelCParams.func = reinterpret_cast<void*>(kernelC);
|
||||
kernelCParams.gridDim = numOfBlocks;
|
||||
kernelCParams.blockDim = threadsPerBlock;
|
||||
kernelCParams.sharedMemBytes = 0;
|
||||
kernelCParams.kernelParams = kernelCArgs;
|
||||
kernelCParams.extra = nullptr;
|
||||
|
||||
// Add the node for kernelC. It has to wait on both kernelA and kernelB to finish, as it depends on their results.
|
||||
hipGraphNode_t kernelCNode;
|
||||
hipGraphNode_t kernelCDependencies[] = {kernelANode, kernelBNode};
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelCNode, graph, kernelCDependencies, 1, &kernelCParams));
|
||||
|
||||
// Copy the results back to the host. Has to wait for kernelC to finish.
|
||||
hipGraphNode_t cpyToHostNode;
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&cpyToHostNode, graph, &kernelCNode, 1, h_array.data(), allocArrayAParams.dptr, arraySize * sizeof(double), hipMemcpyDeviceToHost));
|
||||
|
||||
// Free array of allocNodeA. It needs to wait for the copy to finish, as kernelC stores its results in it.
|
||||
hipGraphNode_t freeNodeA;
|
||||
HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &cpyToHostNode, 1, allocArrayAParams.dptr));
|
||||
// Free array of allocNodeB. It only needs to wait for kernelC to finish, as it is not written back to the host.
|
||||
hipGraphNode_t freeNodeB;
|
||||
HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeB, graph, &kernelCNode, 1, allocArrayBParams.dptr));
|
||||
|
||||
// Instantiate the graph in order to execute it
|
||||
hipGraphExec_t graphExec;
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
|
||||
// The graph can be freed after the instantiation if it's not needed for other purposes
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
|
||||
// Actually launch the graph
|
||||
hipStream_t graphStream;
|
||||
HIP_CHECK(hipStreamCreate(&graphStream));
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, graphStream));
|
||||
|
||||
// Verify results
|
||||
constexpr double expected = initValue * 2.0 + 3;
|
||||
bool passed = true;
|
||||
for(size_t i = 0; i < arraySize; ++i){
|
||||
if(h_array[i] != expected){
|
||||
passed = false;
|
||||
std::cerr << "Validation failed! Expected " << expected << " got " << h_array[0] << std::endl;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if(passed){
|
||||
std::cerr << "Validation passed." << std::endl;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipStreamDestroy(graphStream));
|
||||
}
|
||||
@@ -0,0 +1,107 @@
|
||||
.. meta::
|
||||
:description: Initialization.
|
||||
:keywords: AMD, ROCm, HIP, initialization
|
||||
|
||||
.. _initialization:
|
||||
|
||||
********************************************************************************
|
||||
Initialization
|
||||
********************************************************************************
|
||||
|
||||
The initialization involves setting up the environment and resources needed for
|
||||
using GPUs. The following steps are covered with the initialization:
|
||||
|
||||
- Setting up the HIP runtime
|
||||
|
||||
This includes reading the environment variables set during init, setting up
|
||||
the active or visible devices, loading necessary libraries, setting up
|
||||
internal buffers for memory copies or cooperative launches, initialize the
|
||||
compiler as well as HSA runtime and checks any errors due to lack of resources
|
||||
or no active devices.
|
||||
|
||||
- Querying and setting GPUs
|
||||
|
||||
Identifying and querying the available GPU devices on the system.
|
||||
|
||||
- Setting up contexts
|
||||
|
||||
Creating contexts for each GPU device, which are essential for managing
|
||||
resources and executing kernels. For further details, check the :ref:`context
|
||||
section <context_driver_api>`.
|
||||
|
||||
Initialize the HIP runtime
|
||||
================================================================================
|
||||
|
||||
The HIP runtime is initialized automatically when the first HIP API call is
|
||||
made. However, you can explicitly initialize it using :cpp:func:`hipInit`,
|
||||
to be able to control the timing of the initialization. The manual
|
||||
initialization can be useful to ensure that the GPU is initialized and
|
||||
ready, or to isolate GPU initialization time from other parts of
|
||||
your program.
|
||||
|
||||
.. note::
|
||||
|
||||
You can use :cpp:func:`hipDeviceReset` to delete all streams created, memory
|
||||
allocated, kernels running and events created by the current process. Any new
|
||||
HIP API call initializes the HIP runtime again.
|
||||
|
||||
Querying and setting GPUs
|
||||
================================================================================
|
||||
|
||||
If multiple GPUs are available in the system, you can query and select the
|
||||
desired GPU(s) to use based on device properties, such as size of global memory,
|
||||
size shared memory per block, support of cooperative launch and support of
|
||||
managed memory.
|
||||
|
||||
Querying GPUs
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The properties of a GPU can be queried using :cpp:func:`hipGetDeviceProperties`,
|
||||
which returns a struct of :cpp:struct:`hipDeviceProp_t`. The properties in the
|
||||
struct can be used to identify a device or give an overview of hardware
|
||||
characteristics, that might make one GPU better suited for the task than others.
|
||||
|
||||
The :cpp:func:`hipGetDeviceCount` function returns the number of available GPUs,
|
||||
which can be used to loop over the available GPUs.
|
||||
|
||||
Example code of querying GPUs:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
int main() {
|
||||
|
||||
int deviceCount;
|
||||
if (hipGetDeviceCount(&deviceCount) == hipSuccess){
|
||||
for (int i = 0; i < deviceCount; ++i){
|
||||
hipDeviceProp_t prop;
|
||||
if ( hipGetDeviceProperties(&prop, i) == hipSuccess)
|
||||
std::cout << "Device" << i << prop.name << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Setting the GPU
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
:cpp:func:`hipSetDevice` function select the GPU to be used for subsequent HIP
|
||||
operations. This function performs several key tasks:
|
||||
|
||||
- Context Binding
|
||||
|
||||
Binds the current thread to the context of the specified GPU device. This
|
||||
ensures that all subsequent operations are executed on the selected device.
|
||||
|
||||
- Resource Allocation
|
||||
|
||||
Prepares the device for resource allocation, such as memory allocation and
|
||||
stream creation.
|
||||
|
||||
- Check device availability
|
||||
|
||||
Checks for errors in device selection and returns error if the specified
|
||||
device is not available or not capable of executing HIP operations.
|
||||
@@ -0,0 +1,52 @@
|
||||
.. meta::
|
||||
:description: Memory management and its usage
|
||||
:keywords: AMD, ROCm, HIP, CUDA, memory management
|
||||
|
||||
.. _memory_management:
|
||||
|
||||
********************************************************************************
|
||||
Memory management
|
||||
********************************************************************************
|
||||
|
||||
Memory management is an important part of the HIP runtime API, when creating
|
||||
high-performance applications. Both allocating and copying memory can result in
|
||||
bottlenecks, which can significantly impact performance.
|
||||
|
||||
The programming model is based on a system with a host and a device, each having
|
||||
its own distinct memory. Kernels operate on :ref:`device_memory`, while host functions
|
||||
operate on :ref:`host_memory`.
|
||||
|
||||
The runtime offers functions for allocating, freeing, and copying device memory,
|
||||
along with transferring data between host and device memory.
|
||||
|
||||
Here are the various memory management techniques:
|
||||
|
||||
* :ref:`coherence_control`
|
||||
* :ref:`unified_memory`
|
||||
* :ref:`virtual_memory`
|
||||
* :ref:`stream_ordered_memory_allocator_how-to`
|
||||
|
||||
Memory allocation
|
||||
================================================================================
|
||||
|
||||
The API calls and the resulting allocations are listed here:
|
||||
|
||||
.. list-table:: Memory coherence control
|
||||
:header-rows: 1
|
||||
:align: center
|
||||
|
||||
* - API
|
||||
- Data location
|
||||
- Allocation
|
||||
* - System allocated
|
||||
- Host
|
||||
- :ref:`Pageable <pageable_host_memory>`
|
||||
* - :cpp:func:`hipMallocManaged`
|
||||
- Host
|
||||
- :ref:`Managed <unified_memory>`
|
||||
* - :cpp:func:`hipHostMalloc`
|
||||
- Host
|
||||
- :ref:`Pinned <pinned_host_memory>`
|
||||
* - :cpp:func:`hipMalloc`
|
||||
- Device
|
||||
- Pinned
|
||||
@@ -0,0 +1,178 @@
|
||||
.. meta::
|
||||
:description: HIP coherence control
|
||||
ecosystem ROCm software.
|
||||
:keywords: AMD, ROCm, HIP, host memory
|
||||
|
||||
.. _coherence_control:
|
||||
|
||||
*******************************************************************************
|
||||
Coherence control
|
||||
*******************************************************************************
|
||||
|
||||
Memory coherence describes how memory of a specific part of the system is
|
||||
visible to the other parts of the system. For example, how GPU memory is visible
|
||||
to the CPU and vice versa. In HIP, host and device memory can be allocated with
|
||||
two different types of coherence:
|
||||
|
||||
* **Coarse-grained coherence:** The memory is considered up-to-date only after
|
||||
synchronization performed using :cpp:func:`hipDeviceSynchronize`,
|
||||
:cpp:func:`hipStreamSynchronize`, or any blocking operation that acts on the
|
||||
null stream such as :cpp:func:`hipMemcpy`. To avoid the cache from being
|
||||
accessed by a part of the system while simultaneously being written by
|
||||
another, the memory is made visible only after the caches have been flushed.
|
||||
|
||||
* **Fine-grained coherence:** The memory is coherent even while being modified
|
||||
by a part of the system. Fine-grained coherence ensures that up-to-date data
|
||||
is visible to others regardless of kernel boundaries. This can be useful if
|
||||
both host and device operate on the same data.
|
||||
|
||||
.. note::
|
||||
|
||||
To achieve fine-grained coherence, many AMD GPUs use a limited cache policy,
|
||||
such as leaving these allocations uncached by the GPU or making them read-only.
|
||||
|
||||
Mi200 accelerator's hardware based floating point instructions work on
|
||||
coarse-grained memory regions. Coarse-grained coherence is typically useful in
|
||||
reducing host-device interconnect communication.
|
||||
|
||||
To check the availability of fine- and coarse-grained memory pools, use
|
||||
``rocminfo``:
|
||||
|
||||
.. code-block:: sh
|
||||
|
||||
$ rocminfo
|
||||
...
|
||||
*******
|
||||
Agent 1
|
||||
*******
|
||||
Name: AMD EPYC 7742 64-Core Processor
|
||||
...
|
||||
Pool Info:
|
||||
Pool 1
|
||||
Segment: GLOBAL; FLAGS: FINE GRAINED
|
||||
...
|
||||
Pool 3
|
||||
Segment: GLOBAL; FLAGS: COARSE GRAINED
|
||||
...
|
||||
*******
|
||||
Agent 9
|
||||
*******
|
||||
Name: gfx90a
|
||||
...
|
||||
Pool Info:
|
||||
Pool 1
|
||||
Segment: GLOBAL; FLAGS: COARSE GRAINED
|
||||
...
|
||||
|
||||
The APIs, flags and respective memory coherence control are listed in the
|
||||
following table:
|
||||
|
||||
.. list-table:: Memory coherence control
|
||||
:widths: 25, 35, 20, 20
|
||||
:header-rows: 1
|
||||
:align: center
|
||||
|
||||
* - API
|
||||
- Flag
|
||||
- :cpp:func:`hipMemAdvise` call with argument
|
||||
- Coherence
|
||||
* - ``hipHostMalloc`` :sup:`1`
|
||||
- ``hipHostMallocDefault``
|
||||
-
|
||||
- Fine-grained
|
||||
* - ``hipHostMalloc`` :sup:`1`
|
||||
- ``hipHostMallocNonCoherent``
|
||||
-
|
||||
- Coarse-grained
|
||||
* - ``hipExtMallocWithFlags``
|
||||
- ``hipDeviceMallocDefault``
|
||||
-
|
||||
- Coarse-grained
|
||||
* - ``hipExtMallocWithFlags``
|
||||
- ``hipDeviceMallocFinegrained``
|
||||
-
|
||||
- Fine-grained
|
||||
* - ``hipMallocManaged``
|
||||
-
|
||||
-
|
||||
- Fine-grained
|
||||
* - ``hipMallocManaged``
|
||||
-
|
||||
- ``hipMemAdviseSetCoarseGrain``
|
||||
- Coarse-grained
|
||||
* - ``malloc``
|
||||
-
|
||||
-
|
||||
- Fine-grained
|
||||
* - ``malloc``
|
||||
-
|
||||
- ``hipMemAdviseSetCoarseGrain``
|
||||
- Coarse-grained
|
||||
|
||||
:sup:`1` The :cpp:func:`hipHostMalloc` memory allocation coherence mode can be
|
||||
affected by the ``HIP_HOST_COHERENT`` environment variable, if the
|
||||
``hipHostMallocCoherent``, ``hipHostMallocNonCoherent``, and
|
||||
``hipHostMallocMapped`` are unset. If neither these flags nor the
|
||||
``HIP_HOST_COHERENT`` environment variable is set, or set as 0, the host memory
|
||||
allocation is coarse-grained.
|
||||
|
||||
.. note::
|
||||
|
||||
* When ``hipHostMallocMapped`` flag is set, the allocated host memory is
|
||||
fine-grained and the ``hipHostMallocNonCoherent`` flag is ignored.
|
||||
* Setting both the ``hipHostMallocCoherent`` and
|
||||
``hipHostMallocNonCoherent`` flags leads to an illegal state.
|
||||
|
||||
Visibility of synchronization functions
|
||||
================================================================================
|
||||
|
||||
The fine-grained coherence memory is visible at the synchronization points,
|
||||
however the visibility of coarse-grained memory depends on the synchronization
|
||||
function used. The effect and visibility of various synchronization functions on
|
||||
fine- and coarse-grained memory types are listed here:
|
||||
|
||||
.. list-table:: HIP synchronize functions effect and visibility
|
||||
|
||||
* - HIP API
|
||||
- :cpp:func:`hipStreamSynchronize`
|
||||
- :cpp:func:`hipDeviceSynchronize`
|
||||
- :cpp:func:`hipEventSynchronize`
|
||||
- :cpp:func:`hipStreamWaitEvent`
|
||||
* - Synchronization effect
|
||||
- Host waits for all commands in the specified stream to complete
|
||||
- Host waits for all commands in all streams on the specified device to complete
|
||||
- Host waits for the specified event to complete
|
||||
- Stream waits for the specified event to complete
|
||||
* - Fence
|
||||
- System-scope release
|
||||
- System-scope release
|
||||
- System-scope release
|
||||
- None
|
||||
* - Fine-grained host memory visibility
|
||||
- Yes
|
||||
- Yes
|
||||
- Yes
|
||||
- Yes
|
||||
* - Coarse-grained host memory visibility
|
||||
- Yes
|
||||
- Yes
|
||||
- Depends on the used event.
|
||||
- No
|
||||
|
||||
You can control the release scope for ``hipEvents``. By default, the GPU
|
||||
performs a device-scope acquire and release operation with each recorded event.
|
||||
This makes the host and device memory visible to other commands executing on the
|
||||
same device.
|
||||
|
||||
:cpp:func:`hipEventCreateWithFlags`: You can specify a stronger system-level
|
||||
fence by creating the event with ``hipEventCreateWithFlags``:
|
||||
|
||||
* ``hipEventReleaseToSystem``: Performs a system-scope release operation when
|
||||
the event is recorded. This makes both fine-grained and coarse-grained host
|
||||
memory visible to other agents in the system, which might also involve
|
||||
heavyweight operations such as cache flushing. Fine-grained memory typically
|
||||
uses lighter-weight in-kernel synchronization mechanisms such as an atomic
|
||||
operation and thus doesn't need to use ``hipEventReleaseToSystem``.
|
||||
|
||||
* ``hipEventDisableTiming``: Events created with this flag don't record
|
||||
profiling data, which significantly improves synchronization performance.
|
||||
@@ -0,0 +1,52 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the device memory of the HIP ecosystem
|
||||
ROCm software.
|
||||
:keywords: AMD, ROCm, HIP, device memory
|
||||
|
||||
.. _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.
|
||||
|
||||
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.
|
||||
|
||||
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.
|
||||
|
||||
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.
|
||||
|
||||
The :ref:`texture management module <texture_management_reference>` of the HIP
|
||||
runtime API reference contains the functions of texture memory.
|
||||
|
||||
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.
|
||||
|
||||
The :ref:`surface objects module <surface_object_reference>` of HIP runtime API
|
||||
contains the functions for creating, destroying and reading surface memory.
|
||||
@@ -0,0 +1,214 @@
|
||||
.. meta::
|
||||
:description: This chapter describes the texture fetching modes of the HIP ecosystem
|
||||
ROCm software.
|
||||
:keywords: AMD, ROCm, HIP, Texture, Texture Fetching
|
||||
|
||||
.. _texture_fetching:
|
||||
|
||||
*******************************************************************************
|
||||
Texture fetching
|
||||
*******************************************************************************
|
||||
|
||||
`Textures <../../../../doxygen/html/group___texture.html>`_ are more than just a buffer
|
||||
interpreted as a 1D, 2D, or 3D array.
|
||||
|
||||
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].
|
||||
|
||||
Depending on the index, texture sampling or texture addressing is performed,
|
||||
which decides the return value.
|
||||
|
||||
**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
|
||||
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
|
||||
:align: center
|
||||
|
||||
Texture used as example
|
||||
|
||||
Texture sampling
|
||||
===============================================================================
|
||||
|
||||
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.
|
||||
|
||||
The various texture sampling methods are discussed in the following sections.
|
||||
|
||||
.. _texture_fetching_nearest:
|
||||
|
||||
Nearest point sampling
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
In this method, the modulo of index is calculated as:
|
||||
|
||||
``tex(x) = T[floor(x)]``
|
||||
|
||||
This is also applicable for 2D and 3D variants.
|
||||
|
||||
This doesn't interpolate between neighboring values, which results in a
|
||||
pixelated look.
|
||||
|
||||
The following image shows a texture stretched to a 4x4 pixel quad but still
|
||||
indexed in the [0 to 1] range. The in-between values are the same as the values
|
||||
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
|
||||
:align: center
|
||||
|
||||
Texture upscaled with nearest point sampling
|
||||
|
||||
.. _texture_fetching_linear:
|
||||
|
||||
Linear filtering
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
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:
|
||||
|
||||
* 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]``
|
||||
* For three dimensional textures: ``tex(x,y,z) = (1-α)(1-β)(1-γ)T[i,j,k] + α(1-β)(1-γ)T[i+1,j,k] + (1-α)β(1-γ)T[i,j+1,k] + αβ(1-γ)T[i+1,j+1,k] + (1-α)(1-β)γT[i,j,k+1] + α(1-β)γT[i+1,j,k+1] + (1-α)βγT[i,j+1,k+1] + αβγT[i+1,j+1,k+1]``
|
||||
|
||||
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
|
||||
still indexed in the [0 to 1] range. The in-between values are interpolated
|
||||
between the neighboring texels.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/linear.png
|
||||
:width: 300
|
||||
:alt: Texture upscaled with linear filtering
|
||||
:align: center
|
||||
|
||||
Texture upscaled with linear filtering
|
||||
|
||||
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 following sections describe the various texture addressing methods.
|
||||
|
||||
.. _texture_fetching_border:
|
||||
|
||||
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.
|
||||
|
||||
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.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/border.png
|
||||
:width: 300
|
||||
:alt: Texture with yellow border color
|
||||
:align: center
|
||||
|
||||
Texture with yellow border color.
|
||||
|
||||
The purple lines are not part of the texture. They only denote the edge, where
|
||||
the addressing begins.
|
||||
|
||||
.. _texture_fetching_clamp:
|
||||
|
||||
Address mode clamp
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
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.
|
||||
|
||||
The following image shows the texture on a 4x4 pixel quad, indexed in the
|
||||
[0 to 3] range. The out-of-bounds values are repeating the values at the edge of
|
||||
the texture.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/clamp.png
|
||||
:width: 300
|
||||
:alt: Texture with clamp addressing
|
||||
:align: center
|
||||
|
||||
Texture with clamp addressing
|
||||
|
||||
The purple lines are not part of the texture. They only denote the edge, where
|
||||
the addressing begins.
|
||||
|
||||
.. _texture_fetching_wrap:
|
||||
|
||||
Address mode wrap
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
Wrap mode addressing is only available for normalized texture coordinates. In
|
||||
this addressing mode, the fractional part of the index is used:
|
||||
|
||||
``tex(frac(x))``
|
||||
|
||||
This creates a repeating image effect.
|
||||
|
||||
The following image shows the texture on a 4x4 pixel quad, indexed in the
|
||||
[0 to 3] range. The out-of-bounds values are repeating the original texture.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/wrap.png
|
||||
:width: 300
|
||||
:alt: Texture with wrap addressing
|
||||
:align: center
|
||||
|
||||
Texture with wrap addressing.
|
||||
|
||||
The purple lines are not part of the texture. They only denote the edge, where
|
||||
the addressing begins.
|
||||
|
||||
.. _texture_fetching_mirror:
|
||||
|
||||
Address mode mirror
|
||||
-------------------------------------------------------------------------------
|
||||
|
||||
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.
|
||||
|
||||
The formula is the following:
|
||||
|
||||
``tex(frac(x))``, if ``floor(x)`` is even,
|
||||
|
||||
``tex(1 - frac(x))``, if ``floor(x)`` is odd.
|
||||
|
||||
The following image shows the texture on a 4x4 pixel quad, indexed in The
|
||||
[0 to 3] range. The out-of-bounds values are repeating the original texture, but
|
||||
mirrored.
|
||||
|
||||
.. figure:: ../../../../data/how-to/hip_runtime_api/memory_management/textures/mirror.png
|
||||
:width: 300
|
||||
:alt: Texture with mirror addressing
|
||||
:align: center
|
||||
|
||||
Texture with mirror addressing
|
||||
|
||||
The purple lines are not part of the texture. They only denote the edge, where
|
||||
the addressing begins.
|
||||
@@ -0,0 +1,239 @@
|
||||
.. meta::
|
||||
:description: Host memory of the HIP ecosystem
|
||||
:keywords: AMD, ROCm, HIP, host memory
|
||||
|
||||
.. _host_memory:
|
||||
|
||||
********************************************************************************
|
||||
Host memory
|
||||
********************************************************************************
|
||||
|
||||
Host memory is the "normal" memory residing in the host RAM and allocated by C
|
||||
or C++. Host memory can be allocated in two different ways:
|
||||
|
||||
* Pageable memory
|
||||
|
||||
* Pinned memory
|
||||
|
||||
The following figure explains how data is transferred in pageable and pinned
|
||||
memory.
|
||||
|
||||
.. figure:: ../../../data/how-to/hip_runtime_api/memory_management/pageable_pinned.svg
|
||||
|
||||
The pageable and pinned memory allow you to exercise direct control over
|
||||
memory operations, which is known as explicit memory management. When using the
|
||||
unified memory, you get a simplified memory model with less control over
|
||||
low level memory operations.
|
||||
|
||||
The difference in memory transfers between explicit and unified memory
|
||||
management is highlighted in the following figure:
|
||||
|
||||
.. figure:: ../../../data/how-to/hip_runtime_api/memory_management/unified_memory/um.svg
|
||||
|
||||
For more details on unified memory management, see :doc:`/how-to/hip_runtime_api/memory_management/unified_memory`.
|
||||
|
||||
.. _pageable_host_memory:
|
||||
|
||||
Pageable memory
|
||||
================================================================================
|
||||
|
||||
Pageable memory exists on memory blocks known as "pages" that can be migrated to
|
||||
other memory storage. For example, migrating memory between CPU sockets on a
|
||||
motherboard or in a system whose RAM runs out of space and starts dumping pages
|
||||
into the swap partition of the hard drive.
|
||||
|
||||
Pageable memory is usually allocated with a call to ``malloc`` or ``new`` in a
|
||||
C++ application.
|
||||
|
||||
**Example:** Using pageable host memory in HIP
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
} \
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
const int element_number = 100;
|
||||
|
||||
int *host_input, *host_output;
|
||||
// Host allocation
|
||||
host_input = new int[element_number];
|
||||
host_output = new int[element_number];
|
||||
|
||||
// Host data preparation
|
||||
for (int i = 0; i < element_number; i++) {
|
||||
host_input[i] = i;
|
||||
}
|
||||
memset(host_output, 0, element_number * sizeof(int));
|
||||
|
||||
int *device_input, *device_output;
|
||||
|
||||
// Device allocation
|
||||
HIP_CHECK(hipMalloc((int **)&device_input, element_number * sizeof(int)));
|
||||
HIP_CHECK(hipMalloc((int **)&device_output, element_number * sizeof(int)));
|
||||
|
||||
// Device data preparation
|
||||
HIP_CHECK(hipMemcpy(device_input, host_input, element_number * sizeof(int), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemset(device_output, 0, element_number * sizeof(int)));
|
||||
|
||||
// Run the kernel
|
||||
// ...
|
||||
|
||||
HIP_CHECK(hipMemcpy(device_input, host_input, element_number * sizeof(int), hipMemcpyHostToDevice));
|
||||
|
||||
// Free host memory
|
||||
delete[] host_input;
|
||||
delete[] host_output;
|
||||
|
||||
// Free device memory
|
||||
HIP_CHECK(hipFree(device_input));
|
||||
HIP_CHECK(hipFree(device_output));
|
||||
}
|
||||
|
||||
.. note::
|
||||
|
||||
:cpp:func:`hipMalloc` and :cpp:func:`hipFree` are blocking calls. However, HIP
|
||||
also provides non-blocking versions :cpp:func:`hipMallocAsync` and
|
||||
:cpp:func:`hipFreeAsync`, which require a stream as an additional argument.
|
||||
|
||||
.. _pinned_host_memory:
|
||||
|
||||
Pinned memory
|
||||
================================================================================
|
||||
|
||||
Pinned memory or page-locked memory is stored in pages that are locked in
|
||||
specific sectors in RAM and can't be migrated. The pointer can be used on both
|
||||
host and device. Accessing host-resident pinned memory in device kernels is
|
||||
generally not recommended for performance, as it can force the data to traverse
|
||||
the host-device interconnect such as PCIe, which is much slower than the
|
||||
on-device bandwidth.
|
||||
|
||||
The advantage of pinned memory is the improved transfer time between host and
|
||||
device. For transfer operations, such as :cpp:func:`hipMemcpy` or :cpp:func:`hipMemcpyAsync`,
|
||||
using pinned memory instead of pageable memory on the host can lead to a three times
|
||||
improvement in bandwidth.
|
||||
|
||||
The disadvantage of pinned memory is the reduced availability of RAM for other
|
||||
processes, which can negatively impact the overall performance of the host.
|
||||
|
||||
**Example:** Using pinned memory in HIP
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
} \
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
const int element_number = 100;
|
||||
|
||||
int *host_input, *host_output;
|
||||
// Host allocation
|
||||
HIP_CHECK(hipHostMalloc((int **)&host_input, element_number * sizeof(int)));
|
||||
HIP_CHECK(hipHostMalloc((int **)&host_output, element_number * sizeof(int)));
|
||||
|
||||
// Host data preparation
|
||||
for (int i = 0; i < element_number; i++) {
|
||||
host_input[i] = i;
|
||||
}
|
||||
memset(host_output, 0, element_number * sizeof(int));
|
||||
|
||||
int *device_input, *device_output;
|
||||
|
||||
// Device allocation
|
||||
HIP_CHECK(hipMalloc((int **)&device_input, element_number * sizeof(int)));
|
||||
HIP_CHECK(hipMalloc((int **)&device_output, element_number * sizeof(int)));
|
||||
|
||||
// Device data preparation
|
||||
HIP_CHECK(hipMemcpy(device_input, host_input, element_number * sizeof(int), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemset(device_output, 0, element_number * sizeof(int)));
|
||||
|
||||
// Run the kernel
|
||||
// ...
|
||||
|
||||
HIP_CHECK(hipMemcpy(device_input, host_input, element_number * sizeof(int), hipMemcpyHostToDevice));
|
||||
|
||||
// Free host memory
|
||||
delete[] host_input;
|
||||
delete[] host_output;
|
||||
|
||||
// Free device memory
|
||||
HIP_CHECK(hipFree(device_input));
|
||||
HIP_CHECK(hipFree(device_output));
|
||||
}
|
||||
|
||||
.. _memory_allocation_flags:
|
||||
|
||||
Memory allocation flags for pinned memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The memory allocation for pinned memory can be controlled using ``hipHostMalloc`` flags:
|
||||
|
||||
* ``hipHostMallocPortable``: The memory allocation is not restricted to the
|
||||
context making the allocation.
|
||||
|
||||
* ``hipHostMallocMapped``: The memory is allocated into the address space for
|
||||
the current device and the device pointer can be obtained with
|
||||
:cpp:func:`hipHostGetDevicePointer`.
|
||||
|
||||
* ``hipHostMallocNumaUser``: The host memory allocation follows Numa policy
|
||||
specified by the user. Target of Numa policy is to select a CPU that is
|
||||
closest to each GPU. Numa distance is the distance between GPU and CPU
|
||||
devices.
|
||||
|
||||
* ``hipHostMallocWriteCombined``: The memory is allocated as write-combined.
|
||||
Although lacking read efficiency by most CPUs, write-combined allocation might
|
||||
be transferred faster across the PCIe bus on some system configurations. It's
|
||||
a good option for data transfer from host to device via mapped pinned memory.
|
||||
|
||||
* ``hipHostMallocCoherent``: Fine-grained memory is allocated. Overrides
|
||||
``HIP_HOST_COHERENT`` environment variable for specific allocation. For
|
||||
details, see :ref:`coherence_control`.
|
||||
|
||||
* ``hipHostMallocNonCoherent``: Coarse-grained memory is allocated. Overrides
|
||||
``HIP_HOST_COHERENT`` environment variable for specific allocation. For
|
||||
details, see :ref:`coherence_control`.
|
||||
|
||||
All allocation flags are independent and can be set in any combination. The only
|
||||
exception is setting ``hipHostMallocCoherent`` and ``hipHostMallocNonCoherent``
|
||||
together, which leads to an illegal state. An example of a valid flag
|
||||
combination is calling :cpp:func:`hipHostMalloc` with both
|
||||
``hipHostMallocPortable`` and ``hipHostMallocMapped`` flags set. Both the flags
|
||||
use the same model and differentiate only between how the surrounding code uses
|
||||
the host memory.
|
||||
|
||||
.. note::
|
||||
|
||||
By default, each GPU selects a Numa CPU node with the least Numa distance
|
||||
between them. This implies that the host memory is automatically allocated on
|
||||
the closest memory pool of the current GPU device's Numa node. Using
|
||||
:cpp:func:`hipSetDevice` API to set a different GPU increases the Numa
|
||||
distance but still allows you to access the host allocation.
|
||||
|
||||
Numa policy is implemented on Linux and is under development on Microsoft
|
||||
Windows.
|
||||
@@ -0,0 +1,580 @@
|
||||
.. meta::
|
||||
:description:
|
||||
:keywords: stream, memory allocation, SOMA, stream ordered memory allocator
|
||||
|
||||
.. _stream_ordered_memory_allocator_how-to:
|
||||
|
||||
*******************************************************************************
|
||||
Stream Ordered Memory Allocator
|
||||
*******************************************************************************
|
||||
|
||||
The Stream Ordered Memory Allocator (SOMA) is part of the HIP runtime API. SOMA provides an asynchronous memory allocation mechanism with stream-ordering semantics. You can use SOMA to allocate and free memory in stream order, which ensures that all asynchronous accesses occur between the stream executions of allocation and deallocation. Compliance with stream order prevents use-before-allocation or use-after-free errors, which helps to avoid an undefined behavior.
|
||||
|
||||
Advantages of SOMA:
|
||||
|
||||
- Efficient reuse: Enables efficient memory reuse across streams, which reduces unnecessary allocation overhead.
|
||||
- Fine-grained control: Allows you to set attributes and control caching behavior for memory pools.
|
||||
- Inter-process sharing: Enables secure sharing of allocations between processes.
|
||||
- Optimizations: Allows driver to optimize based on its awareness of SOMA and other stream management APIs.
|
||||
|
||||
Disadvantages of SOMA:
|
||||
|
||||
- Temporal constraints: Requires you to adhere strictly to stream order to avoid errors.
|
||||
- Complexity: Involves memory management in stream order, which can be intricate.
|
||||
- Learning curve: Requires you to put additional efforts to understand and utilize SOMA effectively.
|
||||
|
||||
Using SOMA
|
||||
=====================================
|
||||
|
||||
You can allocate memory using ``hipMallocAsync()`` with stream-ordered
|
||||
semantics. This restricts the asynchronous access to the memory between the stream executions of the allocation and deallocation. Accessing
|
||||
memory if the compliant memory accesses won't overlap
|
||||
temporally. ``hipFreeAsync()`` frees memory from the pool with stream-ordered
|
||||
semantics.
|
||||
|
||||
Here is how to use stream ordered memory allocation:
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: Stream Ordered Memory Allocation
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
// Kernel to perform some computation on allocated memory.
|
||||
__global__ void myKernel(int* data, size_t numElements) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
if (tid < numElements) {
|
||||
data[tid] = tid * 2;
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
// Initialize HIP.
|
||||
hipInit(0);
|
||||
|
||||
// Stream 0.
|
||||
constexpr hipStream_t streamId = 0;
|
||||
|
||||
// Allocate memory with stream ordered semantics.
|
||||
constexpr size_t numElements = 1024;
|
||||
int* devData;
|
||||
hipMallocAsync(&devData, numElements * sizeof(*devData), streamId);
|
||||
|
||||
// Launch the kernel to perform computation.
|
||||
dim3 blockSize(256);
|
||||
dim3 gridSize((numElements + blockSize.x - 1) / blockSize.x);
|
||||
myKernel<<<gridSize, blockSize>>>(devData, numElements);
|
||||
|
||||
// Copy data back to host.
|
||||
int* hostData = new int[numElements];
|
||||
hipMemcpy(hostData, devData, numElements * sizeof(*devData), hipMemcpyDeviceToHost);
|
||||
|
||||
// Print the array.
|
||||
for (size_t i = 0; i < numElements; ++i) {
|
||||
std::cout << "Element " << i << ": " << hostData[i] << std::endl;
|
||||
}
|
||||
|
||||
// Free memory with stream ordered semantics.
|
||||
hipFreeAsync(devData, streamId);
|
||||
delete[] hostData;
|
||||
|
||||
// Synchronize to ensure completion.
|
||||
hipDeviceSynchronize();
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: Ordinary Allocation
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
// Kernel to perform some computation on allocated memory.
|
||||
__global__ void myKernel(int* data, size_t numElements) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
if (tid < numElements) {
|
||||
data[tid] = tid * 2;
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
// Initialize HIP.
|
||||
hipInit(0);
|
||||
|
||||
// Allocate memory.
|
||||
constexpr size_t numElements = 1024;
|
||||
int* devData;
|
||||
hipMalloc(&devData, numElements * sizeof(*devData));
|
||||
|
||||
// Launch the kernel to perform computation.
|
||||
dim3 blockSize(256);
|
||||
dim3 gridSize((numElements + blockSize.x - 1) / blockSize.x);
|
||||
myKernel<<<gridSize, blockSize>>>(devData, numElements);
|
||||
|
||||
// Copy data back to host.
|
||||
int* hostData = new int[numElements];
|
||||
hipMemcpy(hostData, devData, numElements * sizeof(*devData), hipMemcpyDeviceToHost);
|
||||
|
||||
// Print the array.
|
||||
for (size_t i = 0; i < numElements; ++i) {
|
||||
std::cout << "Element " << i << ": " << hostData[i] << std::endl;
|
||||
}
|
||||
|
||||
// Free memory.
|
||||
hipFree(devData);
|
||||
delete[] hostData;
|
||||
|
||||
// Synchronize to ensure completion.
|
||||
hipDeviceSynchronize();
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
For more details, see :ref:`stream_ordered_memory_allocator_reference`.
|
||||
|
||||
Memory pools
|
||||
============
|
||||
|
||||
Memory pools provide a way to manage memory with stream-ordered behavior while ensuring proper synchronization and avoiding memory access errors. Division of a single memory system into separate pools facilitates querying the access path properties for each partition. Memory pools are used for host memory, device memory, and unified memory.
|
||||
|
||||
Set pools
|
||||
---------
|
||||
|
||||
The ``hipMallocAsync()`` function uses the current memory pool and also provides the opportunity to create and access different pools using ``hipMemPoolCreate()`` and ``hipMallocFromPoolAsync()`` functions respectively.
|
||||
|
||||
Unlike NVIDIA CUDA, where stream-ordered memory allocation can be implicit, ROCm HIP is explicit. This requires managing memory allocation for each stream in HIP while ensuring precise control over memory usage and synchronization.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
// Kernel to perform some computation on allocated memory.
|
||||
__global__ void myKernel(int* data, size_t numElements) {
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
if (tid < numElements) {
|
||||
data[tid] = tid * 2;
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
// Create a stream.
|
||||
hipStream_t stream;
|
||||
hipStreamCreate(&stream);
|
||||
|
||||
// Create a memory pool with default properties.
|
||||
hipMemPoolProps poolProps = {};
|
||||
poolProps.allocType = hipMemAllocationTypePinned;
|
||||
poolProps.handleTypes = hipMemHandleTypePosixFileDescriptor;
|
||||
poolProps.location.type = hipMemLocationTypeDevice;
|
||||
poolProps.location.id = 0; // Assuming device 0.
|
||||
|
||||
hipMemPool_t memPool;
|
||||
hipMemPoolCreate(&memPool, &poolProps);
|
||||
|
||||
// Allocate memory from the pool asynchronously.
|
||||
constexpr size_t numElements = 1024;
|
||||
int* devData = nullptr;
|
||||
hipMallocFromPoolAsync(&devData, numElements * sizeof(*devData), memPool, stream);
|
||||
|
||||
// Define grid and block sizes.
|
||||
dim3 blockSize(256);
|
||||
dim3 gridSize((numElements + blockSize.x - 1) / blockSize.x);
|
||||
|
||||
// Launch the kernel to perform computation.
|
||||
myKernel<<<gridSize, blockSize, 0, stream>>>(devData, numElements);
|
||||
|
||||
// Synchronize the stream.
|
||||
hipStreamSynchronize(stream);
|
||||
|
||||
// Copy data back to host.
|
||||
int* hostData = new int[numElements];
|
||||
hipMemcpy(hostData, devData, numElements * sizeof(*devData), hipMemcpyDeviceToHost);
|
||||
|
||||
// Print the array.
|
||||
for (size_t i = 0; i < numElements; ++i) {
|
||||
std::cout << "Element " << i << ": " << hostData[i] << std::endl;
|
||||
}
|
||||
|
||||
// Free the allocated memory.
|
||||
hipFreeAsync(devData, stream);
|
||||
|
||||
// Synchronize the stream again to ensure all operations are complete.
|
||||
hipStreamSynchronize(stream);
|
||||
|
||||
// Destroy the memory pool and stream.
|
||||
hipMemPoolDestroy(memPool);
|
||||
hipStreamDestroy(stream);
|
||||
|
||||
// Free host memory.
|
||||
delete[] hostData;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Trim pools
|
||||
----------
|
||||
|
||||
The memory allocator allows you to allocate and free memory in stream order. To control memory usage, set the release threshold attribute using ``hipMemPoolAttrReleaseThreshold``. This threshold specifies the amount of reserved memory in bytes to hold onto.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
uint64_t threshold = UINT64_MAX;
|
||||
hipMemPoolSetAttribute(memPool, hipMemPoolAttrReleaseThreshold, &threshold);
|
||||
|
||||
When the amount of memory held in the memory pool exceeds the threshold, the allocator tries to release memory back to the operating system during the next call to stream, event, or context synchronization.
|
||||
|
||||
To improve performance, it is a good practice to adjust the memory pool size using ``hipMemPoolTrimTo()``. It helps to reclaim memory from an excessive memory pool, which optimizes memory usage for your application.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
int main() {
|
||||
hipMemPool_t memPool;
|
||||
hipDevice_t device = 0; // Specify the device index.
|
||||
|
||||
// Initialize the device.
|
||||
hipSetDevice(device);
|
||||
|
||||
// Get the default memory pool for the device.
|
||||
hipDeviceGetDefaultMemPool(&memPool, device);
|
||||
|
||||
// Allocate memory from the pool (e.g., 1 MB).
|
||||
size_t allocSize = 1 * 1024 * 1024;
|
||||
void* ptr;
|
||||
hipMalloc(&ptr, allocSize);
|
||||
|
||||
// Free the allocated memory.
|
||||
hipFree(ptr);
|
||||
|
||||
// Trim the memory pool to a specific size (e.g., 512 KB).
|
||||
size_t newSize = 512 * 1024;
|
||||
hipMemPoolTrimTo(memPool, newSize);
|
||||
|
||||
// Clean up.
|
||||
hipMemPoolDestroy(memPool);
|
||||
|
||||
std::cout << "Memory pool trimmed to " << newSize << " bytes." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
Resource usage statistics
|
||||
-------------------------
|
||||
|
||||
Resource usage statistics help in optimization. Here is the list of pool attributes used to query memory usage:
|
||||
|
||||
- ``hipMemPoolAttrReservedMemCurrent``: Returns the total physical GPU memory currently held in the pool.
|
||||
- ``hipMemPoolAttrUsedMemCurrent``: Returns the total size of all the memory allocated from the pool.
|
||||
- ``hipMemPoolAttrReservedMemHigh``: Returns the total physical GPU memory held in the pool since the last reset.
|
||||
- ``hipMemPoolAttrUsedMemHigh``: Returns the total size of all the memory allocated from the pool since the last reset.
|
||||
|
||||
To reset these attributes to the current value, use ``hipMemPoolSetAttribute()``.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
// Sample helper functions for getting the usage statistics in bulk.
|
||||
struct usageStatistics {
|
||||
uint64_t reservedMemCurrent;
|
||||
uint64_t reservedMemHigh;
|
||||
uint64_t usedMemCurrent;
|
||||
uint64_t usedMemHigh;
|
||||
};
|
||||
|
||||
void getUsageStatistics(hipMemPool_t memPool, struct usageStatistics *statistics) {
|
||||
hipMemPoolGetAttribute(memPool, hipMemPoolAttrReservedMemCurrent, &statistics->reservedMemCurrent);
|
||||
hipMemPoolGetAttribute(memPool, hipMemPoolAttrReservedMemHigh, &statistics->reservedMemHigh);
|
||||
hipMemPoolGetAttribute(memPool, hipMemPoolAttrUsedMemCurrent, &statistics->usedMemCurrent);
|
||||
hipMemPoolGetAttribute(memPool, hipMemPoolAttrUsedMemHigh, &statistics->usedMemHigh);
|
||||
}
|
||||
|
||||
// Resetting the watermarks resets them to the current value.
|
||||
void resetStatistics(hipMemPool_t memPool) {
|
||||
uint64_t value = 0;
|
||||
hipMemPoolSetAttribute(memPool, hipMemPoolAttrReservedMemHigh, &value);
|
||||
hipMemPoolSetAttribute(memPool, hipMemPoolAttrUsedMemHigh, &value);
|
||||
}
|
||||
|
||||
int main() {
|
||||
hipMemPool_t memPool;
|
||||
hipDevice_t device = 0; // Specify the device index.
|
||||
|
||||
// Initialize the device.
|
||||
hipSetDevice(device);
|
||||
|
||||
// Get the default memory pool for the device.
|
||||
hipDeviceGetDefaultMemPool(&memPool, device);
|
||||
|
||||
// Allocate memory from the pool (e.g., 1 MB).
|
||||
size_t allocSize = 1 * 1024 * 1024;
|
||||
void* ptr;
|
||||
hipMalloc(&ptr, allocSize);
|
||||
|
||||
// Free the allocated memory.
|
||||
hipFree(ptr);
|
||||
|
||||
// Trim the memory pool to a specific size (e.g., 512 KB).
|
||||
size_t newSize = 512 * 1024;
|
||||
hipMemPoolTrimTo(memPool, newSize);
|
||||
|
||||
// Get and print usage statistics before resetting.
|
||||
usageStatistics statsBefore;
|
||||
getUsageStatistics(memPool, &statsBefore);
|
||||
std::cout << "Before resetting statistics:" << std::endl;
|
||||
std::cout << "Reserved Memory Current: " << statsBefore.reservedMemCurrent << " bytes" << std::endl;
|
||||
std::cout << "Reserved Memory High: " << statsBefore.reservedMemHigh << " bytes" << std::endl;
|
||||
std::cout << "Used Memory Current: " << statsBefore.usedMemCurrent << " bytes" << std::endl;
|
||||
std::cout << "Used Memory High: " << statsBefore.usedMemHigh << " bytes" << std::endl;
|
||||
|
||||
// Reset the statistics.
|
||||
resetStatistics(memPool);
|
||||
|
||||
// Get and print usage statistics after resetting.
|
||||
usageStatistics statsAfter;
|
||||
getUsageStatistics(memPool, &statsAfter);
|
||||
std::cout << "After resetting statistics:" << std::endl;
|
||||
std::cout << "Reserved Memory Current: " << statsAfter.reservedMemCurrent << " bytes" << std::endl;
|
||||
std::cout << "Reserved Memory High: " << statsAfter.reservedMemHigh << " bytes" << std::endl;
|
||||
std::cout << "Used Memory Current: " << statsAfter.usedMemCurrent << " bytes" << std::endl;
|
||||
std::cout << "Used Memory High: " << statsAfter.usedMemHigh << " bytes" << std::endl;
|
||||
|
||||
// Clean up.
|
||||
hipMemPoolDestroy(memPool);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Memory reuse policies
|
||||
---------------------
|
||||
|
||||
The allocator might reallocate memory as long as the compliant memory accesses will not to overlap temporally. To optimize the memory usage, disable or enable the following memory pool reuse policy attribute flags:
|
||||
|
||||
- ``hipMemPoolReuseFollowEventDependencies``: Checks event dependencies before allocating additional GPU memory.
|
||||
- ``hipMemPoolReuseAllowOpportunistic``: Checks freed allocations to determine if the stream order semantic indicated by the free operation has been met.
|
||||
- ``hipMemPoolReuseAllowInternalDependencies``: Manages reuse based on internal dependencies in runtime. If the driver fails to allocate and map additional physical memory, it searches for memory waiting for another stream's progress and reuses it.
|
||||
|
||||
Device accessibility for multi-GPU support
|
||||
------------------------------------------
|
||||
|
||||
Allocations are initially accessible from the device where they reside.
|
||||
|
||||
Interprocess memory handling
|
||||
=============================
|
||||
|
||||
Interprocess capable (IPC) memory pools facilitate efficient and secure sharing of GPU memory between processes.
|
||||
|
||||
To achieve interprocess memory sharing, you can use either :ref:`device pointer <device-pointer>` or :ref:`shareable handle <shareable-handle>`. Both provide allocator (export) and consumer (import) interfaces.
|
||||
|
||||
.. _device-pointer:
|
||||
|
||||
Device pointer
|
||||
--------------
|
||||
|
||||
To export data to share a memory pool pointer directly between processes, use ``hipMemPoolExportPointer()``. It allows you to share a memory allocation with another process.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <sys/stat.h>
|
||||
|
||||
int main() {
|
||||
// Allocate memory.
|
||||
void* devPtr;
|
||||
hipMalloc(&devPtr, sizeof(int));
|
||||
|
||||
// Export the memory pool pointer.
|
||||
hipMemPoolPtrExportData exportData;
|
||||
hipError_t result = hipMemPoolExportPointer(&exportData, devPtr);
|
||||
if (result != hipSuccess) {
|
||||
std::cerr << "Error exporting memory pool pointer: " << hipGetErrorString(result) << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Create a named pipe (FIFO).
|
||||
const char* fifoPath = "/tmp/myfifo"; // Change this to a unique path.
|
||||
mkfifo(fifoPath, 0666);
|
||||
|
||||
// Write the exported data to the named pipe.
|
||||
std::ofstream fifoStream(fifoPath, std::ios::out | std::ios::binary);
|
||||
fifoStream.write(reinterpret_cast<char*>(&exportData), sizeof(hipMemPoolPtrExportData));
|
||||
fifoStream.close();
|
||||
|
||||
// Clean up.
|
||||
hipFree(devPtr);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
To import a memory pool pointer directly from another process, use ``hipMemPoolImportPointer()``.
|
||||
|
||||
Here is how to read the pool exported in the preceding example:
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
int main() {
|
||||
// Considering that you have exported the memory pool pointer already.
|
||||
// Now, let's simulate reading the exported data from a named pipe (FIFO).
|
||||
const char* fifoPath = "/tmp/myfifo"; // Change this to a unique path.
|
||||
std::ifstream fifoStream(fifoPath, std::ios::in | std::ios::binary);
|
||||
|
||||
if (!fifoStream.is_open()) {
|
||||
std::cerr << "Error opening FIFO file: " << fifoPath << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Read the exported data.
|
||||
hipMemPoolPtrExportData importData;
|
||||
fifoStream.read(reinterpret_cast<char*>(&importData), sizeof(hipMemPoolPtrExportData));
|
||||
fifoStream.close();
|
||||
|
||||
if (fifoStream.fail()) {
|
||||
std::cerr << "Error reading from FIFO file." << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Create a memory pool with default properties.
|
||||
hipMemPoolProps poolProps = {};
|
||||
poolProps.allocType = hipMemAllocationTypePinned;
|
||||
poolProps.handleTypes = hipMemHandleTypePosixFileDescriptor;
|
||||
poolProps.location.type = hipMemLocationTypeDevice;
|
||||
poolProps.location.id = 0; // Assuming device 0.
|
||||
|
||||
hipMemPool_t memPool;
|
||||
hipMemPoolCreate(&memPool, &poolProps);
|
||||
|
||||
// Import the memory pool pointer.
|
||||
void* importedDevPtr;
|
||||
hipError_t result = hipMemPoolImportPointer(&importedDevPtr, memPool, &importData);
|
||||
if (result != hipSuccess) {
|
||||
std::cerr << "Error imported memory pool pointer: " << hipGetErrorString(result) << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Now you can use the importedDevPtr for your computations.
|
||||
|
||||
// Clean up (free the memory).
|
||||
hipFree(importedDevPtr);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. _shareable-handle:
|
||||
|
||||
Shareable handle
|
||||
----------------
|
||||
|
||||
To export a memory pool pointer to a shareable handle, use ``hipMemPoolExportToSharedHandle()``. This handle could be a file descriptor or a handle obtained from another process. The exported handle contains information about the memory pool, such as size, location, and other relevant details.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <sys/stat.h>
|
||||
|
||||
int main() {
|
||||
// Create a memory pool with default properties.
|
||||
hipMemPoolProps poolProps = {};
|
||||
poolProps.allocType = hipMemAllocationTypePinned;
|
||||
poolProps.handleTypes = hipMemHandleTypePosixFileDescriptor;
|
||||
poolProps.location.type = hipMemLocationTypeDevice;
|
||||
poolProps.location.id = 0; // Assuming device 0.
|
||||
|
||||
hipMemPool_t memPool;
|
||||
hipError_t poolResult = hipMemPoolCreate(&memPool, &poolProps);
|
||||
if (poolResult != hipSuccess) {
|
||||
std::cerr << "Error creating memory pool: " << hipGetErrorString(poolResult) << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Allocate memory from the memory pool.
|
||||
void* devPtr;
|
||||
hipMallocFromPoolAsync(&devPtr, sizeof(int), memPool, 0);
|
||||
|
||||
// Export the memory pool pointer.
|
||||
int descriptor;
|
||||
hipError_t result = hipMemPoolExportToShareableHandle(&descriptor, memPool, hipMemHandleTypePosixFileDescriptor, 0);
|
||||
if (result != hipSuccess) {
|
||||
std::cerr << "Error exporting memory pool pointer: " << hipGetErrorString(result) << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Create a named pipe (FIFO).
|
||||
const char* fifoPath = "/tmp/myfifo"; // Change this to a unique path.
|
||||
mkfifo(fifoPath, 0666);
|
||||
|
||||
// Write the exported data to the named pipe.
|
||||
std::ofstream fifoStream(fifoPath, std::ios::out | std::ios::binary);
|
||||
fifoStream.write(reinterpret_cast<char*>(&descriptor), sizeof(int));
|
||||
fifoStream.close();
|
||||
|
||||
// Clean up.
|
||||
hipFree(devPtr);
|
||||
hipMemPoolDestroy(memPool);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
To import and restore a memory pool pointer from a shareable handle, which could be a file descriptor or a handle obtained from another process, use ``hipMemPoolImportFromShareableHandle()``. The exported shareable handle data contains information about the memory pool, including its size, location, and other relevant details. Importing the handle provides a valid memory pointer to the same memory, which allows you to share memory across different contexts.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
int main() {
|
||||
// Considering that you have exported the memory pool pointer already.
|
||||
// Now, let's simulate reading the exported data from a named pipe (FIFO).
|
||||
const char* fifoPath = "/tmp/myfifo"; // Change this to a unique path
|
||||
std::ifstream fifoStream(fifoPath, std::ios::in | std::ios::binary);
|
||||
|
||||
if (!fifoStream.is_open()) {
|
||||
std::cerr << "Error opening FIFO file: " << fifoPath << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Read the exported data.
|
||||
int descriptor;
|
||||
fifoStream.read(reinterpret_cast<char*>(&descriptor), sizeof(int));
|
||||
fifoStream.close();
|
||||
|
||||
if (fifoStream.fail()) {
|
||||
std::cerr << "Error reading from FIFO file." << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Import the memory pool.
|
||||
hipMemPool_t memPool;
|
||||
hipError_t result = hipMemPoolImportFromShareableHandle(&memPool, &descriptor, hipMemHandleTypePosixFileDescriptor, 0);
|
||||
if (result != hipSuccess) {
|
||||
std::cerr << "Error importing memory pool: " << hipGetErrorString(result) << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Allocate memory from the imported memory pool.
|
||||
void* importedDevPtr;
|
||||
hipMallocFromPoolAsync(&importedDevPtr, sizeof(int), memPool, 0);
|
||||
|
||||
// Now you can use the importedDevPtr for your computations.
|
||||
|
||||
// Clean up (free the memory).
|
||||
hipFree(importedDevPtr);
|
||||
hipMemPoolDestroy(memPool);
|
||||
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,740 @@
|
||||
.. meta::
|
||||
:description: This chapter describes Unified Memory and shows
|
||||
how to use it in AMD HIP.
|
||||
:keywords: AMD, ROCm, HIP, CUDA, unified memory, unified, memory
|
||||
|
||||
.. _unified_memory:
|
||||
|
||||
*******************************************************************************
|
||||
Unified memory management
|
||||
*******************************************************************************
|
||||
|
||||
In conventional architectures CPUs and attached devices have their own memory
|
||||
space and dedicated physical memory backing it up, e.g. normal RAM for CPUs and
|
||||
VRAM on GPUs. This way each device can have physical memory optimized for its
|
||||
use case. GPUs usually have specialized memory whose bandwidth is a
|
||||
magnitude higher than the RAM attached to CPUs.
|
||||
|
||||
While providing exceptional performance, this setup typically requires explicit
|
||||
memory management, as memory needs to be allocated, copied and freed on the used
|
||||
devices and on the host. Additionally, this makes using more than the physically
|
||||
available memory on the devices complicated.
|
||||
|
||||
Modern GPUs circumvent the problem of having to explicitly manage the memory,
|
||||
while still keeping the benefits of the dedicated physical memories, by
|
||||
supporting the concept of unified memory. This enables the CPU and the GPUs in
|
||||
the system to access host and other GPUs' memory without explicit memory
|
||||
management.
|
||||
|
||||
Unified memory
|
||||
================================================================================
|
||||
|
||||
Unified Memory is a single memory address space accessible from any processor
|
||||
within a system. This setup simplifies memory management and enables
|
||||
applications to allocate data that can be read or written on both CPUs and GPUs
|
||||
without explicitly copying it to the specific CPU or GPU. The Unified memory
|
||||
model is shown in the following figure.
|
||||
|
||||
.. figure:: ../../../data/how-to/hip_runtime_api/memory_management/unified_memory/um.svg
|
||||
|
||||
Unified memory enables the access to memory located on other devices via
|
||||
several methods, depending on whether hardware support is available or has to be
|
||||
managed by the driver.
|
||||
|
||||
Hardware supported on-demand page migration
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
When a kernel on the device tries to access a memory address that is not in its
|
||||
memory, a page-fault is triggered. The GPU then in turn requests the page from
|
||||
the host or an other device, on which the memory is located. The page is then
|
||||
unmapped from the source, sent to the device and mapped to the device's memory.
|
||||
The requested memory is then available to the processes running on the device.
|
||||
|
||||
In case the device's memory is at capacity, a page is unmapped from the device's
|
||||
memory first and sent and mapped to host memory. This enables more memory to be
|
||||
allocated and used for a GPU, than the GPU itself has physically available.
|
||||
|
||||
This level of unified memory support can be very beneficial for sparse accesses
|
||||
to an array, that is not often used on the device.
|
||||
|
||||
Driver managed page migration
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
If the hardware does not support on-demand page migration, then all the pages
|
||||
accessed by a kernel have to be resident on the device, so they have to be
|
||||
migrated before the kernel is running. Since the driver can not know beforehand,
|
||||
what parts of an array are going to be accessed, all pages of all accessed
|
||||
arrays have to be migrated. This can lead to significant delays on the first run
|
||||
of a kernel, on top of possibly copying more memory than is actually accessed by
|
||||
the kernel.
|
||||
|
||||
.. _unified memory system requirements:
|
||||
|
||||
System requirements
|
||||
================================================================================
|
||||
|
||||
Unified memory is supported on Linux by all modern AMD GPUs from the Vega
|
||||
series onward, as shown in the following table. Unified memory management can
|
||||
be achieved by explicitly allocating managed memory using
|
||||
:cpp:func:`hipMallocManaged` or marking variables with the ``__managed__``
|
||||
attribute. For the latest GPUs, with a Linux kernel that supports
|
||||
`Heterogeneous Memory Management (HMM)
|
||||
<https://www.kernel.org/doc/html/latest/mm/hmm.html>`_, the normal system
|
||||
allocator can be used.
|
||||
|
||||
.. list-table:: Supported Unified Memory Allocators by GPU architecture
|
||||
:widths: 40, 25, 25
|
||||
:header-rows: 1
|
||||
:align: center
|
||||
|
||||
* - Architecture
|
||||
- :cpp:func:`hipMallocManaged()`, ``__managed__``
|
||||
- ``new``, ``malloc()``
|
||||
* - CDNA3
|
||||
- ✅
|
||||
- ✅ :sup:`1`
|
||||
* - CDNA2
|
||||
- ✅
|
||||
- ✅ :sup:`1`
|
||||
* - CDNA1
|
||||
- ✅
|
||||
- ✅ :sup:`1`
|
||||
* - RDNA1
|
||||
- ✅
|
||||
- ❌
|
||||
* - GCN5
|
||||
- ✅
|
||||
- ❌
|
||||
|
||||
✅: **Supported**
|
||||
|
||||
❌: **Unsupported**
|
||||
|
||||
:sup:`1` Works only with ``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>`_.
|
||||
|
||||
.. _unified memory allocators:
|
||||
|
||||
Unified memory allocators
|
||||
================================================================================
|
||||
|
||||
Support for the different unified memory allocators depends on the GPU
|
||||
architecture and on the system. For more information, see :ref:`unified memory
|
||||
system requirements` and :ref:`checking unified memory support`.
|
||||
|
||||
- **HIP allocated managed memory and variables**
|
||||
|
||||
:cpp:func:`hipMallocManaged()` is a dynamic memory allocator available on
|
||||
all GPUs with unified memory support. For more details, visit
|
||||
:ref:`unified_memory_reference`.
|
||||
|
||||
The ``__managed__`` declaration specifier, which serves as its counterpart,
|
||||
can be utilized for static allocation.
|
||||
|
||||
- **System allocated unified memory**
|
||||
|
||||
Starting with CDNA2, the ``new`` and ``malloc()`` system allocators allow
|
||||
you to reserve unified memory. The system allocator is more versatile and
|
||||
offers an easy transition for code written for CPUs to HIP code as the
|
||||
same system allocation API is used.
|
||||
|
||||
To ensure the proper functioning of system allocated unified memory on supported
|
||||
GPUs, it is essential to configure the environment variable ``XNACK=1`` and use
|
||||
a kernel 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>`_.
|
||||
|
||||
The table below illustrates the expected behavior of managed and unified memory
|
||||
functions on ROCm and CUDA, both with and without HMM support.
|
||||
|
||||
.. tab-set::
|
||||
.. tab-item:: ROCm allocation behaviour
|
||||
:sync: original-block
|
||||
|
||||
.. list-table:: Comparison of expected behavior of managed and unified memory functions in ROCm
|
||||
:widths: 26, 17, 20, 17, 20
|
||||
:header-rows: 1
|
||||
|
||||
* - call
|
||||
- Allocation origin without HMM or ``XNACK=0``
|
||||
- Access outside the origin without HMM or ``XNACK=0``
|
||||
- Allocation origin with HMM and ``XNACK=1``
|
||||
- Access outside the origin with HMM and ``XNACK=1``
|
||||
* - ``new``, ``malloc()``
|
||||
- host
|
||||
- not accessible on device
|
||||
- host
|
||||
- page-fault migration
|
||||
* - :cpp:func:`hipMalloc()`
|
||||
- device
|
||||
- zero copy [zc]_
|
||||
- device
|
||||
- zero copy [zc]_
|
||||
* - :cpp:func:`hipMallocManaged()`, ``__managed__``
|
||||
- pinned host
|
||||
- zero copy [zc]_
|
||||
- host
|
||||
- page-fault migration
|
||||
* - :cpp:func:`hipHostRegister()`
|
||||
- undefined behavior
|
||||
- undefined behavior
|
||||
- host
|
||||
- page-fault migration
|
||||
* - :cpp:func:`hipHostMalloc()`
|
||||
- pinned host
|
||||
- zero copy [zc]_
|
||||
- pinned host
|
||||
- zero copy [zc]_
|
||||
|
||||
.. tab-item:: CUDA allocation behaviour
|
||||
:sync: cooperative-groups
|
||||
|
||||
.. list-table:: Comparison of expected behavior of managed and unified memory functions in CUDA
|
||||
:widths: 26, 17, 20, 17, 20
|
||||
:header-rows: 1
|
||||
|
||||
* - call
|
||||
- Allocation origin without HMM
|
||||
- Access outside the origin without HMM
|
||||
- Allocation origin with HMM
|
||||
- Access outside the origin with HMM
|
||||
* - ``new``, ``malloc()``
|
||||
- host
|
||||
- not accessible on device
|
||||
- first touch
|
||||
- page-fault migration
|
||||
* - ``cudaMalloc()``
|
||||
- device
|
||||
- not accessible on host
|
||||
- device
|
||||
- page-fault migration
|
||||
* - ``cudaMallocManaged()``, ``__managed__``
|
||||
- host
|
||||
- page-fault migration
|
||||
- first touch
|
||||
- page-fault migration
|
||||
* - ``cudaHostRegister()``
|
||||
- host
|
||||
- page-fault migration
|
||||
- host
|
||||
- page-fault migration
|
||||
* - ``cudaMallocHost()``
|
||||
- pinned host
|
||||
- zero copy [zc]_
|
||||
- pinned host
|
||||
- zero copy [zc]_
|
||||
|
||||
.. [zc] Zero copy is a feature, where the memory is pinned to either the device
|
||||
or the host, and won't be transferred when accessed by another device or
|
||||
the host. Instead only the requested memory is transferred, without
|
||||
making an explicit copy, like a normal memory access, hence the term
|
||||
"zero copy".
|
||||
|
||||
.. _checking unified memory support:
|
||||
|
||||
Checking unified memory support
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The following device attributes can offer information about which :ref:`unified
|
||||
memory allocators` are supported. The attribute value is 1 if the functionality
|
||||
is supported, and 0 if it is not supported.
|
||||
|
||||
.. list-table:: Device attributes for unified memory management
|
||||
:widths: 40, 60
|
||||
:header-rows: 1
|
||||
:align: center
|
||||
|
||||
* - Attribute
|
||||
- Description
|
||||
* - :cpp:enumerator:`hipDeviceAttributeManagedMemory`
|
||||
- Device supports allocating managed memory on this system
|
||||
* - :cpp:enumerator:`hipDeviceAttributePageableMemoryAccess`
|
||||
- Device supports coherently accessing pageable memory without calling :cpp:func:`hipHostRegister()` on it.
|
||||
* - :cpp:enumerator:`hipDeviceAttributeConcurrentManagedAccess`
|
||||
- Full unified memory support. Device can coherently access managed memory concurrently with the CPU
|
||||
|
||||
For details on how to get the attributes of a specific device see :cpp:func:`hipDeviceGetAttribute()`.
|
||||
|
||||
Example for unified memory management
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The following example shows how to use unified memory with
|
||||
:cpp:func:`hipMallocManaged()` for dynamic allocation, the ``__managed__`` attribute
|
||||
for static allocation and the standard ``new`` allocation. For comparison, the
|
||||
explicit memory management example is presented in the last tab.
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: hipMallocManaged()
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 22-25
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
int main() {
|
||||
int *a, *b, *c;
|
||||
|
||||
// Allocate memory for a, b and c that is accessible to both device and host codes.
|
||||
HIP_CHECK(hipMallocManaged(&a, sizeof(*a)));
|
||||
HIP_CHECK(hipMallocManaged(&b, sizeof(*b)));
|
||||
HIP_CHECK(hipMallocManaged(&c, sizeof(*c)));
|
||||
|
||||
// Setup input values.
|
||||
*a = 1;
|
||||
*b = 2;
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, a, b, c);
|
||||
|
||||
// Wait for GPU to finish before accessing on host.
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Print the result.
|
||||
std::cout << *a << " + " << *b << " = " << *c << std::endl;
|
||||
|
||||
// Cleanup allocated memory.
|
||||
HIP_CHECK(hipFree(a));
|
||||
HIP_CHECK(hipFree(b));
|
||||
HIP_CHECK(hipFree(c));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: __managed__
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 19-20
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
// Declare a, b and c as static variables.
|
||||
__managed__ int a, b, c;
|
||||
|
||||
int main() {
|
||||
// Setup input values.
|
||||
a = 1;
|
||||
b = 2;
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, &a, &b, &c);
|
||||
|
||||
// Wait for GPU to finish before accessing on host.
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Prints the result.
|
||||
std::cout << a << " + " << b << " = " << c << std::endl;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: new
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 20-23
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int* a, int* b, int* c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
// This example requires HMM support and the environment variable HSA_XNACK needs to be set to 1
|
||||
int main() {
|
||||
// Allocate memory for a, b, and c.
|
||||
int *a = new int[1];
|
||||
int *b = new int[1];
|
||||
int *c = new int[1];
|
||||
|
||||
// Setup input values.
|
||||
*a = 1;
|
||||
*b = 2;
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, a, b, c);
|
||||
|
||||
// Wait for GPU to finish before accessing on host.
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Prints the result.
|
||||
std::cout << *a << " + " << *b << " = " << *c << std::endl;
|
||||
|
||||
// Cleanup allocated memory.
|
||||
delete[] a;
|
||||
delete[] b;
|
||||
delete[] c;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: Explicit Memory Management
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 27-34, 39-40
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
int main() {
|
||||
int a, b, c;
|
||||
int *d_a, *d_b, *d_c;
|
||||
|
||||
// Setup input values.
|
||||
a = 1;
|
||||
b = 2;
|
||||
|
||||
// Allocate device copies of a, b and c.
|
||||
HIP_CHECK(hipMalloc(&d_a, sizeof(*d_a)));
|
||||
HIP_CHECK(hipMalloc(&d_b, sizeof(*d_b)));
|
||||
HIP_CHECK(hipMalloc(&d_c, sizeof(*d_c)));
|
||||
|
||||
// Copy input values to device.
|
||||
HIP_CHECK(hipMemcpy(d_a, &a, sizeof(*d_a), hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(d_b, &b, sizeof(*d_b), hipMemcpyHostToDevice));
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, d_a, d_b, d_c);
|
||||
|
||||
// Copy the result back to the host.
|
||||
HIP_CHECK(hipMemcpy(&c, d_c, sizeof(*d_c), hipMemcpyDeviceToHost));
|
||||
|
||||
// Cleanup allocated memory.
|
||||
HIP_CHECK(hipFree(d_a));
|
||||
HIP_CHECK(hipFree(d_b));
|
||||
HIP_CHECK(hipFree(d_c));
|
||||
|
||||
// Prints the result.
|
||||
std::cout << a << " + " << b << " = " << c << std::endl;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. _using unified memory:
|
||||
|
||||
Using unified memory
|
||||
================================================================================
|
||||
|
||||
Unified memory can simplify the complexities of memory management in GPU
|
||||
computing, by not requiring explicit copies between the host and the devices. It
|
||||
can be particularly useful in use cases with sparse memory accesses from both
|
||||
the CPU and the GPU, as only the parts of the memory region that are actually
|
||||
accessed need to be transferred to the corresponding processor, not the whole
|
||||
memory region. This reduces the amount of memory sent over the PCIe bus or other
|
||||
interfaces.
|
||||
|
||||
In HIP, pinned memory allocations are coherent by default. Pinned memory is
|
||||
host memory mapped into the address space of all GPUs, meaning that the pointer
|
||||
can be used on both host and device. Additionally, using pinned memory instead of
|
||||
pageable memory on the host can improve bandwidth for transfers between the host
|
||||
and the GPUs.
|
||||
|
||||
While unified memory can provide numerous benefits, it's important to be aware
|
||||
of the potential performance overhead associated with unified memory. You must
|
||||
thoroughly test and profile your code to ensure it's the most suitable choice
|
||||
for your use case.
|
||||
|
||||
.. _unified memory runtime hints:
|
||||
|
||||
Performance optimizations for unified memory
|
||||
================================================================================
|
||||
|
||||
There are several ways, in which the developer can guide the runtime to reduce
|
||||
copies between devices, in order to improve performance.
|
||||
|
||||
Data prefetching
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Data prefetching is a technique used to improve the performance of your
|
||||
application by moving data to the desired device before it's actually
|
||||
needed. ``hipCpuDeviceId`` is a special constant to specify the CPU as target.
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 33-36,41-42
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
int main() {
|
||||
int *a, *b, *c;
|
||||
int deviceId;
|
||||
HIP_CHECK(hipGetDevice(&deviceId)); // Get the current device ID
|
||||
|
||||
// Allocate memory for a, b and c that is accessible to both device and host codes.
|
||||
HIP_CHECK(hipMallocManaged(&a, sizeof(*a)));
|
||||
HIP_CHECK(hipMallocManaged(&b, sizeof(*b)));
|
||||
HIP_CHECK(hipMallocManaged(&c, sizeof(*c)));
|
||||
|
||||
// Setup input values.
|
||||
*a = 1;
|
||||
*b = 2;
|
||||
|
||||
// Prefetch the data to the GPU device.
|
||||
HIP_CHECK(hipMemPrefetchAsync(a, sizeof(*a), deviceId, 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(b, sizeof(*b), deviceId, 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(c, sizeof(*c), deviceId, 0));
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, a, b, c);
|
||||
|
||||
// Prefetch the result back to the CPU.
|
||||
HIP_CHECK(hipMemPrefetchAsync(c, sizeof(*c), hipCpuDeviceId, 0));
|
||||
|
||||
// Wait for the prefetch operations to complete.
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Prints the result.
|
||||
std::cout << *a << " + " << *b << " = " << *c << std::endl;
|
||||
|
||||
// Cleanup allocated memory.
|
||||
HIP_CHECK(hipFree(a));
|
||||
HIP_CHECK(hipFree(b));
|
||||
HIP_CHECK(hipFree(c));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Memory advice
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
Unified memory runtime hints can be set with :cpp:func:`hipMemAdvise()` to help
|
||||
improve the performance of your code if you know the memory usage pattern. There
|
||||
are several different types of hints as specified in the enum
|
||||
:cpp:enum:`hipMemoryAdvise`, for example, whether a certain device mostly reads
|
||||
the memory region, where it should ideally be located, and even whether that
|
||||
specific memory region is accessed by a specific device.
|
||||
|
||||
For the best performance, profile your application to optimize the
|
||||
utilization of HIP runtime hints.
|
||||
|
||||
The effectiveness of :cpp:func:`hipMemAdvise()` comes from its ability to inform
|
||||
the runtime of the developer's intentions regarding memory usage. When the
|
||||
runtime has knowledge of the expected memory access patterns, it can make better
|
||||
decisions about data placement, leading to less transfers via the interconnect
|
||||
and thereby reduced latency and bandwidth requirements. However, the actual
|
||||
impact on performance can vary based on the specific use case and the system.
|
||||
|
||||
The following is the updated version of the example above with memory advice
|
||||
instead of prefetching.
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 29-41
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
int main() {
|
||||
int deviceId;
|
||||
HIP_CHECK(hipGetDevice(&deviceId));
|
||||
int *a, *b, *c;
|
||||
|
||||
// Allocate memory for a, b, and c accessible to both device and host codes.
|
||||
HIP_CHECK(hipMallocManaged(&a, sizeof(*a)));
|
||||
HIP_CHECK(hipMallocManaged(&b, sizeof(*b)));
|
||||
HIP_CHECK(hipMallocManaged(&c, sizeof(*c)));
|
||||
|
||||
// Set memory advice for a and b to be read, located on and accessed by the GPU.
|
||||
HIP_CHECK(hipMemAdvise(a, sizeof(*a), hipMemAdviseSetPreferredLocation, deviceId));
|
||||
HIP_CHECK(hipMemAdvise(a, sizeof(*a), hipMemAdviseSetAccessedBy, deviceId));
|
||||
HIP_CHECK(hipMemAdvise(a, sizeof(*a), hipMemAdviseSetReadMostly, deviceId));
|
||||
|
||||
HIP_CHECK(hipMemAdvise(b, sizeof(*b), hipMemAdviseSetPreferredLocation, deviceId));
|
||||
HIP_CHECK(hipMemAdvise(b, sizeof(*b), hipMemAdviseSetAccessedBy, deviceId));
|
||||
HIP_CHECK(hipMemAdvise(b, sizeof(*b), hipMemAdviseSetReadMostly, deviceId));
|
||||
|
||||
// Set memory advice for c to be read, located on and accessed by the CPU.
|
||||
HIP_CHECK(hipMemAdvise(c, sizeof(*c), hipMemAdviseSetPreferredLocation, hipCpuDeviceId));
|
||||
HIP_CHECK(hipMemAdvise(c, sizeof(*c), hipMemAdviseSetAccessedBy, hipCpuDeviceId));
|
||||
HIP_CHECK(hipMemAdvise(c, sizeof(*c), hipMemAdviseSetReadMostly, hipCpuDeviceId));
|
||||
|
||||
// Setup input values.
|
||||
*a = 1;
|
||||
*b = 2;
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, a, b, c);
|
||||
|
||||
// Wait for GPU to finish before accessing on host.
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Prints the result.
|
||||
std::cout << *a << " + " << *b << " = " << *c << std::endl;
|
||||
|
||||
// Cleanup allocated memory.
|
||||
HIP_CHECK(hipFree(a));
|
||||
HIP_CHECK(hipFree(b));
|
||||
HIP_CHECK(hipFree(c));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Memory range attributes
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
:cpp:func:`hipMemRangeGetAttribute()` allows you to query attributes of a given
|
||||
memory range. The attributes are given in :cpp:enum:`hipMemRangeAttribute`.
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 44-49
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
#define HIP_CHECK(expression) \
|
||||
{ \
|
||||
const hipError_t err = expression; \
|
||||
if(err != hipSuccess){ \
|
||||
std::cerr << "HIP error: " \
|
||||
<< hipGetErrorString(err) \
|
||||
<< " at " << __LINE__ << "\n"; \
|
||||
} \
|
||||
}
|
||||
|
||||
// Addition of two values.
|
||||
__global__ void add(int *a, int *b, int *c) {
|
||||
*c = *a + *b;
|
||||
}
|
||||
|
||||
int main() {
|
||||
int *a, *b, *c;
|
||||
unsigned int attributeValue;
|
||||
constexpr size_t attributeSize = sizeof(attributeValue);
|
||||
|
||||
int deviceId;
|
||||
HIP_CHECK(hipGetDevice(&deviceId));
|
||||
|
||||
// Allocate memory for a, b and c that is accessible to both device and host codes.
|
||||
HIP_CHECK(hipMallocManaged(&a, sizeof(*a)));
|
||||
HIP_CHECK(hipMallocManaged(&b, sizeof(*b)));
|
||||
HIP_CHECK(hipMallocManaged(&c, sizeof(*c)));
|
||||
|
||||
// Setup input values.
|
||||
*a = 1;
|
||||
*b = 2;
|
||||
|
||||
HIP_CHECK(hipMemAdvise(a, sizeof(*a), hipMemAdviseSetReadMostly, deviceId));
|
||||
|
||||
// Launch add() kernel on GPU.
|
||||
hipLaunchKernelGGL(add, dim3(1), dim3(1), 0, 0, a, b, c);
|
||||
|
||||
// Wait for GPU to finish before accessing on host.
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Query an attribute of the memory range.
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&attributeValue,
|
||||
attributeSize,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
a,
|
||||
sizeof(*a)));
|
||||
|
||||
// Prints the result.
|
||||
std::cout << *a << " + " << *b << " = " << *c << std::endl;
|
||||
std::cout << "The array a is" << (attributeValue == 1 ? "" : " NOT") << " set to hipMemRangeAttributeReadMostly" << std::endl;
|
||||
|
||||
// Cleanup allocated memory.
|
||||
HIP_CHECK(hipFree(a));
|
||||
HIP_CHECK(hipFree(b));
|
||||
HIP_CHECK(hipFree(c));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Asynchronously attach memory to a stream
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The :cpp:func:`hipStreamAttachMemAsync()` function attaches memory to a stream,
|
||||
which can reduce the amount of memory transferred, when managed memory is used.
|
||||
When the memory is attached to a stream using this function, it only gets
|
||||
transferred between devices, when a kernel that is launched on this stream needs
|
||||
access to the memory.
|
||||
@@ -0,0 +1,154 @@
|
||||
.. meta::
|
||||
:description: This chapter describes introduces Virtual Memory (VM) and shows
|
||||
how to use it in AMD HIP.
|
||||
:keywords: AMD, ROCm, HIP, CUDA, virtual memory, virtual, memory, UM, APU
|
||||
|
||||
.. _virtual_memory:
|
||||
|
||||
********************************************************************************
|
||||
Virtual memory management
|
||||
********************************************************************************
|
||||
|
||||
Memory management is important when creating high-performance applications in
|
||||
the HIP ecosystem. Both allocating and copying memory can result in bottlenecks,
|
||||
which can significantly impact performance.
|
||||
|
||||
Global memory allocation in HIP uses the C language style allocation function.
|
||||
This works fine for simple cases but can cause problems if your memory needs
|
||||
change. If you need to increase the size of your memory, you must allocate a
|
||||
second larger buffer and copy the data to it before you can free the original
|
||||
buffer. This increases overall memory usage and causes unnecessary ``memcpy``
|
||||
calls. Another solution is to allocate a larger buffer than you initially need.
|
||||
However, this isn't an efficient way to handle resources and doesn't solve the
|
||||
issue of reallocation when the extra buffer runs out.
|
||||
|
||||
Virtual memory management solves these memory management problems. It helps to
|
||||
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
|
||||
is separated into multiple steps using the :cpp:func:`hipMemCreate`,
|
||||
:cpp:func:`hipMemAddressReserve`, :cpp:func:`hipMemMap`, and
|
||||
:cpp:func:`hipMemSetAccess` functions. This guide explains what these functions
|
||||
do and how you can use them for virtual memory management.
|
||||
|
||||
Allocate physical memory
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
The first step is to allocate the physical memory itself with 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`
|
||||
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.
|
||||
|
||||
.. 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;
|
||||
prop.location.id = currentDev;
|
||||
hipMemGetAllocationGranularity(&granularity, &prop, HIP_MEM_ALLOC_GRANULARITY_MINIMUM);
|
||||
padded_size = ROUND_UP(size, granularity);
|
||||
hipMemCreate(&allocHandle, padded_size, &prop, 0);
|
||||
|
||||
Reserve virtual address range
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
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.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAddressReserve(&ptr, padded_size, 0, 0, 0);
|
||||
hipMemMap(ptr, padded_size, 0, allocHandle, 0);
|
||||
|
||||
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
|
||||
:cpp:struct:`hipMemAccessDesc` descriptor as parameters. In a multi-GPU
|
||||
environment, you can map the device memory of one GPU to another. This feature
|
||||
also works with the traditional memory management system, but isn't as scalable
|
||||
as with virtual memory. When memory is allocated with :cpp:func:`hipMalloc`,
|
||||
:cpp:func:`hipDeviceEnablePeerAccess` is used to enable peer access. This
|
||||
function enables access between two devices, but it means that every call to
|
||||
:cpp:func:`hipMalloc` takes more time to perform the checks and the mapping
|
||||
between the devices. When using virtual memory management, peer access is
|
||||
enabled by :cpp:func:`hipMemSetAccess`, which provides a finer level of
|
||||
control over what is shared. This has no performance impact on memory allocation
|
||||
and gives you more control over what memory buffers are shared with which
|
||||
devices.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemAccessDesc accessDesc = {};
|
||||
accessDesc.location.type = HIP_MEM_LOCATION_TYPE_DEVICE;
|
||||
accessDesc.location.id = currentDev;
|
||||
accessDesc.flags = HIP_MEM_ACCESS_FLAGS_PROT_READWRITE;
|
||||
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.
|
||||
|
||||
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
|
||||
virtual address range, use :cpp:func:`hipMemAddressFree`. Finally, to release
|
||||
the physical memory, use :cpp:func:`hipMemRelease`. A side effect of these
|
||||
functions is the lack of synchronization when memory is released. If you call
|
||||
:cpp:func:`hipFree` when you have multiple streams running in parallel, it
|
||||
synchronizes the device. This causes worse resource usage and performance.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
hipMemUnmap(ptr, size);
|
||||
hipMemRelease(allocHandle);
|
||||
hipMemAddressFree(ptr, size);
|
||||
|
||||
.. _usage_virtual_memory:
|
||||
|
||||
Memory usage
|
||||
================================================================================
|
||||
|
||||
Dynamically increase allocation size
|
||||
--------------------------------------------------------------------------------
|
||||
|
||||
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.
|
||||
|
||||
.. 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.
|
||||
@@ -0,0 +1,420 @@
|
||||
.. meta::
|
||||
:description: This chapter describes how to use multiple devices on one host.
|
||||
:keywords: ROCm, HIP, multi-device, multiple, GPUs, devices
|
||||
|
||||
.. _multi-device:
|
||||
|
||||
*******************************************************************************
|
||||
Multi-device management
|
||||
*******************************************************************************
|
||||
|
||||
Device enumeration
|
||||
===============================================================================
|
||||
|
||||
Device enumeration involves identifying all the available GPUs connected to the
|
||||
host system. A single host machine can have multiple GPUs, each with its own
|
||||
unique identifier. By listing these devices, you can decide which GPU to use
|
||||
for computation. The host queries the system to count and list all connected
|
||||
GPUs that support the chosen ``HIP_PLATFORM``, ensuring that the application
|
||||
can leverage the full computational power available. Typically, applications
|
||||
list devices and their properties for deployment planning, and also make
|
||||
dynamic selections during runtime to ensure optimal performance.
|
||||
|
||||
If the application does not define a specific GPU, device 0 is selected.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
int main()
|
||||
{
|
||||
int deviceCount;
|
||||
hipGetDeviceCount(&deviceCount);
|
||||
std::cout << "Number of devices: " << deviceCount << std::endl;
|
||||
|
||||
for (int deviceId = 0; deviceId < deviceCount; ++deviceId)
|
||||
{
|
||||
hipDeviceProp_t deviceProp;
|
||||
hipGetDeviceProperties(&deviceProp, deviceId);
|
||||
std::cout << "Device " << deviceId << std::endl << " Properties:" << std::endl;
|
||||
std::cout << " Name: " << deviceProp.name << std::endl;
|
||||
std::cout << " Total Global Memory: " << deviceProp.totalGlobalMem / (1024 * 1024) << " MiB" << std::endl;
|
||||
std::cout << " Shared Memory per Block: " << deviceProp.sharedMemPerBlock / 1024 << " KiB" << std::endl;
|
||||
std::cout << " Registers per Block: " << deviceProp.regsPerBlock << std::endl;
|
||||
std::cout << " Warp Size: " << deviceProp.warpSize << std::endl;
|
||||
std::cout << " Max Threads per Block: " << deviceProp.maxThreadsPerBlock << std::endl;
|
||||
std::cout << " Max Threads per Multiprocessor: " << deviceProp.maxThreadsPerMultiProcessor << std::endl;
|
||||
std::cout << " Number of Multiprocessors: " << deviceProp.multiProcessorCount << std::endl;
|
||||
std::cout << " Max Threads Dimensions: ["
|
||||
<< deviceProp.maxThreadsDim[0] << ", "
|
||||
<< deviceProp.maxThreadsDim[1] << ", "
|
||||
<< deviceProp.maxThreadsDim[2] << "]" << std::endl;
|
||||
std::cout << " Max Grid Size: ["
|
||||
<< deviceProp.maxGridSize[0] << ", "
|
||||
<< deviceProp.maxGridSize[1] << ", "
|
||||
<< deviceProp.maxGridSize[2] << "]" << std::endl;
|
||||
std::cout << std::endl;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. _multi_device_selection:
|
||||
|
||||
Device selection
|
||||
===============================================================================
|
||||
|
||||
Once you have enumerated the available GPUs, the next step is to select a
|
||||
specific device for computation. This involves setting the active GPU that will
|
||||
execute subsequent operations. This step is crucial in multi-GPU systems where
|
||||
different GPUs might have different capabilities or workloads. By selecting the
|
||||
appropriate device, you ensure that the computational tasks are directed to the
|
||||
correct GPU, optimizing performance and resource utilization.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
exit(status); \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void simpleKernel(double *data)
|
||||
{
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
data[idx] = idx * 2.0;
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
double* deviceData0;
|
||||
double* deviceData1;
|
||||
size_t size = 1024 * sizeof(*deviceData0);
|
||||
|
||||
int deviceId0 = 0;
|
||||
int deviceId1 = 1;
|
||||
|
||||
// Set device 0 and perform operations
|
||||
HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
|
||||
HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
|
||||
simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Set device 1 and perform operations
|
||||
HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
|
||||
HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
|
||||
simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Copy result from device 0
|
||||
double hostData0[1024];
|
||||
HIP_CHECK(hipSetDevice(deviceId0));
|
||||
HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Copy result from device 1
|
||||
double hostData1[1024];
|
||||
HIP_CHECK(hipSetDevice(deviceId1));
|
||||
HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Display results from both devices
|
||||
std::cout << "Device 0 data: " << hostData0[0] << std::endl;
|
||||
std::cout << "Device 1 data: " << hostData1[0] << std::endl;
|
||||
|
||||
// Free device memory
|
||||
HIP_CHECK(hipFree(deviceData0));
|
||||
HIP_CHECK(hipFree(deviceData1));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
Stream and event behavior
|
||||
===============================================================================
|
||||
|
||||
In a multi-device system, streams and events are essential for efficient
|
||||
parallel computation and synchronization. Streams enable asynchronous task
|
||||
execution, allowing multiple devices to process data concurrently without
|
||||
blocking one another. Events provide a mechanism for synchronizing operations
|
||||
across streams and devices, ensuring that tasks on one device are completed
|
||||
before dependent tasks on another device begin. This coordination prevents race
|
||||
conditions and optimizes data flow in multi-GPU systems. Together, streams and
|
||||
events maximize performance by enabling parallel execution, load balancing, and
|
||||
effective resource utilization across heterogeneous hardware.
|
||||
|
||||
.. code-block:: cpp
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream>
|
||||
|
||||
__global__ void simpleKernel(double *data)
|
||||
{
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
data[idx] = idx * 2.0;
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
int numDevices;
|
||||
hipGetDeviceCount(&numDevices);
|
||||
|
||||
if (numDevices < 2) {
|
||||
std::cerr << "This example requires at least two GPUs." << std::endl;
|
||||
return -1;
|
||||
}
|
||||
|
||||
double *deviceData0, *deviceData1;
|
||||
size_t size = 1024 * sizeof(*deviceData0);
|
||||
|
||||
// Create streams and events for each device
|
||||
hipStream_t stream0, stream1;
|
||||
hipEvent_t startEvent0, stopEvent0, startEvent1, stopEvent1;
|
||||
|
||||
// Initialize device 0
|
||||
hipSetDevice(0);
|
||||
hipStreamCreate(&stream0);
|
||||
hipEventCreate(&startEvent0);
|
||||
hipEventCreate(&stopEvent0);
|
||||
hipMalloc(&deviceData0, size);
|
||||
|
||||
// Initialize device 1
|
||||
hipSetDevice(1);
|
||||
hipStreamCreate(&stream1);
|
||||
hipEventCreate(&startEvent1);
|
||||
hipEventCreate(&stopEvent1);
|
||||
hipMalloc(&deviceData1, size);
|
||||
|
||||
// Record the start event on device 0
|
||||
hipSetDevice(0);
|
||||
hipEventRecord(startEvent0, stream0);
|
||||
|
||||
// Launch the kernel asynchronously on device 0
|
||||
simpleKernel<<<1000, 128, 0, stream0>>>(deviceData0);
|
||||
|
||||
// Record the stop event on device 0
|
||||
hipEventRecord(stopEvent0, stream0);
|
||||
|
||||
// Wait for the stop event on device 0 to complete
|
||||
hipEventSynchronize(stopEvent0);
|
||||
|
||||
// Record the start event on device 1
|
||||
hipSetDevice(1);
|
||||
hipEventRecord(startEvent1, stream1);
|
||||
|
||||
// Launch the kernel asynchronously on device 1
|
||||
simpleKernel<<<1000, 128, 0, stream1>>>(deviceData1);
|
||||
|
||||
// Record the stop event on device 1
|
||||
hipEventRecord(stopEvent1, stream1);
|
||||
|
||||
// Wait for the stop event on device 1 to complete
|
||||
hipEventSynchronize(stopEvent1);
|
||||
|
||||
// Calculate elapsed time between the events for both devices
|
||||
float milliseconds0 = 0, milliseconds1 = 0;
|
||||
hipEventElapsedTime(&milliseconds0, startEvent0, stopEvent0);
|
||||
hipEventElapsedTime(&milliseconds1, startEvent1, stopEvent1);
|
||||
|
||||
std::cout << "Elapsed time on GPU 0: " << milliseconds0 << " ms" << std::endl;
|
||||
std::cout << "Elapsed time on GPU 1: " << milliseconds1 << " ms" << std::endl;
|
||||
|
||||
// Cleanup for device 0
|
||||
hipSetDevice(0);
|
||||
hipEventDestroy(startEvent0);
|
||||
hipEventDestroy(stopEvent0);
|
||||
hipStreamSynchronize(stream0);
|
||||
hipStreamDestroy(stream0);
|
||||
hipFree(deviceData0);
|
||||
|
||||
// Cleanup for device 1
|
||||
hipSetDevice(1);
|
||||
hipEventDestroy(startEvent1);
|
||||
hipEventDestroy(stopEvent1);
|
||||
hipStreamSynchronize(stream1);
|
||||
hipStreamDestroy(stream1);
|
||||
hipFree(deviceData1);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
Peer-to-peer memory access
|
||||
===============================================================================
|
||||
|
||||
In multi-GPU systems, peer-to-peer memory access enables one GPU to directly
|
||||
read or write to the memory of another GPU. This capability reduces data
|
||||
transfer times by allowing GPUs to communicate directly without involving the
|
||||
host. Enabling peer-to-peer access can significantly improve the performance of
|
||||
applications that require frequent data exchange between GPUs, as it eliminates
|
||||
the need to transfer data through the host memory.
|
||||
|
||||
By adding peer-to-peer access to the example referenced in
|
||||
:ref:`multi_device_selection`, data can be copied between devices:
|
||||
|
||||
.. tab-set::
|
||||
|
||||
.. tab-item:: with peer-to-peer
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 31-37, 51-55
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
exit(status); \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void simpleKernel(double *data)
|
||||
{
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
data[idx] = idx * 2.0;
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
double* deviceData0;
|
||||
double* deviceData1;
|
||||
size_t size = 1024 * sizeof(*deviceData0);
|
||||
|
||||
int deviceId0 = 0;
|
||||
int deviceId1 = 1;
|
||||
|
||||
// Enable peer access to the memory (allocated and future) on the peer device.
|
||||
// Ensure the device is active before enabling peer access.
|
||||
hipSetDevice(deviceId0);
|
||||
hipDeviceEnablePeerAccess(deviceId1, 0);
|
||||
|
||||
hipSetDevice(deviceId1);
|
||||
hipDeviceEnablePeerAccess(deviceId0, 0);
|
||||
|
||||
// Set device 0 and perform operations
|
||||
HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
|
||||
HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
|
||||
simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Set device 1 and perform operations
|
||||
HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
|
||||
HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
|
||||
simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Use peer-to-peer access
|
||||
hipSetDevice(deviceId0);
|
||||
|
||||
// Now device 0 can access memory allocated on device 1
|
||||
hipMemcpy(deviceData0, deviceData1, size, hipMemcpyDeviceToDevice);
|
||||
|
||||
// Copy result from device 0
|
||||
double hostData0[1024];
|
||||
HIP_CHECK(hipSetDevice(deviceId0));
|
||||
HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Copy result from device 1
|
||||
double hostData1[1024];
|
||||
HIP_CHECK(hipSetDevice(deviceId1));
|
||||
HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Display results from both devices
|
||||
std::cout << "Device 0 data: " << hostData0[0] << std::endl;
|
||||
std::cout << "Device 1 data: " << hostData1[0] << std::endl;
|
||||
|
||||
// Free device memory
|
||||
HIP_CHECK(hipFree(deviceData0));
|
||||
HIP_CHECK(hipFree(deviceData1));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
.. tab-item:: without peer-to-peer
|
||||
|
||||
.. code-block:: cpp
|
||||
:emphasize-lines: 43-49, 53, 58
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#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; \
|
||||
exit(status); \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void simpleKernel(double *data)
|
||||
{
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
data[idx] = idx * 2.0;
|
||||
}
|
||||
|
||||
int main()
|
||||
{
|
||||
double* deviceData0;
|
||||
double* deviceData1;
|
||||
size_t size = 1024 * sizeof(*deviceData0);
|
||||
|
||||
int deviceId0 = 0;
|
||||
int deviceId1 = 1;
|
||||
|
||||
// Set device 0 and perform operations
|
||||
HIP_CHECK(hipSetDevice(deviceId0)); // Set device 0 as current
|
||||
HIP_CHECK(hipMalloc(&deviceData0, size)); // Allocate memory on device 0
|
||||
simpleKernel<<<1000, 128>>>(deviceData0); // Launch kernel on device 0
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Set device 1 and perform operations
|
||||
HIP_CHECK(hipSetDevice(deviceId1)); // Set device 1 as current
|
||||
HIP_CHECK(hipMalloc(&deviceData1, size)); // Allocate memory on device 1
|
||||
simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
// Attempt to use deviceData0 on device 1 (This will not work as deviceData0 is allocated on device 0)
|
||||
HIP_CHECK(hipSetDevice(deviceId1));
|
||||
hipError_t err = hipMemcpy(deviceData1, deviceData0, size, hipMemcpyDeviceToDevice); // This should fail
|
||||
if (err != hipSuccess)
|
||||
{
|
||||
std::cout << "Error: Cannot access deviceData0 from device 1, deviceData0 is on device 0" << std::endl;
|
||||
}
|
||||
|
||||
// Copy result from device 0
|
||||
double hostData0[1024];
|
||||
HIP_CHECK(hipSetDevice(deviceId0));
|
||||
HIP_CHECK(hipMemcpy(hostData0, deviceData0, size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Copy result from device 1
|
||||
double hostData1[1024];
|
||||
HIP_CHECK(hipSetDevice(deviceId1));
|
||||
HIP_CHECK(hipMemcpy(hostData1, deviceData1, size, hipMemcpyDeviceToHost));
|
||||
|
||||
// Display results from both devices
|
||||
std::cout << "Device 0 data: " << hostData0[0] << std::endl;
|
||||
std::cout << "Device 1 data: " << hostData1[0] << std::endl;
|
||||
|
||||
// Free device memory
|
||||
HIP_CHECK(hipFree(deviceData0));
|
||||
HIP_CHECK(hipFree(deviceData1));
|
||||
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,94 @@
|
||||
.. meta::
|
||||
:description: HIP provides an OpenGL interoperability API that allows
|
||||
efficient data sharing between HIP's computing power and
|
||||
OpenGL's graphics rendering.
|
||||
:keywords: AMD, ROCm, HIP, OpenGL, interop, interoperability
|
||||
|
||||
*******************************************************************************
|
||||
OpenGL interoperability
|
||||
*******************************************************************************
|
||||
|
||||
The HIP--OpenGL interoperation involves mapping OpenGL resources, such as
|
||||
buffers and textures, for HIP to interact with OpenGL. This mapping process
|
||||
enables HIP to utilize these resources directly, bypassing the need for costly
|
||||
data transfers between the CPU and GPU. This capability is useful in
|
||||
applications that require both intensive GPU computation and real-time
|
||||
visualization.
|
||||
|
||||
The graphics resources must be registered using functions like
|
||||
:cpp:func:`hipGraphicsGLRegisterBuffer` or :cpp:func:`hipGraphicsGLRegisterImage`
|
||||
then they can be mapped to HIP with :cpp:func:`hipGraphicsMapResources`
|
||||
function.
|
||||
|
||||
After mapping, the :cpp:func:`hipGraphicsResourceGetMappedPointer` or
|
||||
:cpp:func:`hipGraphicsSubResourceGetMappedArray` functions used to retrieve a
|
||||
device pointer to the mapped resource, which can then be used in HIP kernels.
|
||||
|
||||
Unmapping resources with :cpp:func:`hipGraphicsUnmapResources` after
|
||||
computations ensure proper resource management.
|
||||
|
||||
Example
|
||||
===============================================================================
|
||||
|
||||
ROCm examples have a `HIP--OpenGL interoperation example <https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/opengl_interop>`_,
|
||||
where a simple HIP kernel is used to simulate a sine wave and rendered to a
|
||||
window as a grid of triangles using OpenGL. For a working example, there are
|
||||
multiple initialization steps needed like creating and opening a window,
|
||||
initializing OpenGL or selecting the OpenGL-capable device. After the
|
||||
initialization in the example, the kernel simulates the sinewave and updates
|
||||
the window's framebuffer in a cycle until the window is closed.
|
||||
|
||||
.. note::
|
||||
|
||||
The more recent OpenGL functions are loaded with `OpenGL loader <https://github.com/ROCm/rocm-examples/tree/develop/External/glad>`_,
|
||||
as these are not loaded by default on all platforms. The use of a custom
|
||||
loader is shown in the following example
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/opengl_interop.hip
|
||||
:start-after: // [Sphinx opengl functions load start]
|
||||
:end-before: // [Sphinx opengl functions load end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
The OpenGL buffer is imported to HIP in the following way:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/opengl_interop.hip
|
||||
:start-after: // [Sphinx buffer register and get start]
|
||||
:end-before: // [Sphinx buffer register and get end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
The imported pointer is manipulated in the sinewave kernel as shown in the
|
||||
following example:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/opengl_interop.hip
|
||||
:start-after: /// [Sphinx sinewave kernel start]
|
||||
:end-before: /// [Sphinx sinewave kernel end]
|
||||
:language: cpp
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/opengl_interop.hip
|
||||
:start-after: // [Sphinx buffer use in kernel start]
|
||||
:end-before: // [Sphinx buffer use in kernel end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
|
||||
The HIP graphics resource that is imported from the OpenGL buffer and is not
|
||||
needed anymore should be unmapped and unregistered as shown in the following way:
|
||||
|
||||
.. <!-- spellcheck-disable -->
|
||||
|
||||
.. literalinclude:: ../../tools/example_codes/opengl_interop.hip
|
||||
:start-after: // [Sphinx unregister start]
|
||||
:end-before: // [Sphinx unregister end]
|
||||
:language: cpp
|
||||
|
||||
.. <!-- spellcheck-enable -->
|
||||
Reference in New Issue
Block a user