AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-490062 - Update documentation

Telusuri Sumber 
Change-Id: Ib5297fdda2e05795b3b20436cc1de962e310b08b
This commit is contained in:
Istvan Kiss
2024-10-21 16:50:09 +02:00
committed by Istvan Kiss
orang tua 36739655e4
melakukan 3d60bd3a64
108 mengubah file dengan 14036 tambahan dan 2987 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
docs/tutorial/cooperative_groups_tutorial.rst
+240
Melihat File
@@ -0,0 +1,240 @@
.. meta::
:description: HIP cooperative groups tutorial
:keywords: AMD, ROCm, HIP, cooperative groups, tutorial
*******************************************************************************
Cooperative groups
*******************************************************************************
This tutorial demonstrates the basic concepts of cooperative groups in the HIP (Heterogeneous-computing Interface for Portability) programming model and the most essential tooling supporting it. This topic also reviews the commonalities of heterogeneous APIs. Familiarity with the C/C++ compilation model and the language is assumed.
Prerequisites
=============
To follow this tutorial, you'll need properly installed drivers and a HIP compiler toolchain to compile your code. Because ROCm HIP supports compiling and running on Linux and Microsoft Windows with AMD and NVIDIA GPUs, review the HIP development package installation before starting this tutorial. For more information, see :doc:`/install/install`.
Simple HIP Code
===============
To become familiar with heterogeneous programming, review the :doc:`SAXPY tutorial <saxpy>` and the first HIP code subsection. Compiling is also described in that tutorial.
Tiled partition
===============
You can use tiled partition to calculate the sum of ``partition_size`` length sequences and the sum of ``result_size``/ ``BlockSize`` length sequences. The host-side reference implementation is the following:
.. code-block:: cpp
// Host-side function to perform the same reductions as executed on the GPU
std::vector<unsigned int> ref_reduced(const unsigned int partition_size,
std::vector<unsigned int> input)
{
const unsigned int input_size = input.size();
const unsigned int result_size = input_size / partition_size;
std::vector<unsigned int> result(result_size);
for(unsigned int i = 0; i < result_size; i++)
{
unsigned int partition_result = 0;
for(unsigned int j = 0; j < partition_size; j++)
{
partition_result += input[partition_size * i + j];
}
result[i] = partition_result;
}
return result;
}
Device-side code
----------------
To calculate the sum of the sets of numbers, the tutorial uses the shared memory-based reduction on the device side. The warp level intrinsics usage is not covered in this tutorial, unlike in the :doc:`reduction tutorial. <reduction>` ``x`` input variable is a shared pointer, which needs to be synchronized after every value change. The ``thread_group`` input parameter can be ``thread_block_tile`` or ``thread_block`` because the ``thread_group`` is the parent class of these types. The ``val`` are the numbers to calculate the sum of. The returned results of this function return the final results of the reduction on thread ID 0 of the ``thread_group``, and for every other thread, the function results are 0.
.. code-block:: cuda
/// \brief Summation of `unsigned int val`'s in `thread_group g` using shared memory `x`
__device__ unsigned int reduce_sum(thread_group g, unsigned int* x, unsigned int val)
{
// Rank of this thread in the group
const unsigned int group_thread_id = g.thread_rank();
// We start with half the group size as active threads
// 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 for this thread in a shared, temporary array
x[group_thread_id] = val;
// Synchronize all threads in the group
g.sync();
// If our thread is still active, sum with its counterpart in the other half
if(group_thread_id < i)
{
val += x[group_thread_id + i];
}
// Synchronize all threads in the group
g.sync();
}
// Only the first thread returns a valid value
if(g.thread_rank() == 0)
return val;
else
return 0;
}
The ``reduce_sum`` device function is reused to calculate the block and custom
partition sum of the input numbers. The kernel has three sections:
1. Initialization of the reduction function variables.
2. The reduction of thread block and store the results in global memory.
3. The reduction of custom partition and store the results in global memory.
1. Initialization of the reduction function variables
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
In this code section, the shared memory is declared, the thread_block_group and
custom_partition are defined, and the input variables are loaded from global
memory.
.. code-block:: cuda
// threadBlockGroup consists of all threads in the block
thread_block thread_block_group = this_thread_block();
// Workspace array in shared memory required for reduction
__shared__ unsigned int workspace[2048];
unsigned int output;
// Input to reduce
const unsigned int input = d_vector[thread_block_group.thread_rank()];
// ...
// Every custom_partition group consists of 16 threads
thread_block_tile<PartitionSize> custom_partition
= tiled_partition<PartitionSize>(thread_block_group);
2. The reduction of thread block
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
In this code section, the sum is calculated on ``thread_block_group`` level, then the results are stored in global memory.
.. code-block:: cuda
// Perform reduction
output = reduce_sum(thread_block_group, workspace, input);
// Only the first thread returns a valid value
if(thread_block_group.thread_rank() == 0)
{
d_block_reduced_vector[0] = output;
}
3. The reduction of custom partition
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
In this code section, the sum is calculated on the custom partition level, then the results are stored in global memory. The custom partition is a partial block of the thread block, it means the reduction calculates on a shorter sequence of input numbers than at the ``thread_block_group`` case.
.. code-block:: cuda
// Perform reduction
output = reduce_sum(custom_partition, &workspace[group_offset], input);
// Only the first thread in each partition returns a valid value
if(custom_partition.thread_rank() == 0)
{
const unsigned int partition_id = thread_block_group.thread_rank() / PartitionSize;
d_partition_reduced_vector[partition_id] = output;
}
Host-side code
--------------
On the host-side, the following steps are done in the example:
1. Confirm the cooperative group support on AMD GPUs.
2. Initialize the cooperative group configuration.
3. Allocate and copy input to global memory.
4. Launch the cooperative kernel.
5. Save the results from global memory.
6. Free the global memory.
Only the first, second and fourth steps are important from the cooperative groups aspect, that's why those steps are detailed further.
1. Confirm the cooperative group support on AMD GPUs
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
Not all AMD GPUs support cooperative groups. You can confirm support with the following code:
.. code-block:: cpp
#ifdef __HIP_PLATFORM_AMD__
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;
}
#endif
2. Initialize the cooperative group configuration
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
In the example, there is only one block in the grid, and the ``threads_per_block`` must be dividable with ``partition_size``.
.. code-block:: cpp
// Number of blocks to launch.
constexpr unsigned int num_blocks = 1;
// Number of threads in each kernel block.
constexpr unsigned int threads_per_block = 64;
// Total element count of the input vector.
constexpr unsigned int size = num_blocks * threads_per_block;
// Total elements count of a tiled_partition.
constexpr unsigned int partition_size = 16;
// Total size (in bytes) of the input vector.
constexpr size_t size_bytes = sizeof(unsigned int) * size;
static_assert(threads_per_block % partition_size == 0,
"threads_per_block must be a multiple of partition_size");
4. Launch the kernel
~~~~~~~~~~~~~~~~~~~~
The kernel launch is done with the ``hipLaunchCooperativeKernel`` of the cooperative groups API.
.. code-block:: cpp
void* params[] = {&d_vector, &d_block_reduced, &d_partition_reduced};
// Launching kernel from host.
HIP_CHECK(hipLaunchCooperativeKernel(vector_reduce_kernel<partition_size>,
dim3(num_blocks),
dim3(threads_per_block),
params,
0,
hipStreamDefault));\
// Check if the kernel launch was successful.
HIP_CHECK(hipGetLastError());
Conclusion
==========
With cooperative groups, you can easily use custom partitions to create custom tiles for custom solutions. You can find the complete code at `cooperative groups ROCm example. <https://github.com/ROCm/rocm-examples/tree/develop/HIP-Basic/cooperative_groups>`_
docs/tutorial/reduction.rst
+722
Melihat File
@@ -0,0 +1,722 @@
.. meta::
:description: HIP reduction tutorial
:keywords: AMD, ROCm, HIP, reduction, tutorial
*************************************************************
Reduction
*************************************************************
Reduction is a common algorithmic operation used in parallel programming to reduce an array of elements into a shorter array of elements or a single value. This document exploits reduction to introduce some key considerations while designing and optimizing GPU algorithms.
This document is a rejuvenation and extension of the invaluable `work of Mark Harris <https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf>`_. While the author approaches the topic with a less naive approach, reviewing some original material is valuable to see how much the underlying hardware has changed. This document provides a greater insight to demonstrate progress.
The algorithm
=============
Reduction has many names depending on the domain; in functional programming it's referred to as `fold <https://en.wikipedia.org/wiki/Fold_(higher-order_function)>`_, in C++, it's called ``std::accumulate`` and in C++17, as ``std::reduce``. A reduction takes a range of inputs and "reduces" the given range with a binary operation to a singular or scalar output. Canonically, a reduction requires a "zero" element that bootstraps the algorithm and serves as one of the initial operands to the binary operation. The "zero" element is generally called `identity or neutral <https://en.wikipedia.org/wiki/Identity_element>`_ element in the group theory, which implies that it is an operand that doesn't change the result. Some typical use cases are: calculating a sum or normalizing a dataset and finding the maximum value in the dataset. The latter use case is discussed further in this tutorial.
.. figure:: ../data/tutorial/reduction/foldl.svg
:alt: Diagram demonstrating fold left
There are multiple variations of reduction that allow parallel processing. The approach taken by ``std::reduce`` requires the user-provided binary operator to operate on any combination of identity and input range elements, or even exclusively on any of them. This allows you to insert any number of identities to facilitate parallel processing and then combine the partial results of parallel execution.
.. figure:: ../data/tutorial/reduction/parallel_foldl.svg
:alt: Diagram demonstrating parallel fold left
Reduction on GPUs
=================
Implementing reductions on GPUs requires a basic understanding of the :doc:`/understand/programming_model`. The document explores aspects of low-level optimization best discussed through the :ref:`inherent_thread_model`, and refrains from using cooperative groups.
Synchronizing parallel threads of execution across a GPU is crucial for correctness as the partial results can't be synchronized before they manifest. Synchronizing all the threads running on a GPU at any given time is possible, however, it is a costly and intricate operation. If synchronization is not absolutely necessary, map the parallel algorithm so that multiprocessors and blocks can make independent progress and need not sync frequently.
There are ten reduction implementations in the `rocm-examples <https://github.com/ROCm/rocm-examples/tree/develop/Tutorials/reduction/include/Reduction>`_, which are described in the following sections.
Naive shared reduction
----------------------
The naive algorithm takes a tree-like shape, where the computational domain is purposefully distributed among blocks. In all blocks, all threads participate in loading data from persistent (from the kernel's perspective) global memory into the shared memory. This helps to perform tree-like reduction for a single thread by writing the partial result to global, in a location unique to the block, which allows the block to make independent progress. The partial results are combined in subsequent launches of the same kernel until a scalar result is reached.
.. figure:: ../data/tutorial/reduction/naive_reduction.svg
:alt: Diagram demonstrating naive reduction
This approach requires temporary storage based on the number of blocks launched, as each block outputs a scalar partial result. Depending on the need to store or destroy the input, a second temporary storage might be needed, which could be large enough to store the results of the second kernel launch. Alternatively, you can reuse the storage of the larger than necessary original input. These implementations differ so slightly that the document only considers the use case where the input could be destroyed.
.. code-block:: C++
std::size_t factor = block_size; // block_size from hipGetDeviceProperties()
auto new_size = [factor](const std::size_t actual)
{
// Every pass reduces input length by 'factor'. If actual size is not divisible by factor,
// an extra output element is produced using some number of zero_elem inputs.
return actual / factor + (actual % factor == 0 ? 0 : 1);
};
For threads that don't have unique inputs, feed ``zero_elem`` instances to threads. The backing of double-buffering is allocated as such:
.. code-block:: C++
// Initialize host-side storage
std::vector<unsigned> input(input_count);
std::iota(input.begin(), input.end(), 0);
// Initialize device-side storage
unsigned *front,
*back;
hipMalloc((void**)&front, sizeof(unsigned) * input_count);
hipMalloc((void**)&back, sizeof(unsigned) * new_size(input_count));
hipMemcpy(front, input.data(), input.size() * sizeof(unsigned), hipMemcpyHostToDevice);
Data is initialized on the host and dispatched to the device followed by the commencement of device-side reduction. The swapping of the double-buffer on the last iteration is omitted, therefore the result is in the back-buffer irrespective of the input size.
.. code-block:: C++
for (uint32_t curr = input_count; curr > 1;)
{
hipLaunchKernelGGL(
kernel,
dim3(new_size(curr)),
dim3(block_size),
factor * sizeof(unsigned),
hipStreamDefault,
front,
back,
kernel_op,
zero_elem,
curr);
curr = new_size(curr);
if (curr > 1)
std::swap(front, back);
}
This structure persists in the kernel throughout all the variations of reduction with slight modifications to ``factor`` and shared memory allocation:
.. code-block:: C++
template<typename T, typename F>
__global__ void kernel(
T* front,
T* back,
F op,
T zero_elem,
uint32_t front_size)
{
extern __shared__ T shared[];
// Overindex-safe read of input
auto read_global_safe = [&](const uint32_t i)
{
return i < front_size ? front[i] : zero_elem;
};
const uint32_t tid = threadIdx.x,
bid = blockIdx.x,
gid = bid * blockDim.x + tid;
// Read input from front buffer to shared
shared[tid] = read_global_safe(gid);
__syncthreads();
// Shared reduction
for (uint32_t i = 1; i < blockDim.x; i *= 2)
{
if (tid % (2 * i) == 0)
shared[tid] = op(shared[tid], shared[tid + i]);
__syncthreads();
}
// Write result from shared to back buffer
if (tid == 0)
back[bid] = shared[0];
}
While the ``tid % (2 * i) == 0`` indexing scheme yields correct results, it also leads to high thread divergence. Thread divergence indicates the event when the threads in a warp diverge, which implies that the threads have to execute different instructions in a given clock cycle. This is easily manifested using ``if-else`` statements as shown here, but can also be manifested as ``for`` loop dependent on thread ID lengths. Even though the number of active threads participating in the reduction reduces, warps remain active longer than necessary, as at least one lane in a warp hits the ``if`` statement.
Reducing thread divergence
--------------------------
You can reduce divergence by keeping dataflow between memory addresses identical but reassigning the thread ids.
.. figure:: ../data/tutorial/reduction/reduced_divergence_reduction.svg
:alt: Diagram demonstrating reduced divergence reduction
.. code-block:: diff
:emphasize-lines: 4-7
// Shared reduction
for (uint32_t i = 1; i < blockDim.x; i *= 2)
{
- if (tid % (2 * i) == 0)
- shared[tid] = op(shared[tid], shared[tid + i]);
+ if (uint32_t j = 2 * i * tid; j < blockDim.x)
+ shared[j] = op(shared[j], shared[j + i]);
__syncthreads();
}
This way inactive threads start accumulating uniformly towards the higher thread ID index range and might uniformly skip to ``__syncthreads()``. However, this introduces a bank conflicts issue.
Resolving bank conflicts
------------------------
Both AMD and NVIDIA implement shared memory in the hardware by organizing storage into banks of various sizes. This hardware element is known as Local Data Share (LDS) on AMD hardware. On NVIDIA hardware, it's implemented using the same silicon as the L1 data cache. You can think of shared memory as a striped 2-dimensional range of memory. Shared memory bank's count, width, and depth depend on the architecture. A bank conflict occurs when different threads in a warp access the same bank during the same operation. In this case, the hardware prevents the attempted concurrent accesses to the same bank by converting them into serial accesses.
- `"AMD Instinct MI200" Instruction Set Architecture, Chapter 11.1 <https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/instinct-mi200-cdna2-instruction-set-architecture.pdf>`_
- `"RDNA 2" Instruction Set Architecture, Chapter 10.1 <https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna2-shader-instruction-set-architecture.pdf>`_
A notable exception is when the shared read uniformly broadcasts to the same address across the entire warp. A better implementation of the naive algorithm is to form continuous ranges of the threads activities and their memory accesses.
.. code-block:: diff
:emphasize-lines: 2-7
// Shared reduction
-for (uint32_t i = 1; i < blockDim.x; i *= 2)
-{
- if (tid % (2 * i) == 0)
+for (uint32_t i = blockDim.x / 2; i != 0; i /= 2)
+{
+ if (tid < i)
shared[tid] = op(shared[tid], shared[tid + i]);
__syncthreads();
}
.. figure:: ../data/tutorial/reduction/conflict_free_reduction.svg
:alt: Diagram demonstrating bank conflict free reduction
.. note::
To avoid bank conflicts, read shared memory in a coalesced manner, which implies that reads/writes of each lane in a warp evaluate to consecutive locations. Analyzing the read/write patterns could help you to understand the cause of bank conflicts. For more details, check `CDNA3 ISA <https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-mi300-cdna3-instruction-set-architecture.pdf>`_ or `RDNA3 ISA <https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/instruction-set-architectures/rdna3-shader-instruction-set-architecture-feb-2023_0.pdf>`_ data share operations chapter.
Utilize upper half of the block
-------------------------------
The preceding implementation is free of low-level GPU-specific anti-patterns. However, it still exhibits some common shortcomings. The loop performing the reduction in the shared memory starts from ``i = blockDim.x / 2`` and the first predicate ``if (tid < i)`` immediately disables half of the block, which only helps load the data into the shared memory. You can change the kernel along with the calculation of ``factor`` on the host, as shown here:
.. code-block:: diff
:emphasize-lines: 3,4
const uint32_t tid = threadIdx.x,
bid = blockIdx.x,
- gid = bid * blockDim.x + tid;
+ gid = bid * (blockDim.x * 2) + tid;
// Read input from front buffer to shared
-shared[tid] = read_global_safe(gid);
+shared[tid] = op(read_global_safe(gid), read_global_safe(gid + blockDim.x));
__syncthreads();
By eliminating half of the threads and giving meaningful work to all the threads by unconditionally performing a binary ``op``, you can prevent the wastage of half of the threads.
Even though global memory is read in a coalesced fashion, as preferred by the memory controller, optimal performance is still limited by the instruction throughput.
Omit superfluous synchronization
--------------------------------
Warps are known to execute in a strict lockstep fashion. Therefore, once shared reduction reaches a point where only a single warp participates meaningfully, you can cut short the loop and let the rest of the warps terminate. Moreover, you can also unroll the loop without syncing the entire block.
The ``tmp`` namespace used beyond this point in this document holds a handful of template meta-programmed utilities to facilitate writing flexible and optimal code.
:code:`tmp::static_for` is not just a constant folding within the optimizer but a variation of the language :code:`for` loop, where the running index is a compile-time constant and is eligible for use in compile-time evaluated contexts.
Consider the following code:
.. code-block:: C++
constexpr int size = 4;
for (int i = 0 ; i < size ; ++i)
{
printf("%d", i);
}
This compiles to the following binaries:
**LLVM Block**
.. code-block::
main:
push rbx
lea rbx, [rip + .L.str]
mov rdi, rbx
xor esi, esi
xor eax, eax
call printf@PLT
mov rdi, rbx
mov esi, 1
xor eax, eax
call printf@PLT
mov rdi, rbx
mov esi, 2
xor eax, eax
call printf@PLT
mov rdi, rbx
mov esi, 3
xor eax, eax
call printf@PLT
xor eax, eax
pop rbx
ret
.L.str:
.asciz "%d"
**GCC**
.. code-block:: asm
.LC0:
.string "%d"
main:
push rbx
xor ebx, ebx
.L2:
mov esi, ebx
mov edi, OFFSET FLAT:.LC0
xor eax, eax
add ebx, 1
call printf
cmp ebx, 4
jne .L2
xor eax, eax
pop rbx
ret
**MSVC**
.. code-block::
main PROC
$LN12:
push rbx
sub rsp, 32
xor ebx, ebx
npad 8
$LL4@main:
mov edx, ebx
lea rcx, OFFSET FLAT:'string'
call printf
inc ebx
cmp ebx, 4
jl SHORT $LL4@main
xor eax, eax
add rsp, 32
pop rbx
ret 0
main ENDP
LLVM unrolls the loop and compiles to a flat series of ``printf`` invocations, while both GCC and MSVC keep the loop intact, as visible from the compare (``cmp``) and the jump (``jne``, ``jl``) instructions. LLVM code generation is identical to manually writing the unrolled loop:
.. code-block:: C++
printf("%d", 0);
printf("%d", 1);
printf("%d", 2);
printf("%d", 3);
While various non-standard pragmas are available to hint or force the compiler to unroll the loop, we instead use template meta-programming to force feed the compiler the unrolled loop.
.. code-block:: C++
constexpr int size = 4;
// Maybe unrolled loop
for (int i = 0 ; i < size ; ++i)
{
printf("%d", i);
}
// Force unrolled loop
using namespace tmp;
static_for<0, less_than<size>, increment<1>>([]<int i>()
{
printf("%d", i);
});
The most notable structural difference is that in the language ``for`` loop, the loop variable is given a name in the beginning, while in the ``static_for`` utility, the loop variable is given a name in the end. An important bonus is that in the loop's body, you can use the running index ``i`` in contexts requiring constant expressions such as template arguments or inside ``if constexpr``.
:code:`tmp::static_switch` takes runtime value and runtime dispatches to a range of set of tabulated functions, where said value is a compile-time constant and is eligible for use in compile-time evaluated contexts.
Consider the following code:
.. code-block:: C++
int warp_size = device_props.warpSize;
switch (warp_size)
{
case 32:
hipLaunchKernelGGL(kernel<32>, ...);
break;
case 64:
hipLaunchKernelGGL(kernel<64>, ...);
break;
}
In the preceding code, note the code repetition for all possible values of ``warp_size``, the code is prepared to handle. To avoid this, use ``tmp::static_switch``, as shown:
.. code-block:: C++
tmp::static_switch<std::array{32, 64}>(warp_size, [&]<int WarpSize>()
{
hipLaunchKernelGGL(kernel<WarpSize>, ...);
});
.. code-block:: diff
:emphasize-lines: 1,2,9,10,16-24
-template<typename T, typename F>
+template<uint32_t WarpSize, typename T, typename F>
__global__ void kernel(
...
)
{
...
// Shared reduction
-for (uint32_t i = blockDim.x / 2; i != 0; i /= 2)
+for (uint32_t i = blockDim.x / 2; i > WarpSize; i /= 2)
{
if (tid < i)
shared[tid] = op(shared[tid], shared[tid + i]);
__syncthreads();
}
+// Warp reduction
+tmp::static_for<WarpSize, tmp::not_equal<0>, tmp::divide<2>>([&]<int I>()
+{
+ if (tid < I)
+ shared[tid] = op(shared[tid], shared[tid + I]);
+#ifdef __HIP_PLATFORM_NVIDIA__
+ __syncwarp(0xffffffff >> (WarpSize - I));
+#endif
+});
Because HIP typically targets hardware with warp sizes of 32 (NVIDIA GPUs and RDNA AMD GPUs) and 64 (CDNA AMD GPUs), portable HIP code must handle both. That is why instead of assuming a warp size of 32, make the warp size a template argument of the kernel. This allows you to unroll the final loop using ``tmp::static_for`` in a parametric way but still having the code read much like an ordinary loop.
Promoting the warp size to being a compile-time constant also requires you to handle it similarly on the host-side. You can sandwich the kernel launch with ``tmp::static_switch``, promoting the snake-case run-time ``warp_size`` variable to a camel-case compile-time constant ``WarpSize``.
.. code-block:: diff
:emphasize-lines: 4,5,7,8,18
// Device-side reduction
for (uint32_t curr = input_count; curr > 1;)
{
+ tmp::static_range_switch<std::array{32, 64}>(warp_size, [&]<int WarpSize>() noexcept
+ {
hipLaunchKernelGGL(
- kernel,
+ kernel<WarpSize>,
dim3(new_size(curr)),
dim3(block_size),
factor * sizeof(unsigned),
hipStreamDefault,
front,
back,
kernel_op,
zero_elem,
curr);
+ });
...
}
.. note::
Neither RDNA- nor CDNA-based AMD hardware provides guaranteed independent progress to lanes of the same warp. When targeting NVIDIA hardware, lanes of a warp might execute somewhat independently as long as the programmer assists the compiler using dedicated built-in functions. This feature is called Independent Thread Scheduling. The HIP headers don't expose the necessary warp primitives and their overloads.
Portable applications can still tap into this feature with carefully ``#ifdef`` -ed code, but at this particular optimization level, it's a requirement. The code implicitly relies on the lockstep behavior of an ROCm wavefront, but CUDA warps don't share this property. You must synchronize all the active lanes of a warp to avoid a data race with some lanes progressing faster than others in the same warp.
Unroll all loops
----------------
While the previous step primarily aims to remove unnecessary syncing, it also unrolls the end of the loop. However, you could also force unrolling the first part of the loop. This saves a few scalar registers (values the compiler can prove to be uniform across warps).
.. code-block:: diff
:emphasize-lines: 1-4,11,12,17,18,20-23,26
-template<uint32_t WarpSize, typename T, typename F>
-__global__ void kernel(
+template<uint32_t BlockSize, uint32_t WarpSize, typename T, typename F>
+__global__ __launch_bounds__(BlockSize) void kernel(
T* front,
T* back,
F op,
T zero_elem,
uint32_t front_size)
{
- extern __shared__ T shared[];
+ __shared__ T shared[BlockSize];
...
// Shared reduction
- for (uint32_t i = blockDim.x / 2; i > WarpSize; i /= 2)
+ tmp::static_for<BlockSize / 2, tmp::greater_than<WarpSize>, tmp::divide<2>>([&]<int I>()
{
- if (tid < i)
- shared[tid] = op(shared[tid], shared[tid + i]);
+ if (tid < I)
+ shared[tid] = op(shared[tid], shared[tid + I]);
__syncthreads();
}
+ );
Introducing yet another template argument for the kernel and moving from ``for`` to ``tmp::static_for`` leads to the following two notable improvements:
- Introducing new attribute ``__launch_bounds__(BlockSize)`` to the kernel instructs the compiler that the kernel will only be launched using the designated block size. This implies that the launches of differing block sizes will fail. This allows the optimizer to enroll the ``blockDim.x`` variable in constant folding as well as get information about register usage.
- Turning the block size into a compile-time constant allows you to statically allocate the shared memory.
Communicate using warp-collective functions
-------------------------------------------
Shared memory provides a fast communication path within a block, however when performing reduction within the last warp, you can use faster means of communication, which is warp-collective or cross-lane functions. Instead of using the hardware-backed shared memory, you can directly copy between the local memory (registers) of each lane in a warp. This can be achieve using the shuffle functions.
See how to use ``__shfl_down()``, which is one of the most restrictive but also the most structured communication schemes.
.. code-block:: C++
// Warp reduction
if (tid < WarpSize)
{
T res = op(shared[tid], shared[tid + WarpSize]);
tmp::static_for<WarpSize / 2, tmp::not_equal<0>, tmp::divide<2>>([&]<int Delta>()
{
res = op(res, __shfl_down(res, Delta));
});
// Write result from shared to back buffer
if (tid == 0)
back[bid] = res;
}
Using warp-collective functions for communication requires the control flow to be uniform across warps, as the name warp-collective implies. Therefore, you can see that the thread ID is being checked outside the loop, but the result is written inside due to variable scoping.
Prefer warp communication over shared
-------------------------------------
As mentioned in the previous step, communication between local memory is faster than shared memory. Instead of relying on the local memory only at the end of the tree-like reduction, a better approach is to turn the tree reduction inside out and perform multiple warp reductions in parallel on all active threads, thus communicating only their partial results through the shared memory.
.. figure:: ../data/tutorial/reduction/warp_reduction.svg
:alt: Diagram demonstrating warp reduction
.. figure:: ../data/tutorial/reduction/warp_reduction_with_shared.svg
:alt: Diagram demonstrating warp reduction and results store in shared memory
The kernel versions differ significantly enough to be described using a diff; use afresh instead.
.. code-block:: C++
template<uint32_t BlockSize, uint32_t WarpSize, typename T, typename F>
__global__ __launch_bounds__(BlockSize) void kernel(
T* front,
T* back,
F op,
T zero_elem,
uint32_t front_size)
{
// ...
}
The kernel signature and the reduction factor are the same as in previous cases; only the implementation differs.
.. code-block:: C++
static constexpr uint32_t WarpCount = BlockSize / WarpSize;
__shared__ T shared[WarpCount];
auto read_global_safe =
[&](const uint32_t i) { return i < front_size ? front[i] : zero_elem; };
auto read_shared_safe =
[&](const uint32_t i) { return i < WarpCount ? shared[i] : zero_elem; };
const uint32_t tid = threadIdx.x,
bid = blockIdx.x,
gid = bid * (blockDim.x * 2) + tid,
wid = tid / WarpSize,
lid = tid % WarpSize;
// Read input from front buffer to local
T res = op(read_global_safe(gid), read_global_safe(gid + blockDim.x));
As we communicate the results of warps through shared memory, the same number of elements are required in the shared memory as warps within the block. Similar to how you can only launch kernels at block granularity, you can only warp reduce with ``WarpSize`` granularity due to the collective nature of the cross-lane builtins. To address this, you can use ``read_shared_safe`` to pad overindexing by reading ``zero_elem``. Reading from global remains unaffected.
.. code-block:: C++
// Perform warp reductions and communicate results via shared
// for (uint32_t ActiveWarps = WarpCount;
// ActiveWarps != 0;
// ActiveWarps = ActiveWarps != 1 ?
// divide_ceil(ActiveWarps, WarpSize) :
// ActiveWarps = 0)
tmp::static_for<
WarpCount,
tmp::not_equal<0>,
tmp::select<
tmp::not_equal<1>,
tmp::divide_ceil<WarpSize>,
tmp::constant<0>>>([&]<uint32_t ActiveWarps>()
{
if(wid < ActiveWarps)
{
// Warp reduction
tmp::static_for<WarpSize / 2, tmp::not_equal<0>, tmp::divide<2>>([&]<int Delta>()
{
res = op(res, __shfl_down(res, Delta));
});
// Write warp result from local to shared
if(lid == 0)
shared[wid] = res;
}
__syncthreads();
// Read warp result from shared to local
res = read_shared_safe(tid);
});
// Write result from local to back buffer
if(tid == 0)
back[bid] = res;
``ActiveWarps`` iterates from ``WarpCount`` until it reaches ``0``. Every iteration of ``ActiveWarps`` reduces the ``WarpSize``. In cases where the partial result count isn't a divisor of ``ActiveWarps`` and you need to launch an extra warp, use ``tmp::divide_ceil``, which always rounds to positive infinity. The tertiary ``tmp::select`` is required because such division never reaches ``0``, so you must terminate the loop after the last warp concludes.
In each iteration, if the warp is active, which means it has at least a single valid input, it carries out a pass of warp reduction and writes output based on warp ID. Reading is carried out based on thread ID. Global output continues to be based on block ID.
Amortize bookkeeping variable overhead
--------------------------------------
The previous sections explained how to reduce register usage to improve occupancy. This allows more blocks to execute in parallel on all multiprocessors, leading to more global store/load latency to be hidden. Reducing the number of kernels in flight while still carrying out the same workload reduces the wastage of registers while loading and maintaining bookkeeping variables such as kernel indices.
An example of this optimization is performing one binary ``op`` while loading input from global. Even though the operation is said to be carried out "in flight", the two values are loaded into local memory (registers) before ``op`` is called.
A more general form of this optimization is wrapping most kernel logic in loops that carry out the workload of multiple kernel instances but require storing only a single instance of most of the bookkeeping logic. In code, this multiplicity factor is referred to via the ``ItemsPerThread`` compile-time constant, which is supplied by a template argument to allow for loop unrolling.
This kernel variant utilizes another generally applicable utility known as ``hip::static_array``, which is a more restrictive wrapper over the builtin array than ``std::array``, as it allows indexing only compile-time constants using the usual tuple-like ``template <size_t I> auto get<I>(...)`` interface.
.. note::
On a GPU, there is no stack, and the local memory is provisioned from the register file. This provisioning takes place statically. To paraphrase, the address range of a thread's local memory is determined at compile-time. When an array is defined and used in the local storage, the compiler can only maintain its storage in the register file as long as all accesses to the array are computable by the compiler at compile-time. It doesn't need to be a compile-time constant as long as the compiler can resolve the addresses of the accesses through constant folding or some other means. If the compiler fails to do so, the array will be backed by global memory, which is indicated by allocating a non-zero number of spill registers observable using static analysis tools. However, this is slower by the magnitude of multiple order. ``hip::static_array`` via its ``hip::get<>`` interface ensures that no such spills occur.
.. code-block:: C++
template<uint32_t BlockSize, uint32_t WarpSize, uint32_t ItemsPerThread>
__global__ static __launch_bounds__(BlockSize) void kernel(...)
The kernel now has three compile-time configurable parameters. The only part of the kernel that changes depends on how you load data from global and perform the binary operation on those loaded values. So, the following step to read input from front buffer to global is now split into two steps: :ref:`reading-items` and :ref:`processing-items` .
.. code-block:: C++
// Read input from front buffer to local
T res = op(read_global_safe(gid), read_global_safe(gid + blockDim.x));
.. _reading-items:
Reading ``ItemsPerThread``
^^^^^^^^^^^^^^^^^^^^^^^^^^
The change to reading happens inside `read_global_safe`:
.. code-block:: C++
auto read_global_safe = [&](const int32_t i) -> hip::static_array<T, ItemsPerThread>
{
return [&]<int32_t... I>(std::integer_sequence<int32_t, I...>)
{
if(i + ItemsPerThread < front_size)
return hip::static_array<T, ItemsPerThread>{
front[i + I]...
};
else
return hip::static_array<T, ItemsPerThread>{
(i + I < front_size ? front[i + I] : zero_elem)...
};
}(std::make_integer_sequence<int32_t, ItemsPerThread>());
};
Note that each array element is being loaded consecutively without the flexibility of a configurable ``ItemsPerThread`` property. This is morally equivalent to:
.. code-block:: C++
T arr[4] = {
front[gid + 0],
front[gid + 1],
front[gid + 2],
front[gid + 3]
}
This is exactly what's happening in the ``front[i + I]...`` fold-expression. However, this can only be issued if the entire read operates on real input without padding using ``zero_elem``. If some reads over-index the input, the read turns into:
.. code-block:: C++
T arr[4] = {
i + 0 < front_size ? front[i + 0] : zero_elem,
i + 1 < front_size ? front[i + 1] : zero_elem,
i + 2 < front_size ? front[i + 2] : zero_elem,
i + 3 < front_size ? front[i + 3] : zero_elem
}
This makes it easier for the compiler to recognize vector loads from global. As the performance at large is dominated by how you move the data, it's only natural to utilize dedicated instructions to move more data with less binary. This is evident by the huge performance improvement when loading two values per thread. For more information, see `the compiler explorer <https://godbolt.org/z/b36Eea69q>`_ to learn how loading for AMD (both RDNA and CDNA) compiles to ``global_load_dwordx4``, where ``x4`` denotes the 4-vector variant of the instruction.
.. note::
Note that ``read_global_safe``, which used to take an ``uint32_t`` as the index type, now takes a signed integer. When indexing an array with unsigned integers, the compiler has to handle integer overflows, as the C/C++ standards defined them. It might happen that some part of the vector load indices overflow, thus resulting in a non-contiguous read. If you change the previously linked code to use an unsigned integer as the thread ID, the compiler won't emit a vector load. Signed integer overflow is an undefined behavior, and hence, unknown to the optimizer. To convey the absence of overflow to the compiler with unsigned indices, add ``__builtin_assume(gid + 4 > gid)``, or the more portable ``[[assume]](gid + 4 > gid)``, once ``amdclang++`` supports it.
``read_global_safe`` implementation is an Immediately Invoked Lambda Expression (IILE), because ``ItemsPerThread`` is an integer value, while you need a compile-time ``iota``-like sequence of integers as a pack for the fold-expressions to expand on. This can only occur as part of template argument deduction on the IILE.
.. _processing-items:
Processing ``ItemsPerThread``
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
Once the kernel reads ``ItemsPerThread`` number of inputs to local, it immediately reduces them to a scalar. There is no reason to propagate the input element multiplicity to the warp reduction phase.
.. code-block:: C++
T res = [&]()
{
// Read input from front buffer to local
hip::static_array<T, ItemsPerThread> arr = read_global_safe(gid);
// Reduce ItemsPerThread to scalar
tmp::static_for<1, tmp::less_than<ItemsPerThread>, tmp::increment<1>>([&]<int I>()
{
get<0>(arr) = op(get<0>(arr), get<I>(arr));
});
return get<0>(arr);
}();
Two-pass reduction
------------------
Alter kernel launch and input fetching such that no more blocks are launched than what a subsequent kernel launch's single block can conveniently reduce, while performing multiple passes of input reading from global and combining their results before engaging in the end game tree-like reduction.
With this method, you can save at least one to two kernel launches for large inputs.
Global data share
-----------------
.. warning::
This modification can only be executed on AMD hardware.
Perform the first step of the two-pass reduction, but in the end, instead of writing to global and reading it back in a subsequent kernel, write the partial results to the Global Data Share (GDS). This is an ``N+1`` th shared memory that is accessed by all multiprocessors and is also on-chip memory.
.. note::
The API doesn't guarantee the order in which blocks are scheduled even though all GPUs schedule them in the same monotonically increasing order of block ids. Relying on this implicitly, the last block of a grid is in the optimal position to observe the side effects of all other blocks (using spinlocks or other methods) without occupying a multiprocessor for longer than necessary.
Without launching a second kernel, you can make the last block collect the results of all other blocks from GDS by implicitly exploiting the scheduling behavior or relying on another AMD-specific feature called Global Wave Sync (GWS) to merge them for a final tree-like reduction.
.. note::
GDS and GWS are reserved runtime features that the HIP API doesn’t cover. Invoking these functionalities requires inline AMDGCN assembly. Moreover, the fact that the runtime doesn’t virtualize the GDS, imposes further restrictions on concurrent scheduling of other kernels.
Conclusion
==========
Optimizing code on GPUs, like on any other architecture, requires careful consideration and balancing of resources and costs of various operations to obtain optimal performance. This document explored optimizing reductions much beyond the territory of diminishing returns. This approach introduced multiple optimization techniques and discussed opportunities.
The document focused on reductions when an entire device participates in it. Still, the choice of optimal compile-time constants or even the algorithm itself might not be optimal when its multiple blocks participate in multiple parallel reductions or when each thread performs its reduction. However, when multiple devices participate in the same reduction, other aspects must be considered.
Most solutions, including the ones covered in this document, are given to the end users in a turnkey fashion via algorithm primitive libraries. These solutions might not be the fastest in all cases, but they are close to being the gold standard for carrying out certain operations as reasonably as possible.
docs/tutorial/saxpy.rst
+751
Melihat File
@@ -0,0 +1,751 @@
.. meta::
:description: The SAXPY tutorial on HIP
:keywords: AMD, ROCm, HIP, SAXPY, tutorial
*******************************************************************************
SAXPY - Hello, HIP
*******************************************************************************
This tutorial explains the basic concepts of the single-source
Heterogeneous-computing Interface for Portability (HIP) programming model and
the essential tooling around it. It also reviews some commonalities of
heterogenous APIs in general. This topic assumes basic familiarity with the
C/C++ compilation model and language.
Prerequisites
=============
To follow this tutorial, you'll need installed drivers and a HIP compiler
toolchain to compile your code. Because HIP for ROCm supports compiling and
running on Linux and Windows with AMD and NVIDIA GPUs, the combination of
install instructions is more than worth covering as part of this tutorial. For
more information about installing HIP development packages, see
:doc:`/install/install`.
.. _hip-tutorial-saxpy-heterogeneous-programming:
Heterogeneous programming
=========================
*Heterogeneous programming* and *offloading APIs* are often mentioned together. Heterogeneous programming deals with devices of varying capabilities simultaneously. Offloading focuses on the "remote" and asynchronous aspects of computation. HIP encompasses both. It exposes GPGPU (general-purpose GPU) programming much like ordinary host-side CPU programming and lets you move data across various devices.
When programming in HIP (and other heterogenous APIs for that matter), remember that target devices are built for a specific purpose. They are designed with different tradeoffs than traditional CPUs and therefore have very different performance characteristics. Even subtle changes in code might adversely affect execution time.
Your first lines of HIP code
============================
First, let's do the "Hello, World!" of GPGPU: SAXPY. Single-precision A times X Plus Y (*SAXPY*) is a mathematical acronym; a vector equation :math:`a\cdot x+y=z` where :math:`a\in\mathbb{R}` is a scalar and :math:`x,y,z\in\mathbb{V}` are vector quantities of some large dimensionality. This vector space is defined over the set of reals. Practically speaking, you can compute this using a single ``for`` loop over three arrays.
.. code-block:: C++
for (int i = 0 ; i < N ; ++i)
z[i] = a * x[i] + y[i];
In linear algebra libraries, such as BLAS (Basic Linear Algebra Subsystem) this operation is defined as AXPY "A times X Plus Y". The "S" comes from *single-precision*, meaning that array element is ``float`` -s (IEEE 754 binary32 representation).
To quickly get started, use the set of `HIP samples from GitHub <https://github.com/amd/rocm-examples/>`_. With Git configured on your machine, open a command-line and navigate to your desired working directory, then run:
.. code-block:: shell
git clone https://github.com/amd/rocm-examples.git
A simple implementation of SAXPY resides in the ``HIP-Basic/saxpy/main.hip`` file in this repository. The HIP code here mostly deals with where data has to be and when, and how devices transform this data. The first HIP calls deal with allocating device-side memory and copying data from host-side memory to device side in a C runtime-like fashion.
.. code-block:: C++
// Allocate and copy vectors to device memory.
float* d_x{};
float* d_y{};
HIP_CHECK(hipMalloc(&d_x, size_bytes));
HIP_CHECK(hipMalloc(&d_y, size_bytes));
HIP_CHECK(hipMemcpy(d_x, x.data(), size_bytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(d_y, y.data(), size_bytes, hipMemcpyHostToDevice));
``HIP_CHECK`` is a custom macro borrowed from the examples utilities which checks the error code returned by API functions for errors and reports them to the console. It is not essential to the API, but it is a good practice to check the error codes of the HIP APIs in case you pass on incorrect values to the API, or the API might be out of resources.
The code selects the device to allocate to and to copy to. Commands are issued to the HIP runtime per thread, and every thread has a device set as the target of commands. The default device is ``0``, which is equivalent to calling ``hipSetDevice(0)``.
Launch the calculation on the device after the input data has been prepared.
.. code-block:: C++
__global__ void saxpy_kernel(const float a, const float* d_x, float* d_y, const unsigned int size)
{
// ...
}
int main()
{
// ...
// Launch the kernel on the default stream.
saxpy_kernel<<<dim3(grid_size), dim3(block_size), 0, hipStreamDefault>>>(a, d_x, d_y, size);
}
Analyze at the signature of the offloaded function:
- ``__global__`` instructs the compiler to generate code for this function as an
entrypoint to a device program, such that it can be launched from the host.
- The function does not return anything, because there is no trivial way to
construct a return channel of a parallel invocation. Device-side entrypoints
may not return a value, their results should be communicated using output
parameters.
- Device-side functions are typically called compute kernels, or just kernels
for short. This is to distinguish them from non-graphics-related graphics
shaders, or just shaders for short.
- Arguments are taken by value and all arguments shall be
`TriviallyCopyable <https://en.cppreference.com/w/cpp/named_req/TriviallyCopyable>`_,
meaning they should be `memcpy`-friendly. (Imagine if they had custom copy
constructors. Where would that logic execute? On the host? On the device?)
Pointer arguments are pointers to device memory, one typically backed by
VRAM.
- We said that we'll be computing :math:`a\cdot x+y=z`, however we only pass
two pointers to the function. We'll be canonically reusing one of the inputs
as outputs.
This function is launched from the host using a language extension often called
the triple chevron syntax. Inside the angle brackets, provide the following.
- The number of :ref:`blocks <inherent_thread_hierarchy_block>` to launch (our :ref:`grid <inherent_thread_hierarchy_grid>` size)
- The number of threads in a :ref:`block <inherent_thread_hierarchy_block>` (our :ref:`block <inherent_thread_hierarchy_block>` size)
- The amount of shared memory to allocate by the host
- The device stream to enqueue the operation on
The :ref:`block <inherent_thread_hierarchy_block>` size and shared memory become important later in :doc:`reduction`. For
now, a hardcoded ``256`` is a safe default for simple kernels such as this.
Following the triple chevron is ordinary function argument passing.
Look at how the kernel is implemented.
.. code-block:: C++
__global__ void saxpy_kernel(const float a, const float* d_x, float* d_y, const unsigned int size)
{
// Compute the current thread's index in the grid.
const unsigned int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
// The grid can be larger than the number of items in the vectors. Avoid out-of-bounds addressing.
if(global_idx < size)
{
d_y[global_idx] = a * d_x[global_idx] + d_y[global_idx];
}
}
- The unique linear index identifying the thread is computed from the :ref:`block <inherent_thread_hierarchy_block>` ID
the thread is a member of, the :ref:`block <inherent_thread_hierarchy_block>`'s size and the ID of the thread within
the :ref:`block <inherent_thread_hierarchy_block>`.
- A check is made to avoid overindexing the input.
- The useful part of the computation is carried out.
Retrieval of the result from the device is done much like input data copy. In this current step the results copied from device to host. The opposite direction of the input data copy:
.. code-block:: C++
HIP_CHECK(hipMemcpy(y.data(), d_y, size_bytes, hipMemcpyDeviceToHost));
Compiling on the command line
=============================
.. _setting_up_the_command-line:
Setting up the command line
---------------------------
Strictly speaking there's no such thing as "setting up the command-line
for compilation" on Linux. To make invocations more terse, Linux and Windows
example follow.
.. tab-set::
.. tab-item:: Linux and AMD
:sync: linux-amd
While distro maintainers might package ROCm so that it installs to
system-default locations, AMD's packages aren't installed that way. They need
to be added to the PATH by the user.
.. code-block:: bash
export PATH=/opt/rocm/bin:${PATH}
You should be able to call the compiler on the command line now:
.. code-block:: bash
amdclang++ --version
.. note::
Docker images distributed by AMD, such as
`rocm-terminal <https://hub.docker.com/r/rocm/rocm-terminal/>`_ already
have `/opt/rocm/bin` on the Path for convenience. This subtly affects
CMake package detection logic of ROCm libraries.
.. tab-item:: Linux and NVIDIA
:sync: linux-nvidia
Both distro maintainers and NVIDIA package CUDA so that ``nvcc`` and related
tools are available on the command line by default. You can call the
compiler on the command line with:
.. code-block:: bash
nvcc --version
.. tab-item:: Windows and AMD
:sync: windows-amd
Windows compilers and command line tooling have traditionally relied on
extra environmental variables and PATH entries to function correctly.
Visual Studio refers to command lines with this setup as "Developer
Command Prompt" or "Developer PowerShell" for ``cmd.exe`` and PowerShell
respectively.
The HIP SDK on Windows doesn't include a complete toolchain. You will also
need:
- The Microsoft Windows SDK. It provides the import libs to crucial system
libraries that all executables must link to and some auxiliary compiler
tooling.
- A Standard Template Library (STL). Installed as part of the Microsoft
Visual C++ compiler (MSVC) or with Visual Studio.
If you don't have a version of Visual Studio 2022 installed, for a
minimal command line experience, install the
`Build Tools for Visual Studio 2022 <https://aka.ms/vs/17/release/vs_BuildTools.exe>`_
with the Desktop Developemnt Workload. Under Individual Components select:
- A version of the Windows SDK
- "MSVC v143 - VS 2022 C++ x64/x86 build tools (Latest)"
- "C++ CMake tools for Windows" (optional)
.. note::
The "C++ CMake tools for Windows" individual component is a convenience which
puts both ``cmake.exe`` and ``ninja.exe`` onto the PATH inside developer
command prompts. You can install these manually, but then you must manage
them manually.
Visual Studio 2017 and later are detectable as COM object instances via WMI.
To setup a command line from any shell for the latest Visual Studio's
default Visual C++ toolset issue:
.. code-block:: powershell
$InstallationPath = Get-CimInstance MSFT_VSInstance | Sort-Object -Property Version -Descending | Select-Object -First 1 -ExpandProperty InstallLocation
Import-Module $InstallationPath\Common7\Tools\Microsoft.VisualStudio.DevShell.dll
Enter-VsDevShell -InstallPath $InstallationPath -SkipAutomaticLocation -Arch amd64 -HostArch amd64 -DevCmdArguments '-no_logo'
$env:PATH = "${env:HIP_PATH}bin;${env:PATH}"
You should be able to call the compiler on the command line now:
.. code-block:: powershell
clang++ --version
.. tab-item:: Windows and NVIDIA
:sync: windows-nvidia
Windows compilers and command line tooling have traditionally relied on
extra environmental variables and PATH entries to function correctly.
Visual Studio refers to command lines with this setup as "Developer
Command Prompt" or "Developer PowerShell" for ``cmd.exe`` and PowerShell
respectively.
The HIP and CUDA SDKs on Windows don't include complete toolchains. You will
also need:
- The Microsoft Windows SDK. It provides the import libs to crucial system
libraries that all executables must link to and some auxiliary compiler
tooling.
- A Standard Template Library (STL). Installed as part of the Microsoft
Visual C++ compiler (MSVC) or with Visual Studio.
If you don't have a version of Visual Studio 2022 installed, for a
minimal command line experience, install the
`Build Tools for Visual Studio 2022 <https://aka.ms/vs/17/release/vs_BuildTools.exe>`_
with the Desktop Developemnt Workload. Under Individual Components select:
- A version of the Windows SDK
- "MSVC v143 - VS 2022 C++ x64/x86 build tools (Latest)"
- "C++ CMake tools for Windows" (optional)
.. note::
The "C++ CMake tools for Windows" individual component is a convenience which
puts both ``cmake.exe`` and ``ninja.exe`` onto the PATH inside developer
command prompts. You can install these manually, but then you must manage
them manually.
Visual Studio 2017 and later are detectable as COM object instances via WMI.
To setup a command line from any shell for the latest Visual Studio's
default Visual C++ toolset issue:
.. code-block:: powershell
$InstallationPath = Get-CimInstance MSFT_VSInstance | Sort-Object -Property Version -Descending | Select-Object -First 1 -ExpandProperty InstallLocation
Import-Module $InstallationPath\Common7\Tools\Microsoft.VisualStudio.DevShell.dll
Enter-VsDevShell -InstallPath $InstallationPath -SkipAutomaticLocation -Arch amd64 -HostArch amd64 -DevCmdArguments '-no_logo'
You should be able to call the compiler on the command line now:
.. code-block:: powershell
nvcc --version
Invoking the compiler manually
------------------------------
To compile and link a single-file application, use the following commands:
.. tab-set::
.. tab-item:: Linux and AMD
:sync: linux-amd
.. code-block:: bash
amdclang++ ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -lamdhip64 -L /opt/rocm/lib -O2
.. tab-item:: Linux and NVIDIA
:sync: linux-nvidia
.. code-block:: bash
nvcc ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -I /opt/rocm/include -O2 -x cu
.. tab-item:: Windows and AMD
:sync: windows-amd
.. code-block:: powershell
clang++ .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I .\Common -lamdhip64 -L ${env:HIP_PATH}lib -O2
.. tab-item:: Windows and NVIDIA
:sync: windows-nvidia
.. code-block:: powershell
nvcc .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I ${env:HIP_PATH}include -I .\Common -O2 -x cu
Depending on your computer, the resulting binary might or might not run. If not,
it typically complains about "Invalid device function". That error
(corresponding to the ``hipErrorInvalidDeviceFunction`` entry of ``hipError_t``)
means that the runtime could not find a device program binary of the
appropriate flavor embedded into the executable.
So far, the discussion has covered how data makes it from the host to the
device and back. It has also discussed the device code as source, with the HIP
runtime arguing that the correct binary to dispatch for execution. How can you
find out what device binary flavors are embedded into the executable?
.. tab-set::
.. tab-item:: Linux and AMD
:sync: linux-amd
The utilities included with ROCm help significantly to inspect binary
artifacts on disk. Add the ROCmCC installation folder to your PATH if you
want to use these utilities (the utilities expect them to be on the PATH).
You can list embedded program binaries using ``roc-obj-ls``.
.. code-block:: bash
roc-obj-ls ./saxpy
It should return something like:
.. code-block:: shell
1 host-x86_64-unknown-linux file://./saxpy#offset=12288&size=0
1 hipv4-amdgcn-amd-amdhsa--gfx803 file://./saxpy#offset=12288&size=9760
The compiler embeds a version 4 code object (more on `code
object versions <https://www.llvm.org/docs/AMDGPUUsage.html#code-object-metadata>`_)
and used the LLVM target triple `amdgcn-amd-amdhsa--gfx803` (more on `target triples
<https://www.llvm.org/docs/AMDGPUUsage.html#target-triples>`_). You can
extract that program object in a disassembled fashion for human consumption
via ``roc-obj``.
.. code-block:: bash
roc-obj -t gfx803 -d ./saxpy
This creates two files on disk and ``.s`` extension is of most interest.
Opening this file or dumping it to the console using ``cat``
lets find the disassembled binary of the SAXPY compute kernel, something
similar to:
.. code-block::
Disassembly of section .text:
<_Z12saxpy_kernelfPKfPfj>:
s_load_dword s0, s[4:5], 0x2c // 000000001000: C0020002 0000002C
s_load_dword s1, s[4:5], 0x18 // 000000001008: C0020042 00000018
s_waitcnt lgkmcnt(0) // 000000001010: BF8C007F
s_and_b32 s0, s0, 0xffff // 000000001014: 8600FF00 0000FFFF
s_mul_i32 s6, s6, s0 // 00000000101C: 92060006
v_add_u32_e32 v0, vcc, s6, v0 // 000000001020: 32000006
v_cmp_gt_u32_e32 vcc, s1, v0 // 000000001024: 7D980001
s_and_saveexec_b64 s[0:1], vcc // 000000001028: BE80206A
s_cbranch_execz 22 // 00000000102C: BF880016 <_Z12saxpy_kernelfPKfPfj+0x88>
s_load_dwordx4 s[0:3], s[4:5], 0x8 // 000000001030: C00A0002 00000008
v_mov_b32_e32 v1, 0 // 000000001038: 7E020280
v_lshlrev_b64 v[0:1], 2, v[0:1] // 00000000103C: D28F0000 00020082
s_waitcnt lgkmcnt(0) // 000000001044: BF8C007F
v_mov_b32_e32 v3, s1 // 000000001048: 7E060201
v_add_u32_e32 v2, vcc, s0, v0 // 00000000104C: 32040000
v_addc_u32_e32 v3, vcc, v3, v1, vcc // 000000001050: 38060303
flat_load_dword v2, v[2:3] // 000000001054: DC500000 02000002
v_mov_b32_e32 v3, s3 // 00000000105C: 7E060203
v_add_u32_e32 v0, vcc, s2, v0 // 000000001060: 32000002
v_addc_u32_e32 v1, vcc, v3, v1, vcc // 000000001064: 38020303
flat_load_dword v3, v[0:1] // 000000001068: DC500000 03000000
s_load_dword s0, s[4:5], 0x0 // 000000001070: C0020002 00000000
s_waitcnt vmcnt(0) lgkmcnt(0) // 000000001078: BF8C0070
v_mac_f32_e32 v3, s0, v2 // 00000000107C: 2C060400
flat_store_dword v[0:1], v3 // 000000001080: DC700000 00000300
s_endpgm // 000000001088: BF810000
Alternatively, call the compiler with ``--save-temps`` to dump all device
binary to disk in separate files.
.. code-block:: bash
amdclang++ ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -lamdhip64 -L /opt/rocm/lib -O2 --save-temps
List all the temporaries created while compiling ``main.hip`` with:
.. code-block:: bash
ls main-hip-amdgcn-amd-amdhsa-*
main-hip-amdgcn-amd-amdhsa-gfx803.bc
main-hip-amdgcn-amd-amdhsa-gfx803.cui
main-hip-amdgcn-amd-amdhsa-gfx803.o
main-hip-amdgcn-amd-amdhsa-gfx803.out
main-hip-amdgcn-amd-amdhsa-gfx803.out.resolution.txt
main-hip-amdgcn-amd-amdhsa-gfx803.s
Files with the ``.s`` extension hold the disassembled contents of the binary.
The filename notes the graphics IPs used by the compiler. The contents of
this file are similar to what ``roc-obj`` printed to the console.
.. tab-item:: Linux and NVIDIA
:sync: linux-nvidia
Unlike HIP on AMD, when compiling using the NVIDIA support of HIP the resulting
binary will be a valid CUDA executable as far as the binary goes. Therefor
it'll incorporate PTX ISA (Parallel Thread eXecution Instruction Set
Architecture) instead of AMDGPU binary. As s result, tooling shipping with the
CUDA SDK can be used to inspect which device ISA got compiled into a specific
executable. The tool most useful to us currently is ``cuobjdump``.
.. code-block:: bash
cuobjdump --list-ptx ./saxpy
Which will print something like:
.. code-block::
PTX file 1: saxpy.1.sm_52.ptx
From this we can see that the saxpy kernel is stored as ``sm_52``, which shows
that a compute capability 5.2 ISA got embedded into the executable, so devices
which sport compute capability 5.2 or newer will be able to run this code.
.. tab-item:: Windows and AMD
:sync: windows-amd
The HIP SDK for Windows don't yet sport the ``roc-*`` set of utilities to work
with binary artifacts. To find out what binary formats are embedded into an
executable, one may use ``dumpbin`` tool from the Windows SDK to obtain the
raw data of the ``.hip_fat`` section of an executable. (This binary payload is
what gets parsed by the ``roc-*`` set of utilities on Linux.) Skipping over the
reported header, the rendered raw data as ASCII has ~3 lines per entries.
Depending on how many binaries are embedded, you may need to alter the number
of rendered lines. An invocation such as:
.. code-block:: powershell
dumpbin.exe /nologo /section:.hip_fat /rawdata:8 .\saxpy.exe | select -Skip 20 -First 12
The output may look like:
.. code-block::
000000014004C000: 5F474E414C435F5F 5F44414F4C46464F __CLANG_OFFLOAD_
000000014004C010: 5F5F454C444E5542 0000000000000002 BUNDLE__........
000000014004C020: 0000000000001000 0000000000000000 ................
000000014004C030: 0000000000000019 3638782D74736F68 ........host-x86
000000014004C040: 6E6B6E752D34365F 756E696C2D6E776F _64-unknown-linu
000000014004C050: 0000000000100078 00000000000D9800 x...............
000000014004C060: 0000000000001F00 612D347670696800 .........hipv4-a
000000014004C070: 6D612D6E6367646D 617368646D612D64 mdgcn-amd-amdhsa
000000014004C080: 3630397866672D2D 0000000000000000 --gfx906........
000000014004C090: 0000000000000000 0000000000000000 ................
000000014004C0A0: 0000000000000000 0000000000000000 ................
000000014004C0B0: 0000000000000000 0000000000000000 ................
We can see that the compiler embedded a version 4 code object (more on code
`object versions <https://www.llvm.org/docs/AMDGPUUsage.html#code-object-metadata>`_) and
used the LLVM target triple `amdgcn-amd-amdhsa--gfx906` (more on `target triples
<https://www.llvm.org/docs/AMDGPUUsage.html#target-triples>`_). Don't be
alarmed about linux showing up as a binary format, AMDGPU binaries uploaded to
the GPU for execution are proper linux ELF binaries in their format.
Alternatively we can call the compiler with ``--save-temps`` to dump all device
binary to disk in separate files.
.. code-block:: powershell
clang++ .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I .\Common -lamdhip64 -L ${env:HIP_PATH}lib -O2 --save-temps
Now we can list all the temporaries created while compiling ``main.hip`` via
.. code-block:: powershell
Get-ChildItem -Filter main-hip-* | select -Property Name
Name
----
main-hip-amdgcn-amd-amdhsa-gfx906.bc
main-hip-amdgcn-amd-amdhsa-gfx906.hipi
main-hip-amdgcn-amd-amdhsa-gfx906.o
main-hip-amdgcn-amd-amdhsa-gfx906.out
main-hip-amdgcn-amd-amdhsa-gfx906.out.resolution.txt
main-hip-amdgcn-amd-amdhsa-gfx906.s
Files with the ``.s`` extension hold the disassembled contents of the binary and
the filename directly informs us of the graphics IPs used by the compiler.
.. code-block:: powershell
Get-ChildItem main-hip-*.s | Get-Content
.text
.amdgcn_target "amdgcn-amd-amdhsa--gfx906"
.protected _Z12saxpy_kernelfPKfPfj ; -- Begin function _Z12saxpy_kernelfPKfPfj
.globl _Z12saxpy_kernelfPKfPfj
.p2align 8
.type _Z12saxpy_kernelfPKfPfj,@function
_Z12saxpy_kernelfPKfPfj: ; @_Z12saxpy_kernelfPKfPfj
; %bb.0:
s_load_dword s0, s[4:5], 0x4
s_load_dword s1, s[6:7], 0x18
s_waitcnt lgkmcnt(0)
s_and_b32 s0, s0, 0xffff
s_mul_i32 s8, s8, s0
v_add_u32_e32 v0, s8, v0
v_cmp_gt_u32_e32 vcc, s1, v0
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_2
; %bb.1:
s_load_dwordx4 s[0:3], s[6:7], 0x8
v_mov_b32_e32 v1, 0
v_lshlrev_b64 v[0:1], 2, v[0:1]
s_waitcnt lgkmcnt(0)
v_mov_b32_e32 v3, s1
v_add_co_u32_e32 v2, vcc, s0, v0
v_addc_co_u32_e32 v3, vcc, v3, v1, vcc
global_load_dword v2, v[2:3], off
v_mov_b32_e32 v3, s3
v_add_co_u32_e32 v0, vcc, s2, v0
v_addc_co_u32_e32 v1, vcc, v3, v1, vcc
global_load_dword v3, v[0:1], off
s_load_dword s0, s[6:7], 0x0
s_waitcnt vmcnt(0) lgkmcnt(0)
v_fmac_f32_e32 v3, s0, v2
global_store_dword v[0:1], v3, off
.LBB0_2:
s_endpgm
...
.. tab-item:: Windows and NVIDIA
:sync: windows-nvidia
Unlike HIP on AMD, when compiling using the NVIDIA support for HIP, the resulting
binary will be a valid CUDA executable. Therefore, it'll incorporate PTX ISA
(Parallel Thread eXecution Instruction Set Architecture) instead of AMDGPU
binary. As a result, tooling included with the CUDA SDK can be used to
inspect which device ISA was compiled into a specific executable. The most
helpful to us currently is ``cuobjdump``.
.. code-block:: bash
cuobjdump.exe --list-ptx .\saxpy.exe
Which prints something like:
.. code-block::
PTX file 1: saxpy.1.sm_52.ptx
This example shows that the SAXPY kernel is stored as ``sm_52``. It also shows
that a compute capability 5.2 ISA was embedded into the executable, so devices
that support compute capability 5.2 or newer will be able to run this code.
Now that you've found what binary got embedded into the executable, find which
format our available devices use.
.. tab-set::
.. tab-item:: Linux and AMD
:sync: linux-amd
On Linux a utility called ``rocminfo`` helps us list all the properties of the
devices available on the system, including which version of graphics IP
(``gfxXYZ``) they employ. You can filter the output to have only these lines:
.. code-block:: bash
/opt/rocm/bin/rocminfo | grep gfx
Name: gfx906
Name: amdgcn-amd-amdhsa--gfx906:sramecc+:xnack-
Now that you know which graphics IPs our devices use, recompile your program with
the appropriate parameters.
.. code-block:: bash
amdclang++ ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -lamdhip64 -L /opt/rocm/lib -O2 --offload-arch=gfx906:sramecc+:xnack-
Now the sample will run.
.. code-block::
./saxpy
Calculating y[i] = a * x[i] + y[i] over 1000000 elements.
First 10 elements of the results: [ 3, 5, 7, 9, 11, 13, 15, 17, 19, 21 ]
.. tab-item:: Linux and NVIDIA
:sync: linux-nvidia
On Linux HIP with the NVIDIA back-end, the ``deviceQuery`` CUDA SDK sample
can help us list all the properties of the devices available on the system,
including which version of compute capability a device sports.
``<major>.<minor>`` compute capability is passed to ``nvcc`` on the
command-line as ``sm_<major><minor>``, for eg. ``8.6`` is ``sm_86``.
Because it's not included as a binary, compile the matching
example from ROCm.
.. code-block:: bash
nvcc ./HIP-Basic/device_query/main.cpp -o device_query -I ./Common -I /opt/rocm/include -O2
Filter the output to have only the lines of interest, for example:
.. code-block:: bash
./device_query | grep "major.minor"
major.minor: 8.6
major.minor: 7.0
.. note::
In addition to the ``nvcc`` executable is another tool called ``__nvcc_device_query``
which prints the SM Architecture numbers to standard out as a comma
separated list of numbers. The utility's name suggests it's not a user-facing
executable but is used by ``nvcc`` to determine what devices are in the
system at hand.
Now that you know which graphics IPs our devices use, recompile your program with
the appropriate parameters.
.. code-block:: bash
nvcc ./HIP-Basic/saxpy/main.hip -o saxpy -I ./Common -I /opt/rocm/include -O2 -x cu -arch=sm_70,sm_86
.. note::
If you want to portably target the development machine which is compiling, you
may specify ``-arch=native`` instead.
Now the sample will run.
.. code-block::
./saxpy
Calculating y[i] = a * x[i] + y[i] over 1000000 elements.
First 10 elements of the results: [ 3, 5, 7, 9, 11, 13, 15, 17, 19, 21 ]
.. tab-item:: Windows and AMD
:sync: windows-amd
On Windows, a utility called ``hipInfo.exe`` helps us list all the properties
of the devices available on the system, including which version of graphics IP
(``gfxXYZ``) they employ. Filter the output to have only these lines:
.. code-block:: powershell
& ${env:HIP_PATH}bin\hipInfo.exe | Select-String gfx
gcnArchName: gfx1032
gcnArchName: gfx1035
Now that you know which graphics IPs our devices use, recompile your program with
the appropriate parameters.
.. code-block:: powershell
clang++ .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I .\Common -lamdhip64 -L ${env:HIP_PATH}lib -O2 --offload-arch=gfx1032 --offload-arch=gfx1035
Now the sample will run.
.. code-block::
.\saxpy.exe
Calculating y[i] = a * x[i] + y[i] over 1000000 elements.
First 10 elements of the results: [ 3, 5, 7, 9, 11, 13, 15, 17, 19, 21 ]
.. tab-item:: Windows and NVIDIA
:sync: windows-nvidia
On Windows HIP with the NVIDIA back-end, the ``deviceQuery`` CUDA SDK sample
can help us list all the properties of the devices available on the system,
including which version of compute capability a device sports.
``<major>.<minor>`` compute capability is passed to ``nvcc`` on the
command-line as ``sm_<major><minor>``, for eg. ``8.6`` is ``sm_86``.
Because it's not included as a binary, compile the matching
example from ROCm.
.. code-block:: powershell
nvcc .\HIP-Basic\device_query\main.cpp -o device_query.exe -I .\Common -I ${env:HIP_PATH}include -O2
Filter the output to have only the lines of interest, for example:
.. code-block:: powershell
.\device_query.exe | Select-String "major.minor"
major.minor: 8.6
major.minor: 7.0
.. note::
Next to the ``nvcc`` executable is another tool called ``__nvcc_device_query.exe``
which simply prints the SM Architecture numbers to standard out as a comma
separated list of numbers. The naming of this utility suggests it's not a user
facing executable but is used by ``nvcc`` to determine what devices are in the
system at hand.
Now that you know which graphics IPs our devices use, recompile your program with
the appropriate parameters.
.. code-block:: powershell
nvcc .\HIP-Basic\saxpy\main.hip -o saxpy.exe -I ${env:HIP_PATH}include -I .\Common -O2 -x cu -arch=sm_70,sm_86
.. note::
If you want to portably target the development machine which is compiling, you
may specify ``-arch=native`` instead.
Now the sample will run.
.. code-block::
.\saxpy.exe
Calculating y[i] = a * x[i] + y[i] over 1000000 elements.
First 10 elements of the results: [ 3, 5, 7, 9, 11, 13, 15, 17, 19, 21 ]