AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-451170 - Merging PR#3432 from Github to Gerrit

Browse Source 
Change-Id: Ic6b44cb838dd949cd5b50e8cc0e79395b0451ca8
This commit is contained in:
Julia Jiang
2024-04-08 11:27:13 -04:00
parent 23f226ef0e
commit 4eb0e1fecc
4 changed files with 166 additions and 147 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/data/understand/programming_model/cdna3_cu_dark.png
BIN
View File
Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 141 KiB

docs/data/understand/programming_model/rdna3_cu.drawio
View File
docs/reference/programming_model.rst
+105 -82
View File
@@ -20,9 +20,8 @@ The SIMT nature of HIP is captured by the ability to execute user-provided
device programs, expressed as single-source C/C++ functions or sources compiled
online/offline to binaries in bulk.
Multiple instances of the device program (aka. kernel) may execute in parallel,
all uniquely identified by a set of integral values which are referred to as
thread IDs. The set of integers identifying a thread relate to the hierarchy in
Multiple instances of the device program (aka. kernel), so called threads, may execute in parallel,
all uniquely identified by a set of integral values. The set of integers identifying a thread relate to the hierarchy in
which threads execute.
.. _inherent_thread_model:
@@ -30,25 +29,27 @@ which threads execute.
Inherent Thread Model
-------------------------------------------------------------------------------
The thread hiearchy inherent to how AMD GPUs operate manifest as depicted in
The thread hierarchy inherent to how AMD GPUs operate is depicted in
:numref:`inherent_thread_hierarchy`.
.. _inherent_thread_hierarchy:
.. figure:: ../data/reference/programming_model/thread_hierarchy.svg
:alt: Diagram depicting nested rectangles of varying color. The outermost one
titled "Grid", inside sets of uniform rectangles layered on oneanother
titled "Block". Each "Block" containing sets of uniform rectangles
layered on oneanother titled "Warp". Each of the "Warp" titled
rectangles filled with downward pointing arrows inside.
:alt: Diagram depicting the thread hierarchy in HIP using nested rectangles.
The outermost rectangle is titled Grid, containing uniform rectangles
layered on one another titled Block. Each Block contains sets of uniform
rectangles layered on one another titled Warp. Each of the Warp titled
rectangles is filled with downward pointing arrows, representing single
threads.
Hierarchy of thread groups.
Hierarchy of thread groups. The arrows represent the threads.
* The innermost grouping is called a warp, or a wavefront in ISA terms. A warp
is the most tightly coupled groups of threads, both physically and logically.
Warp
The most tightly coupled group of threads, both physically and logically.
Also known as a wavefront in ISA (Instruction Set Architecture) terms.
When referring to threads inside a warp, they may be called lanes, and the
integral value identifying them the lane ID. Lane IDs aren't quieried like
integral value identifying them the lane ID. Lane IDs aren't queried like
other thread IDs, but are user-calculated. As a consequence they are only as
multi-dimensional as the user interprets the calculated values to be.
@@ -56,20 +57,22 @@ The thread hiearchy inherent to how AMD GPUs operate manifest as depicted in
signified by the set of communication primitives at their disposal, detailed
under :ref:`warp_cross_lane_functions`.
* The middle grouping is called a block or thread block. The defining feature
of a block is that all threads in a block will share an instance of memory
which they may use to share data or synchronize with oneanother.
Block
The defining feature of a block (or thread block) is that all threads in the
same block are executed on the same Compute Unit, share an instance of
memory which they may use to share data, and can synchronize with one another.
The size of a block is user-configurable but is maxmized by the queryable
capabilites of the executing hardware. The unique ID of the thread within a
block is 3-dimensional as provided by the API. When linearizing thread IDs
within a block, assume the "fast index" being dimension ``x``, followed by
the ``y`` and ``z`` dimensions.
The size of a block is user-configurable but is maximized by the queryable
capabilities of the executing hardware. The unique ID of the thread within a
block is 3-dimensional as provided by the `threadIdx` built-in. When
linearizing thread IDs within a block, assume the "fast index" being
dimension ``x``, followed by the ``y`` and ``z`` dimensions.
* The outermost grouping is called a grid. A grid manifests as a single
dispatch of kernels for execution. The unique ID of each block within a grid
is 3-dimensional, as provided by the API and is queryable by every thread
within the block.
Grid
The outermost grouping is called a grid. A grid manifests as a single
dispatch of a kernel for execution. The unique ID of each block within a grid
is 3-dimensional, as provided by the `blockIdx` built-in and is queryable
by every thread within the block.
Cooperative Groups Thread Model
-------------------------------------------------------------------------------
@@ -82,100 +85,120 @@ imposed by the strict 1:1 mapping of architectural details to the programming
model.
The rich set of APIs introduced by Cooperative Groups allow the programmer
to define their own groups based on run-time predicates, but a set of implicit
groups manifest based on kernel launch parameters.
to define their own set of thread groups which may fit their user-cases better
than those defined by the hardware. The set of implicit groups by kernel launch
parameters are still available.
The thread hiearchy abstraction of Cooperative Groups manifest as depicted in
The thread hierarchy abstraction of Cooperative Groups manifest as depicted in
:numref:`coop_thread_hierarchy`.
.. _coop_thread_hierarchy:
.. figure:: ../data/reference/programming_model/thread_hierarchy_coop.svg
:alt: Diagram depicting nested rectangles of varying color. The outermost one
titled "Grid", inside sets of different sized rectangles layered on
oneanother titled "Block". Each "Block" containing sets of uniform
rectangles layered on oneanother titled "Warp". Each of the "Warp"
titled rectangles filled with downward pointing arrows inside.
:alt: Diagram depicting the structure of Cooperative Groups using nested
rectangles. The outermost rectangle is titled Multi Grid, containing
sets of different shaped rectangles titled Grid. Each Grid contains sets
of uniform rectangles layered on oneanother titled Cluster. The Clusters
have different shapes in different Grids. Inside of the Clusters are
uniform rectangles layered on each other titled Block, which include
arrows that represent threads.
Cooperative group thread hierarchy.
* Multi Grid is an abstraction of potentially multiple simultaneous launches of
the same kernel over multiple devices. Grids inside a multi device kernel
launch need not be of uniform size, thus allowing taking into account
different device capabilities and preferences.
Multi Grid
An abstraction of potentially multiple simultaneous launches of the same
kernel over multiple devices. Grids inside a multi device kernel launch need
not be of uniform size, thus allowing taking into account different device
capabilities and preferences.
.. deprecated:: 5.0
* Same as the :ref:`inherent_thread_model` Grid entity. The ability to
.. note::
The performance hit of implementing cooperative groups spanning multiple
devices proved to outweigh the provided convenience, hence the use of
Multi Grid is discouraged.
Grid
Same as the :ref:`inherent_thread_model` Grid entity. The ability to
synchronize over a grid requires the kernel to be launched using the
Cooperative Groups API.
* The defining feature of a cluster or block cluster is that all threads in a
cluster will share a common set of distributed shared memory which they may
use to share data or synchronize with oneanother.
Cluster block
The defining feature of a cluster or block cluster is that all threads in a
cluster will use a common set of distributed shared memory which they may
use to share data or synchronize with one another.
* Same as the :ref:`inherent_thread_model` Block entity.
Block
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, one may use the group partitioning
part of the API, typically but not necessarily ``tiled_partition``.
part of the API, for instance, ``tiled_partition``.
Memory Model
===============================================================================
The hierarchy of threads introduced by :ref:`inherent_thread_model` is induced
by the memory subsystem of GPUs. :numref:`memory_hierarchy` summarizes that memory namespaces and
how they relate to the various levels of the threading model.
The hierarchy of threads introduced by the :ref:`inherent_thread_model` is
induced by the memory subsystem of GPUs. :numref:`memory_hierarchy` summarizes
the memory namespaces and how they relate to the various levels of the threading
model.
.. _memory_hierarchy:
.. figure:: ../data/reference/programming_model/memory_hierarchy.svg
:alt: Diagram depicting nested rectangles of varying color. The outermost one
titled "Grid", inside on the upper half a rectangle titled "Cluster".
Inside it are two identical rectangles titled "Block", inside them are
ones titled "Local" with multiple "Warp" titled rectangles. Blocks have
not just Local inside, but also rectangles titled "Shared". The Shared
rectangles of Blocks in the same Cluster are grouped together with a
translucent halo titled "Cluster shared". Outside the Cluster but
inside the Grid is a rectangle titled "Global" with three others
inside: "Surface", "Texture" (same color) and "Constant" (different
color).
:alt: Diagram depicting the memory hierarchy using nested rectangles. The
outermost is title Grid, containing two rectangles, one titled Cluster
and the other titled Global. Cluster contains two identical rectangles
titled Block, which are partly overlaid and connected by a rectangle
titled Cluster Shared. The Block rectangles each contain a rectangle
titled Local, which in turn contain rectangles titled Warp that include
arrows representing the threads. Cluster shared contains two rectangles
titled Shared, each located within one of the Blocks. Global contains
three rectangles, titled Constant, Texture and Surface.
Memory hierarchy.
* Local or per-thread memory is read-write storage only visible to the
threads defining the given variables. The size of a block for a given kernel,
Local
Read-write storage only visible to the threads defining the given variables,
also called per-threas memory. The size of a block for a given kernel,
the number of concurrent warps are limited by local memory usage.
This relates to an important aspect: occupancy. This is the default memory
namespace.
* Shared memory is read-write storage visible to all the threads in a given
block.
Shared
Read-write storage visible to all the threads in a given block.
* Distributed shared memory is read-write storage visible to all the threads
in a given block cluster.
Cluster shared
Read-write storage visible to all the threads in a given cluster.
* Global memory is read-write storage visible to all threads in a given grid.
There are specialized versions of global memory with different usage
semantics which are typically backed by the same hardware storing global.
Global
Read-write storage visible to all threads in a given grid. There are
specialized versions of global memory with different usage semantics which
are typically backed by the same hardware.
* Constant memory is read-only storage visible to all threads in a given
grid. It is a limited segment of global with queryable size.
Constant
Read-only storage visible to all threads in a given grid. It is a limited
segment of global with queryable size.
* Texture memory is read-only storage visible to all threads in a given grid
and accessible through additional APIs.
Texture
Read-only storage visible to all threads in a given grid and accessible
through additional APIs.
* Surface is a writable version of texture memory.
Surface
Read-write version of texture memory.
Execution Model
===============================================================================
HIP programs consist of two distinct scopes:
* The host-side API running on the host processor. There are to APIs available:
* The host-side running on the host processor.
There are two types of APIs available:
* The HIP runtime API which enables use of the single-source programming
model.
@@ -185,8 +208,9 @@ HIP programs consist of two distinct scopes:
importantly around kernel launching and argument setting. It is geared
towards implementing abstractions atop, such as the runtime API itself.
* The device-side kernels running on GPUs. Both the host and the device-side
APIs have synchronous and asynchronous functions in them.
* The device-side kernels running on GPUs.
Both the host and the device-side APIs have synchronous and asynchronous functions in them.
Host-side execution
-------------------------------------------------------------------------------
@@ -196,7 +220,7 @@ queries are synchronous. All asynchronous APIs, such as kernel execution, data
movement and potentially data allocation/freeing all happen in the context of
device streams.
Streams are FIFO buffers of commands to execute relating to a given device.
Streams are FIFO buffers of commands to execute on a given device.
Commands which enqueue tasks on a stream all return promptly and the command is
executed asynchronously. All side-effects of a command on a stream are visible
to all subsequent commands on the same stream. Multiple streams may point to
@@ -206,10 +230,10 @@ be.
Asynchronous APIs involving a stream all return a stream event which may be
used to synchronize the execution of multiple streams. A user may enqueue a
barrier onto a stream referencing an event. The barrier will will block until
the command related to the event does not complete, at which point all
side-effects of the command shall be visible to commands following the barrier,
even if those side-effects manifest on different devices.
barrier onto a stream referencing an event. The barrier will block until all
commands related to the event complete, at which point all side-effects of
the commands are visible to commands following the barrier, even if those
side-effects manifest on different devices.
Streams also support executing user-defined functions as callbacks on the host.
The stream will not launch subsequent commands until the callback completes.
@@ -242,18 +266,17 @@ intended use-cases.
.. tip::
This name by default is a macro expanding to triple-chevron. In cases where
``hipLaunchKernelGGL()`` by default is a macro expanding to triple-chevron. In cases where
language syntax extensions are undesirable, or where launching templated
and/or overloaded kernel functions define the
``HIP_TEMPLATE_KERNEL_LAUNCH`` preprocessor macro before including the HIP
headers to turn it into a templated function.
* Using the
:doxygen:`launch APIs supporting the triple-chevron syntax <Clang>` directly.
* Using the launch APIs supporting the triple-chevron syntax directly.
.. caution::
These APIs are intended to be used/generated by tools such as the HIP
compiler itself and not intended towards end-user code. Should you be
compiler itself and not intended for end-user code. Should you be
writing a tool having to launch device code using HIP, consider using these
over the alternatives.
docs/understand/programming_model.rst
+61 -65
View File
@@ -8,64 +8,61 @@
Programming Model
*******************************************************************************
The HIP programming model makes it as easy as reasonably possible to map
data-parallel C/C++ algorithms and map them to massively parallel, wide SIMD
architectures, such as GPUs. As a consequence, one needs a basic understanding
of the underlying device architecture to make efficient use of HIP and GPGPU
(General Purpose Graphics Processing Unit) programming in general.
The HIP programming model makes it easy to map data-parallel C/C++ algorithms to
massively parallel, wide SIMD architectures, such as GPUs. As a consequence, one
needs a basic understanding of the underlying device architecture to make efficient
use of HIP and GPGPU (General Purpose Graphics Processing Unit) programming in general.
RDNA & CDNA Architecture Summary
===============================================================================
Most GPU architectures, much like RDNA and CDNA have a hierarchical structure.
The inner-most piece is some Single Instruction Multiple Data (SIMD) enabled
vector Arithmetic Logical Unit (ALU). Most recent GPUs beside the vector ALU
also house some matrix ALU for accelerating algorithms of well defined shapes.
Most GPU architectures, like RDNA and CDNA, have a hierarchical structure.
The inner-most piece is a Single Instruction Multiple Data (SIMD) enabled
vector Arithmetic Logical Unit (ALU). In addition to the vector ALUs, most
recent GPUs also house also house matrix ALUs for accelerating algorithms involving
matrix multiply-accumulate operations. AMD GPUs also contain scalar ALUs, that
can be used to reduce the load on the vector ALU by performing operations which
are uniform for all threads of a warp.
Some number of vector and matrix ALUs comprise a larger block, often referred
to as a Compute Unit (OpenCL, AMD block diagrams) but is referred to as Multi
Processor in HIP terms.
A set of ALUs, together with register files, caches and shared memory, comprise
a larger block, often referred to as a Compute Unit (CU), e.g. in OpenCL and
AMD block diagrams, or as Streaming Multiprocessor (SM).
.. _rdna3_cu:
.. figure:: ../data/understand/programming_model/rdna3_cu.png
:alt: Block diagram showing components mostly duplicated on the upper and
lower halves of the image with some spanning over both parts (hinting
at them being shared). Both the top and the bottom have two sets of
identical hardware blocks. a "Scheduler" having "Vector GPR"s with an
associated Vector ALU with the following noted capabilities:
Float/INT/Matrix SIMD32, Float/Matrix SIMD32, Transcendental SIMD8, AI
MATRIX Accelerator, DPFP (1). Both top and bottom have a single "Ray
Accelerator", "Texture Filters" and "LD/ST/Tex Addr" and L0 blocks.
Shared among top and bottom are "Scalar Cache", "Shader Instruction
Cache" and Shared Memory.
:alt: Block diagram showing the structure of an RDNA3 Compute Unit. It
consists of four SIMD units, each including a vector and scalar register
file, with the corresponding scalar and vector ALUs. All four SIMDs
share a scalar and instruction cache, as well as the shared memory. Two
of the SIMD units each share an L0 cache.
Block Diagram of an RDNA3 Compute Unit.
.. _cdna3_cu:
.. figure:: ../data/understand/programming_model/cdna3_cu.png
:alt: Block diagram showing components: "Scheduler", "Local Data Share",
"Matrix Core Unit", "Shader Core" and "L1 Cache" along with some
unnamed blocks.
:alt: Block diagram showing the structure of a CDNA3 compute unit. It includes
Shader Cores, the Matrix Core Unit, a Local Data Share used for sharing
memory between threads in a block, an L1 Cache and a Scheduler. The
Shader Cores represent the vector ALUs and the Matrix Core Unit the
matrix ALUs. The Local Data Share is used as the shared memory.
Block Diagram of an CDNA3 Compute Unit.
Block Diagram of a CDNA3 Compute Unit.
For hardware implementation's sake, some number of Multi Processors are grouped
together into a Shader Engine or Compute Engine, typically sharing some fixed
function units or memory subsystem resources.
FFor implementation in hardware, multiple Compute Units are grouped together into
a Shader Engine or Compute Engine, typically sharing some fixed function units or
memory subsystem resources.
.. _cdna2_gcd:
.. figure:: ../data/understand/programming_model/cdna2_gcd.png
:alt: Block diagram showing four "Compute Engine"s each with 28 "CU"s
(Compute Unit) inside. These four Compute Engines share one block of
"L2 Cache and Controllers". Around them are four "Memory Controller"s,
each having a "Memory Phy" block next to them. To the top and bottom of
all these are eight blocks of "Infinity Fabric Link" with one of the
eight reading "Infinity Fabric or PCIe". Two sole "VCN" blocks sit in
top corners. At the very bottom spans a colored seaction reading
"Infinity Fabric" along with its logo.
:alt: Block diagram showing four Compute Engines each with 28 Compute Units
inside. These four Compute Engines share one block of L2 Cache. Around
them are four Memory Controllers. To the top and bottom of all these are
eight blocks of Infinity Fabric Links. Two Video Core Next blocks sit in
the top corners. At the very bottom spans a colored section reading
Infinity Fabric.
Block Diagram of a CDNA2 Graphics Compute Die.
@@ -75,14 +72,13 @@ Single Instruction Multiple Threads
The SIMT programming model behind the HIP device-side execution is a
middle-ground between SMT (Simultaneous Multi-Threading) programming known from
multi-core CPUs, and SIMD (Single Instruction, Multiple Data) programming
mostly known from exploiting relevant instruction sets on CPUs (eg.
SSE/AVX/Neon).
mostly known from exploiting relevant instruction sets on CPUs (eg. SSE/AVX/Neon).
A HIP device compiler maps our SIMT code written in HIP C++ to an inherently
SIMD architecture (like GPUs) not by exploiting data parallelism within a
single instance of a kernel and spreading identical instructions over the SIMD
engines at hand, but by scalarizing the entire kernel and issuing the scalar
instructions of multiple kernel instances to each of the SIMD engine lanes.
A HIP device compiler maps SIMT code written in HIP C++ to an inherently SIMD
architecture (like GPUs), by scalarizing the entire kernel, and issuing the scalar
instructions of multiple kernel instances to each of the SIMD engine lanes, rather
than exploiting data parallelism within a single instance of a kernel and spreading
identical instructions over the available SIMD engines.
Consider the following kernel
@@ -93,31 +89,31 @@ Consider the following kernel
int tid = threadIdx.x;
int bid = blockIdx.x;
int dim = blockDim.x;
a[tid] += (tid + bid - dim) * b[tid];
}
The incoming four-vector of floating-point values ``a`` is multiplied by a
scalar and then multiplied element-wise by another four-vector. On modern
SIMD-capable architectures the four-vector ops are expected to compile to a
single SIMD instruction. GPU execution of this kernel however will typically
look the following:
The incoming four-vector of floating-point values ``b`` is multiplied by a
scalar and then added element-wise to the four-vector floating-point values of
``a``. On modern SIMD-capable architectures the four-vector ops are expected to
compile to a single SIMD instruction. GPU execution of this kernel however will
typically look the following:
.. _simt:
.. figure:: ../data/understand/programming_model/simt.svg
:alt: Two large arrows pointing downward with blocks inside and ellipses
between the arrows. Inside the arrows the same series of blocks with
the following texts inside from top to bottom: "ADD", "DIV", "FMA",
"FMA", "FMA" and "FMA".
:alt: Image representing the instruction flow of a SIMT program. Two identical
arrows pointing downward with blocks representing the instructions
inside and ellipsis between the arrows. The instructions represented in
the arrows are, from top to bottom: ADD, DIV, FMA, FMA, FMA and FMA.
Instruction flow of the sample SIMT program.
In HIP, lanes of a SIMD architecture are fed by mapping threads of a SIMT
execution, one thread down each lane of a SIMD engine. Execution parallelism
isn't exploited from the width of the built-in vector types, but via the thread
id constants ``threadIdx.x``, ``blockIdx.x``, etc. For more details, refer to
:ref:`inherent_thread_model`.
usually isn't exploited from the width of the built-in vector types, but via the
thread id constants ``threadIdx.x``, ``blockIdx.x``, etc. For more details,
refer to :ref:`inherent_thread_model`.
Heterogenous Programming
===============================================================================
@@ -136,23 +132,23 @@ a few key differences between the two:
from one means nothing in another. Moreover not all address spaces are
accessible from all contexts.
If one were to look at {ref}`cdna2_gcd` and inside the {ref}`cdna3_cu`,
If one were to look at :ref:`cdna2_gcd` and inside the :ref:`cdna3_cu`,
every Compute Unit has an instance of storage backing the namespace
``__shared__``. Even if the host were to have access to these regions of
memory, the performance benefits of the segmented memory subsystem are
supported by the inability of meaningful asynchronous accesss from the host.
supported by the inability of asynchronous access from the host.
* Not all C++ language features map cleanly to typical device architectures,
some are very expensive (meaning: slow) to implement on GPU devices, therefor
some are very expensive (meaning slow) to implement on GPU devices, therefor
they are forbidden in device contexts to avoid users tapping into features
unexpectedly decimating their program's performance. Offload devices targeted
that unexpectedly decimate their program's performance. Offload devices targeted
by HIP aren't general purpose devices, at least not in the sense a CPU is.
HIP focuses on data parallel computations and as such caters to throughput
optimized architectures, such as GPUs or accelerators derived from GPU
architectures.
* Asynchrony is at the forefront of the HIP API. Computations launched by HIP
execute asynchronously on the device and it is the user's responsibility to
* Asynchrony is at the forefront of the HIP API. Computations launched on the device
execute asynchronously with respect to the host and it is the user's responsibility to
synchronize their data dispatch/fetch with computations on the device. HIP
does perform implicit synchronization on occasions, but unlike some APIs
(OpenCL, SYCL) by and large places the onus of synchronization on the user.
does perform implicit synchronization on occasions, more advanced than other APIs such as
OpenCL or SYCL, in which the responsibility of synchronization mostly depends on the user.