diff --git a/docs/data/understand/programming_model/cdna3_cu_dark.png b/docs/data/understand/programming_model/cdna3_cu_dark.png deleted file mode 100644 index 3fada0d43f..0000000000 Binary files a/docs/data/understand/programming_model/cdna3_cu_dark.png and /dev/null differ diff --git a/docs/data/understand/programming_model/rdna3_cu.drawio b/docs/data/understand/programming_model/rdna3_cu.drawio deleted file mode 100644 index e69de29bb2..0000000000 diff --git a/docs/reference/programming_model.rst b/docs/reference/programming_model.rst index e8d3b2301e..6d9fff1d09 100644 --- a/docs/reference/programming_model.rst +++ b/docs/reference/programming_model.rst @@ -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 ` 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. diff --git a/docs/understand/programming_model.rst b/docs/understand/programming_model.rst index b636f8973e..6f1d5c32fd 100644 --- a/docs/understand/programming_model.rst +++ b/docs/understand/programming_model.rst @@ -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.