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

Docs: refactor and integrate into ROCm docs portal (#362)

Browse Source 
* pip-compile docs/requirements.txt

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Add Sphinx docs config

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Add Sphinx config

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Update docs build config

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* style(conf.py): Apply black formatting to docs/conf.py

Signed-off-by: Sam Wu <22262939+samjwu@users.noreply.github.com>

* Update docs requirements

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Update to rocm-docs-core 1.3.0

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Update docs requirements

Signed-off-by: Peter Jun Park <peter.park@amd.com>

pip-compile requirements

Signed-off-by: Peter Jun Park <peter.park@amd.com>

bump rocm-docs-core to 1.5.0

bump rocm-docs-core to 1.4.1

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* Add dependabot.yml and update CODEOWNERS

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Update toc and conf

Signed-off-by: Peter Jun Park <peter.park@amd.com>

update dependabot

* Port docs to rocm-docs standard

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Add toc and Diataxis cards

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Add basic file structure

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add glossary

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add includes

Signed-off-by: Peter Jun Park <peter.park@amd.com>

Add license.rst

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add compatible hw

Signed-off-by: Peter Jun Park <peter.park@amd.com>

fix spelling and license

Signed-off-by: Peter Jun Park <peter.park@amd.com>

clean up index

Signed-off-by: Peter Jun Park <peter.park@amd.com>

clean up installation guides

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add basic usage (quickstart)

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add ref to global options

update toc

Signed-off-by: Peter Jun Park <peter.park@amd.com>

modularize modes and global options

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add profile mode

Signed-off-by: Peter Jun Park <peter.park@amd.com>

fixes

Signed-off-by: Peter Jun Park <peter.park@amd.com>

reorg and clean up

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add dynamic omniperf version number in installation guide

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add datatemplate

more reorg

Signed-off-by: Peter Jun Park <peter.park@amd.com>

clean up

Signed-off-by: Peter Jun Park <peter.park@amd.com>

reorg images

move profile mode

reorg

reorg

reorg more

fix formatting

fix headings

ref anchor mi2xx note

add extlinks

add extlinks

Signed-off-by: Peter Jun Park <peter.park@amd.com>

black format

fix formatting, anchors

Signed-off-by: Peter Jun Park <peter.park@amd.com>

reorg

fix words and formatting

Signed-off-by: Peter Jun Park <peter.park@amd.com>

formatting

Signed-off-by: Peter Jun Park <peter.park@amd.com>

same

reorg

format

fix formatting

fix toc

Signed-off-by: Peter Jun Park <peter.park@amd.com>

format

* impr internal linking and fix sphinx warnings

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* add spellcheck/linting from rocm-docs-core

Signed-off-by: Peter Jun Park <peter.park@amd.com>

fix rst directives

satisfy spellcheck

fix more spelling

rm unused files

fix spelling and update wordlist

* bump rocm-docs-core to 1.6.0

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* add fixes from @skyreflectedinmirrors and @lpaoletti

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add references to toc

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add more fixes

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* add package manager install section

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* add fixes

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add metadata and fixes

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add fixes

bump to 1.6.1

more fixes

fix fmt in profiling examples

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add missing mem type table

Signed-off-by: Peter Jun Park <peter.park@amd.com>

fix formatting

fmt

* add custom css

Signed-off-by: Peter Jun Park <peter.park@amd.com>

fix css fs

* make images/figs click-to-expand

Signed-off-by: Peter Jun Park <peter.park@amd.com>

add missed image

update

fix link

* update documentation link in README

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* formatting fixes

Signed-off-by: Peter Jun Park <peter.park@amd.com>

more formatting

* fix heading

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* move archived docs

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* exclude archived docs from docs build

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* update archived docs workflow

Signed-off-by: Peter Jun Park <peter.park@amd.com>

move files

update archived docs workflow

Signed-off-by: Peter Jun Park <peter.park@amd.com>

fix version number

clean up workflow

workflow test

workflow test

another workflow test

* rm docs linting

Signed-off-by: Peter Jun Park <peter.park@amd.com>

* Apply cmake-format suggested changes

Signed-off-by: Sam Wu <22262939+samjwu@users.noreply.github.com>

* Apply cmake-format

Signed-off-by: Sam Wu <22262939+samjwu@users.noreply.github.com>

---------

Signed-off-by: Peter Jun Park <peter.park@amd.com>
Signed-off-by: Sam Wu <22262939+samjwu@users.noreply.github.com>
Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com>
This commit is contained in:
Peter Park
2024-07-31 10:42:27 -04:00
committed by David Galiffi
parent 96803e327c
commit a0dc485ceb
307 changed files with 16783 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/conceptual/command-processor.rst
+154
View File
@@ -0,0 +1,154 @@
.. meta::
:description: Omniperf performance model: Command processor (CP)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, command, processor, fetcher, packet processor, CPF, CPC
**********************
Command processor (CP)
**********************
The command processor (CP) is responsible for interacting with the AMDGPU kernel
driver -- the Linux kernel -- on the CPU and for interacting with user-space
HSA clients when they submit commands to HSA queues. Basic tasks of the CP
include reading commands (such as, corresponding to a kernel launch) out of
:hsa-runtime-pdf:`HSA queues <68>`, scheduling work to subsequent parts of the
scheduler pipeline, and marking kernels complete for synchronization events on
the host.
The command processor consists of two sub-components:
* :ref:`Fetcher <cpf-metrics>` (CPF): Fetches commands out of memory to hand
them over to the CPC for processing.
* :ref:`Packet processor <cpc-metrics>` (CPC): Micro-controller running the
command processing firmware that decodes the fetched commands and (for
kernels) passes them to the :ref:`workgroup processors <desc-spi>` for
scheduling.
Before scheduling work to the accelerator, the command processor can
first acquire a memory fence to ensure system consistency
(:hsa-runtime-pdf:`Section 2.6.4 <91>`). After the work is complete, the
command processor can apply a memory-release fence. Depending on the AMD CDNA™
accelerator under question, either of these operations *might* initiate a cache
write-back or invalidation.
Analyzing command processor performance is most interesting for kernels
that you suspect to be limited by scheduling or launch rate. The command
processor’s metrics therefore are focused on reporting, for example:
* Utilization of the fetcher
* Utilization of the packet processor, and decoding processing packets
* Stalls in fetching and processing
.. _cpf-metrics:
Command processor fetcher (CPF)
===============================
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - CPF Utilization
- Percent of total cycles where the CPF was busy actively doing any work.
The ratio of CPF busy cycles over total cycles counted by the CPF.
- Percent
* - CPF Stall
- Percent of CPF busy cycles where the CPF was stalled for any reason.
- Percent
* - CPF-L2 Utilization
- Percent of total cycles counted by the CPF-:doc:`L2 <l2-cache>` interface
where the CPF-L2 interface was active doing any work. The ratio of CPF-L2
busy cycles over total cycles counted by the CPF-L2.
- Percent
* - CPF-L2 Stall
- Percent of CPF-:doc:`L2 <l2-cache>` L2 busy cycles where the CPF-L2
interface was stalled for any reason.
- Percent
* - CPF-UTCL1 Stall
- Percent of CPF busy cycles where the CPF was stalled by address
translation.
- Percent
.. _cpc-metrics:
Command processor packet processor (CPC)
========================================
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - CPC Utilization
- Percent of total cycles where the CPC was busy actively doing any work.
The ratio of CPC busy cycles over total cycles counted by the CPC.
- Percent
* - CPC Stall
- Percent of CPC busy cycles where the CPC was stalled for any reason.
- Percent
* - CPC Packet Decoding Utilization
- Percent of CPC busy cycles spent decoding commands for processing.
- Percent
* - CPC-Workgroup Manager Utilization
- Percent of CPC busy cycles spent dispatching workgroups to the
:ref:`workgroup manager <desc-spi>`.
- Percent
* - CPC-L2 Utilization
- Percent of total cycles counted by the CPC-:doc:`L2 <l2-cache>` interface
where the CPC-L2 interface was active doing any work.
- Percent
* - CPC-UTCL1 Stall
- Percent of CPC busy cycles where the CPC was stalled by address
translation.
- Percent
* - CPC-UTCL2 Utilization
- Percent of total cycles counted by the CPC's :doc:`L2 <l2-cache>` address
translation interface where the CPC was busy doing address translation
work.
- Percent
docs/conceptual/compute-unit.rst
+60
View File
@@ -0,0 +1,60 @@
.. meta::
:description: Omniperf performance model: Compute unit (CU)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, GCN, compute, unit, pipeline, workgroup, wavefront,
CDNA
*****************
Compute unit (CU)
*****************
The compute unit (CU) is responsible for executing a user's kernels on
CDNA™-based accelerators. All :ref:`wavefronts <desc-wavefront>` of a
:ref:`workgroup <desc-workgroup>` are scheduled on the same CU.
.. image:: ../data/performance-model/gcn_compute_unit.png
:align: center
:alt: AMD CDNA accelerator compute unit diagram
:width: 800
The CU consists of several independent execution pipelines and functional units.
The :doc:`/conceptual/pipeline-descriptions` section details the various
execution pipelines -- VALU, SALU, LDS, scheduler, and so forth. The metrics
presented by Omniperf for these pipelines are described in
:doc:`pipeline-metrics`. The :doc:`vL1D <vector-l1-cache>` cache and
:doc:`LDS <local-data-share>` are described in their own sections.
* The :ref:`desc-valu` is composed of multiple SIMD (single
instruction, multiple data) vector processors, vector general purpose
registers (VGPRs) and instruction buffers. The VALU is responsible for
executing much of the computational work on CDNA accelerators, including but
not limited to floating-point operations (FLOPs) and integer operations
(IOPs).
* The vector memory (VMEM) unit is responsible for issuing loads, stores and
atomic operations that interact with the memory system.
* The :ref:`desc-salu` is shared by all threads in a
:ref:`wavefront <desc-wavefront>`, and is responsible for executing
instructions that are known to be uniform across the wavefront at compile
time. The SALU has a memory unit (SMEM) for interacting with memory, but it
cannot issue separately from the SALU.
* The :doc:`local-data-share` is an on-CU software-managed scratchpad memory
that can be used to efficiently share data between all threads in a
:ref:`workgroup <desc-workgroup>`.
* The :ref:`desc-scheduler` is responsible for issuing and decoding instructions
for all the :ref:`wavefronts <desc-wavefront>` on the compute unit.
* The :doc:`vector L1 data cache (vL1D) <vector-l1-cache>` is the first level
cache local to the compute unit. On current CDNA accelerators, the vL1D is
write-through. The vL1D caches from multiple compute units are kept coherent
with one another through software instructions.
* CDNA accelerators -- that is, AMD Instinct™ MI100 and newer -- contain
specialized matrix-multiplication accelerator pipelines known as the
:ref:`desc-mfma`.
For a more in-depth description of a compute unit on a CDNA accelerator, see
:hip-training-pdf:`22` and :gcn-crash-course:`27`.
docs/conceptual/definitions.rst
+152
View File
@@ -0,0 +1,152 @@
.. meta::
:description: Omniperf terminology and definitions
:keywords: Omniperf, ROCm, glossary, definitions, terms, profiler, tool,
Instinct, accelerator, AMD
***********
Definitions
***********
The following table briefly defines some terminology used in Omniperf interfaces
and in this documentation.
.. include:: ./includes/terms.rst
.. include:: ./includes/normalization-units.rst
.. _memory-spaces:
Memory spaces
=============
AMD Instinct™ MI-series accelerators can access memory through multiple address spaces
which may map to different physical memory locations on the system. The
following table provides a view into how various types of memory used
in HIP map onto these constructs:
.. list-table::
:header-rows: 1
* - LLVM Address Space
- Hardware Memory Space
- HIP Terminology
* - Generic
- Flat
- N/A
* - Global
- Global
- Global
* - Local
- LDS
- LDS/Shared
* - Private
- Scratch
- Private
* - Constant
- Same as global
- Constant
The following is a high-level description of the address spaces in the AMDGPU
backend of LLVM:
.. list-table::
:header-rows: 1
* - Address space
- Description
* - Global
- Memory that can be seen by all threads in a process, and may be backed by
the local accelerator's HBM, a remote accelerator's HBM, or the CPU's
DRAM.
* - Local
- Memory that is only visible to a particular workgroup. On AMD's Instinct
accelerator hardware, this is stored in :doc:`LDS <local-data-share>`
memory.
* - Private
- Memory that is only visible to a particular [work-item](workitem)
(thread), stored in the scratch space on AMD's Instinct accelerators.
* - Constant
- Read-only memory that is in the global address space and stored on the
local accelerator's HBM.
* - Generic
- Used when the compiler cannot statically prove that a pointer is
addressing memory in a single (non-generic) address space. Mapped to Flat
on AMD's Instinct accelerators, the pointer could dynamically address
global, local, private or constant memory.
`LLVM's documentation for AMDGPU Backend <https://llvm.org/docs/AMDGPUUsage.html#address-spaces>`_
has the most up-to-date information. Refer to this source for a more complete
explanation.
.. _memory-type:
Memory type
===========
AMD Instinct accelerators contain a number of different memory allocation
types to enable the HIP language's
:doc:`memory coherency model <hip:how-to/programming_manual>`.
These memory types are broadly similar between AMD Instinct accelerator
generations, but may differ in exact implementation.
In addition, these memory types *might* differ between accelerators on the same
system, even when accessing the same memory allocation.
For example, an :ref:`MI2XX <mixxx-note>` accelerator accessing *fine-grained*
memory allocated local to that device may see the allocation as coherently
cacheable, while a remote accelerator might see the same allocation as
*uncached*.
These memory types include:
.. list-table::
:header-rows: 1
* - Memory type
- Description
* - Uncached Memory (UC)
- Memory that will not be cached in this accelerator. On
:ref:`MI2XX <mixxx-note>` accelerators, this corresponds “fine-grained”
(or, “coherent”) memory allocated on a remote accelerator or the host,
for example, using ``hipHostMalloc`` or ``hipMallocManaged`` with default
allocation flags.
* - Non-hardware-Coherent Memory (NC)
- Memory that will be cached by the accelerator, and is only guaranteed to
be consistent at kernel boundaries / after software-driven
synchronization events. On :ref:`MI2XX <mixxx-note>` accelerators, this
type of memory maps to, for example, “coarse-grained” ``hipHostMalloc``’d
memory -- that is, allocated with the ``hipHostMallocNonCoherent``
flag -- or ``hipMalloc``’d memory allocated on a remote accelerator.
* - Coherently Cachable (CC)
- Memory for which only reads from the accelerator where the memory was
allocated will be cached. Writes to CC memory are uncached, and trigger
invalidations of any line within this accelerator. On
:ref:`MI2XX <mixxx-note>` accelerators, this type of memory maps to
“fine-grained” memory allocated on the local accelerator using, for
example, the ``hipExtMallocWithFlags`` API using the
``hipDeviceMallocFinegrained`` flag.
* - Read/Write Coherent Memory (RW)
- Memory that will be cached by the accelerator, but may be invalidated by
writes from remote devices at kernel boundaries / after software-driven
synchronization events. On :ref:`MI2XX <mixxx-note>` accelerators, this
corresponds to “coarse-grained” memory allocated locally to the
accelerator, using for example, the default ``hipMalloc`` allocator.
Find a good discussion of coarse and fine-grained memory allocations and what
type of memory is returned by various combinations of memory allocators, flags
and arguments in the
`Crusher quick-start guide <https://docs.olcf.ornl.gov/systems/crusher_quick_start_guide.html#floating-point-fp-atomic-operations-and-coarse-fine-grained-memory-allocations>`_.
docs/conceptual/includes/normalization-units.rst
+47
View File
@@ -0,0 +1,47 @@
.. _normalization-units:
Normalization units
===================
A user-configurable unit by which you can choose to normalize data. Options
include:
.. list-table::
:header-rows: 1
* - Name
- Description
* - ``per_wave``
- The total value of the measured counter or metric that occurred per
kernel invocation divided by the total number of
:ref:`wavefronts <desc-wavefront>` launched in the kernel.
* - ``per_cycle``
- The total value of the measured counter or metric that occurred per
kernel invocation divided by the
:ref:`kernel cycles <kernel-cycles>`, that is, the total number of
cycles the kernel executed as measured by the
:doc:`command processor <command-processor>`.
* - ``per_kernel``
- The total value of the measured counter or metric that occurred per
kernel invocation.
* - ``per_second``
- The total value of the measured counter or metric that occurred per
kernel invocation divided by the :ref:`kernel time <kernel-time>`,
that is, the total runtime of the kernel in seconds, as measured by the
:doc:`command processor <command-processor>`.
By default, Omniperf uses the ``per_wave`` normalization.
.. tip::
The best normalization may vary depending on your use case. For instance, a
``per_second`` normalization might be useful for FLOP or bandwidth
comparisons, while a ``per_wave`` normalization could be useful to see how many
(and what types) of instructions are used per wavefront. A ``per_kernel``
normalization can be useful to get the total aggregate values of metrics for
comparison between different configurations.
docs/conceptual/includes/terms.rst
+188
View File
@@ -0,0 +1,188 @@
.. _desc-workgroup:
.. _desc-work-item:
.. _desc-wavefront:
.. _desc-divergence:
.. _kernel-time:
.. _kernel-cycles:
.. _total-active-cu-cycles:
.. _total-cu-cycles:
.. _total-se-cycles:
.. _total-simd-cycles:
.. _total-pipe-cycles:
.. _total-l1i-cycles:
.. _total-active-l2-cycles:
.. _total-l2-cycles:
.. _total-sl1d-cycles:
.. _thread-requests:
.. list-table::
:header-rows: 1
* - Name
- Description
- Unit
* - Kernel time
- The number of seconds the accelerator was executing a kernel, from the
:doc:`command processor <command-processor>`'s (CP) start-of-kernel
timestamp (a number of cycles after the CP beings processing the packet)
to the CP's end-of-kernel timestamp (a number of cycles before the CP
stops processing the packet).
- Seconds
* - Kernel cycles
- The number of cycles the accelerator was active doing *any* work, as
measured by the :doc:`command processor <command-processor>` (CP).
- Cycles
* - Total CU cycles
- The number of cycles the accelerator was active doing *any* work
(that is, kernel cycles), multiplied by the number of
:doc:`compute units <compute-unit>` on the accelerator. A
measure of the total possible active cycles the compute units could be
doing work, useful for the normalization of metrics inside the CU.
- Cycles
* - Total active CU cycles
- The number of cycles a CU on the accelerator was active doing *any*
work, summed over all :doc:`compute units <compute-unit>` on the
accelerator.
- Cycles
* - Total SIMD cycles
- The number of cycles the accelerator was active doing *any* work (that
is, kernel cycles), multiplied by the number of
:doc:`SIMDs <compute-unit>` on the accelerator. A measure of the
total possible active cycles the SIMDs could be doing work, useful for
the normalization of metrics inside the CU.
- Cycles
* - Total L2 cycles
- The number of cycles the accelerator was active doing *any* work (that
is, kernel cycles), multiplied by the number of :doc:`L2 <l2-cache>`
channels on the accelerator. A measure of the total possible active
cycles the L2 channels could be doing work, useful for the normalization
of metrics inside the L2.
- Cycles
* - Total active L2 cycles
- The number of cycles a channel of the L2 cache was active doing *any*
work, summed over all :doc:`L2 <l2-cache>` channels on the accelerator.
- Cycles
* - Total sL1D cycles
- The number of cycles the accelerator was active doing *any* work (that
is, kernel cycles), multiplied by the number of
:ref:`scalar L1 data caches <desc-sl1d>` on the accelerator. A measure of
the total possible active cycles the sL1Ds could be doing work, useful
for the normalization of metrics inside the sL1D.
- Cycles
* - Total L1I cycles
- The number of cycles the accelerator was active doing *any* work (that
is, kernel cycles), multiplied by the number of
:ref:`L1 instruction caches <desc-l1i>` (L1I) on the accelerator. A
measure of the total possible active cycles the L1Is could be doing
work, useful for the normalization of metrics inside the L1I.
- Cycles
* - Total scheduler-pipe cycles
- The number of cycles the accelerator was active doing *any* work (that
is, kernel cycles), multiplied by the number of
:doc:`scheduler pipes <command-processor>` on the accelerator. A measure
of the total possible active cycles the scheduler-pipes could be doing
work, useful for the normalization of metrics inside the
:ref:`workgroup manager <desc-spi>` and
:doc:`command processor <command-processor>`.
- Cycles
* - Total shader-engine cycles
- The total number of cycles the accelerator was active doing *any* work,
multiplied by the number of :doc:`shader engines <shader-engine>` on the
accelerator. A measure of the total possible active cycles the shader
engines could be doing work, useful for the normalization of
metrics inside the :ref:`workgroup manager <desc-spi>`.
- Cycles
* - Thread-requests
- The number of unique memory addresses accessed by a single memory
instruction. On AMD Instinct accelerators, this has a maximum of 64
(that is, the size of the :ref:`wavefront <wavefront>`).
- Addresses
* - Work-item
- A single *thread*, or lane, of execution that executes in lockstep with
the rest of the work-items comprising a :ref:`wavefront <wavefront>`
of execution.
- N/A
* - Wavefront
- A group of work-items, or threads, that execute in lockstep on the
:doc:`compute unit <compute-unit>`. On AMD Instinct accelerators, the
wavefront size is always 64 work-items.
- N/A
* - Workgroup
- A group of wavefronts that execute on the same
:doc:`compute unit <compute-unit>`, and can cooperatively execute and
share data via the use of synchronization primitives,
:doc:`LDS <local-data-share>`, atomics, and others.
- N/A
* - Divergence
- Divergence within a wavefront occurs when not all work-items are active
when executing an instruction, that is, due to non-uniform control flow
within a wavefront. Can reduce execution efficiency by causing,
for instance, the :ref:`VALU <desc-valu>` to need to execute both
branches of a conditional with different sets of work-items active.
- N/A
docs/conceptual/l2-cache.rst
+776
View File
@@ -0,0 +1,776 @@
.. meta::
:description: Omniperf performance model: L2 cache (TCC)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, L2, cache, infinity fabric, metrics
**************
L2 cache (TCC)
**************
The L2 cache is the coherence point for current AMD Instinct™ MI-series GCN™
GPUs and CDNA™ accelerators, and is shared by all :doc:`CUs <compute-unit>`
on the device. Besides serving requests from the
:doc:`vector L1 data caches <vector-l1-cache>`, the L2 cache also is responsible
for servicing requests from the :ref:`L1 instruction caches <desc-l1i>`, the
:ref:`scalar L1 data caches <desc-sL1D>` and the
:doc:`command processor <command-processor>`. The L2 cache is composed of a
number of distinct channels (32 on MI100 and :ref:`MI2XX <mixxx-note>` series CDNA
accelerators at 256B address interleaving) which can largely operate
independently. Mapping of incoming requests to a specific L2 channel is
determined by a hashing mechanism that attempts to evenly distribute requests
across the L2 channels. Requests that miss in the L2 cache are passed out to
:ref:`Infinity Fabric™ <l2-fabric>` to be routed to the appropriate memory
location.
The L2 cache metrics reported by Omniperf are broken down into four
categories:
* :ref:`L2 Speed-of-Light <l2-sol>`
* :ref:`L2 cache accesses <l2-cache-accesses>`
* :ref:`L2-Fabric transactions <l2-fabric>`
* :ref:`L2-Fabric stalls <l2-fabric-stalls>`
.. _l2-sol:
L2 Speed-of-Light
=================
.. warning::
The theoretical maximum throughput for some metrics in this section
are currently computed with the maximum achievable clock frequency, as
reported by ``rocminfo``, for an accelerator. This may not be realistic for
all workloads.
The L2 cache’s speed-of-light table contains a few key metrics about the
performance of the L2 cache, aggregated over all the L2 channels, as a
comparison with the peak achievable values of those metrics:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Utilization
- The ratio of the
:ref:`number of cycles an L2 channel was active, summed over all L2 channels on the accelerator <total-active-l2-cycles>`
over the :ref:`total L2 cycles <total-l2-cycles>`.
- Percent
* - Bandwidth
- The number of bytes looked up in the L2 cache, as a percent of the peak
theoretical bandwidth achievable on the specific accelerator. The number
of bytes is calculated as the number of cache lines requested multiplied
by the cache line size. This value does not consider partial requests, so
e.g., if only a single value is requested in a cache line, the data
movement will still be counted as a full cache line.
- Percent
* - Hit Rate
- The ratio of the number of L2 cache line requests that hit in the L2
cache over the total number of incoming cache line requests to the L2
cache.
- Percent
* - L2-Fabric Read BW
- The number of bytes read by the L2 over the
:ref:`Infinity Fabric interface <l2-fabric>` per unit time.
- GB/s
* - L2-Fabric Write and Atomic BW
- The number of bytes sent by the L2 over the
:ref:`Infinity Fabric interface <l2-fabric>` by write and atomic
operations per unit time.
- GB/s
.. note::
The L2 cache on AMD Instinct MI CDNA accelerators uses a "hit-on-miss"
approach to reporting cache hits. That is, if while satisfying a miss,
another request comes in that would hit on the same pending cache line, the
subsequent request will be counted as a 'hit'. Therefore, it is also
important to consider the latency metric in the :ref:`L2-Fabric <l2-fabric>`
section when evaluating the L2 hit rate.
.. _l2-cache-accesses:
L2 cache accesses
=================
This section details the incoming requests to the L2 cache from the
:doc:`vL1D <vector-l1-cache>` and other clients -- for instance, the
:ref:`sL1D <desc-sL1D>` and :ref:`L1I <desc-l1i>` caches.
.. list-table::
:header-rows: 1
:widths: 13 70 17
* - Metric
- Description
- Unit
* - Bandwidth
- The number of bytes looked up in the L2 cache, per
:ref:`normalization unit <normalization-units>`. The number of bytes is
calculated as the number of cache lines requested multiplied by the cache
line size. This value does not consider partial requests, so for example,
if only a single value is requested in a cache line, the data movement
will still be counted as a full cache line.
- Bytes per :ref:`normalization unit <normalization-units>`.
* - Requests
- The total number of incoming requests to the L2 from all clients for all
request types, per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Read Requests
- The total number of read requests to the L2 from all clients.
- Requests per :ref:`normalization unit <normalization-units>`
* - Write Requests
- The total number of write requests to the L2 from all clients.
- Requests per :ref:`normalization unit <normalization-units>`
* - Atomic Requests
- The total number of atomic requests (with and without return) to the L2
from all clients.
- Requests per :ref:`normalization unit <normalization-units>`
* - Streaming Requests
- The total number of incoming requests to the L2 that are marked as
*streaming*. The exact meaning of this may differ depending on the
targeted accelerator, however on an :ref:`MI2XX <mixxx-note>` this
corresponds to
`non-temporal load or stores <https://clang.llvm.org/docs/LanguageExtensions.html#non-temporal-load-store-builtins>`_.
The L2 cache attempts to evict *streaming* requests before normal
requests when the L2 is at capacity.
- Requests per :ref:`normalization unit <normalization-units>`
* - Probe Requests
- The number of coherence probe requests made to the L2 cache from outside
the accelerator. On an :ref:`MI2XX <mixxx-note>`, probe requests may be
generated by, for example, writes to
:ref:`fine-grained device <memory-type>` memory or by writes to
:ref:`coarse-grained <memory-type>` device memory.
- Requests per :ref:`normalization unit <normalization-units>`
* - Hit Rate
- The ratio of the number of L2 cache line requests that hit in the L2
cache over the total number of incoming cache line requests to the L2
cache.
- Percent
* - Hits
- The total number of requests to the L2 from all clients that hit in the
cache. As noted in the :ref:`Speed-of-Light <l2-sol>` section, this
includes hit-on-miss requests.
- Requests per :ref:`normalization unit <normalization-units>`
* - Misses
- The total number of requests to the L2 from all clients that miss in the
cache. As noted in the :ref:`Speed-of-Light <l2-sol>` section, these do
not include hit-on-miss requests.
- Requests per :ref:`normalization unit <normalization-units>`
* - Writebacks
- The total number of L2 cache lines written back to memory for any reason.
Write-backs may occur due to user code (such as HIP kernel calls to
``__threadfence_system`` or atomic built-ins) by the
:doc:`command processor <command-processor>`'s memory acquire/release
fences, or for other internal hardware reasons.
- Cache lines per :ref:`normalization unit <normalization-units>`
* - Writebacks (Internal)
- The total number of L2 cache lines written back to memory for internal
hardware reasons, per :ref:`normalization unit <normalization-units>`.
- Cache lines per :ref:`normalization unit <normalization-units>`.
* - Writebacks (vL1D Req)
- The total number of L2 cache lines written back to memory due to requests
initiated by the :doc:`vL1D cache <vector-l1-cache>`, per
:ref:`normalization unit <normalization-units>`.
- Cache lines per :ref:`normalization unit <normalization-units>`.
* - Evictions (Normal)
- The total number of L2 cache lines evicted from the cache due to capacity
limits, per :ref:`normalization unit <normalization-units>`.
- Cache lines per :ref:`normalization unit <normalization-units>`.
* - Evictions (vL1D Req)
- The total number of L2 cache lines evicted from the cache due to
invalidation requests initiated by the
:doc:`vL1D cache <vector-l1-cache>`, per
:ref:`normalization unit <normalization-units>`.
- Cache lines per :ref:`normalization unit <normalization-units>`.
* - Non-hardware-Coherent Requests
- The total number of requests to the L2 to Not-hardware-Coherent (NC)
memory allocations, per :ref:`normalization unit <normalization-units>`.
See the :ref:`memory-type` for more information.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Uncached Requests
- The total number of requests to the L2 that go to Uncached (UC) memory
allocations. See the :ref:`memory-type` for more information.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Coherently Cached Requests
- The total number of requests to the L2 that go to Coherently Cacheable (CC)
memory allocations. See the :ref:`memory-type` for more information.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Read/Write Coherent Requests
- The total number of requests to the L2 that go to Read-Write coherent memory
(RW) allocations. See the :ref:`memory-type` for more information.
- Requests per :ref:`normalization unit <normalization-units>`.
.. note::
All requests to the L2 are for a single cache line's worth of data. The size
of a cache line may vary depending on the accelerator, however on an AMD
Instinct CDNA2 :ref:`MI2XX <mixxx-note>` accelerator, it is 128B, while on
an MI100, it is 64B.
.. _l2-fabric:
L2-Fabric transactions
======================
Requests/data that miss in the L2 must be routed to memory in order to
service them. The backing memory for a request may be local to this
accelerator (i.e., in the local high-bandwidth memory), in a remote
accelerator’s memory, or even in the CPU’s memory. Infinity Fabric
is responsible for routing these memory requests/data to the correct
location and returning any fetched data to the L2 cache. The
:ref:`l2-request-flow` describes the flow of these requests through
Infinity Fabric in more detail, as described by Omniperf metrics,
while :ref:`l2-request-metrics` give detailed definitions of
individual metrics.
.. _l2-request-flow:
Request flow
------------
The following is a diagram that illustrates how L2↔Fabric requests are reported
by Omniperf:
.. figure:: ../data/performance-model/fabric.png
:align: center
:alt: L2-Fabric transaction flow on AMD Instinct MI-series accelerators
:width: 800
L2↔Fabric transaction flow on AMD Instinct MI-series accelerators.
Requests from the L2 Cache are broken down into two major categories, read
requests and write requests (at this granularity, atomic requests are treated
as writes).
From there, these requests can additionally subdivided in a number of ways.
First, these requests may be sent across Infinity Fabric as different
transaction sizes, 32B or 64B on current CDNA accelerators.
.. note::
On current CDNA accelerators, the 32B read request path is expected to be
unused and so is disconnected in the flow diagram.
In addition, the read and write requests can be further categorized as:
* Uncached read/write requests, for instance: for access to
:ref:`fine-grained memory <memory-type>`
* Atomic requests, for instance: for atomic updates to
:ref:`fine-grained memory <memory-type>`
* HBM read/write requests OR remote read/write requests, for instance: for
requests to the accelerator’s local HBM OR requests to a remote accelerator’s
HBM or the CPU’s DRAM
These classifications are not necessarily *exclusive*. For example, a
write request can be classified as an atomic request to the
accelerator’s local HBM, and an uncached write request. The request-flow
diagram marks *exclusive* classifications as a splitting of the flow,
while *non-exclusive* requests do not split the flow line. For example,
a request is either a 32B Write Request OR a 64B Write request, as the
flow splits at this point:
.. figure:: ../data/performance-model/split.*
:align: center
:alt: Splitting request flow
:width: 800
Splitting request flow
However, continuing along, the same request might be an atomic request and an
uncached write request, as reflected by a non-split flow:
.. figure:: ../data/performance-model/nosplit.*
:align: center
:alt: Non-splitting request flow
:width: 800
Non-splitting request flow
Finally, we note that :ref:`uncached <memory-type>` read requests (e.g., to
:ref:`fine-grained memory <memory-type>`) are handled specially on CDNA
accelerators, as indicated in the request flow diagram. These are
expected to be counted as a 64B Read Request, and *if* they are requests
to uncached memory (denoted by the dashed line), they will also be
counted as *two* uncached read requests (that is, the request is split):
.. figure:: ../data/performance-model/uncached.*
:align: center
:alt: Uncached read-request splitting
:width: 800
Uncached read-request splitting.
.. _l2-request-metrics:
Metrics
-------
The following metrics are reported for the L2-Fabric interface:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - L2-Fabric Read Bandwidth
- The total number of bytes read by the L2 cache from Infinity Fabric per
:ref:`normalization unit <normalization-units>`.
- Bytes per :ref:`normalization unit <normalization-units>`.
* - HBM Read Traffic
- The percent of read requests generated by the L2 cache that are routed to
the accelerator's local high-bandwidth memory (HBM). This breakdown does
not consider the *size* of the request (meaning that 32B and 64B requests
are both counted as a single request), so this metric only *approximates*
the percent of the L2-Fabric Read bandwidth directed to the local HBM.
- Percent
* - Remote Read Traffic
- The percent of read requests generated by the L2 cache that are routed to
any memory location other than the accelerator's local high-bandwidth
memory (HBM) -- for example, the CPU's DRAM or a remote accelerator's
HBM. This breakdown does not consider the *size* of the request (meaning
that 32B and 64B requests are both counted as a single request), so this
metric only *approximates* the percent of the L2-Fabric Read bandwidth
directed to a remote location.
- Percent
* - Uncached Read Traffic
- The percent of read requests generated by the L2 cache that are reading
from an :ref:`uncached memory allocation <memory-type>`. Note, as
described in the :ref:`request flow <l2-request-flow>` section, a single
64B read request is typically counted as two uncached read requests. So,
it is possible for the Uncached Read Traffic to reach up to 200% of the
total number of read requests. This breakdown does not consider the
*size* of the request (i.e., 32B and 64B requests are both counted as a
single request), so this metric only *approximates* the percent of the
L2-Fabric read bandwidth directed to an uncached memory location.
- Percent
* - L2-Fabric Write and Atomic Bandwidth
- The total number of bytes written by the L2 over Infinity Fabric by write
and atomic operations per
:ref:`normalization unit <normalization-units>`. Note that on current
CDNA accelerators, such as the :ref:`MI2XX <mixxx-note>`, requests are
only considered *atomic* by Infinity Fabric if they are targeted at
non-write-cacheable memory, for example,
:ref:`fine-grained memory <memory-type>` allocations or
:ref:`uncached memory <memory-type>` allocations on the
MI2XX.
- Bytes per :ref:`normalization unit <normalization-units>`.
* - HBM Write and Atomic Traffic
- The percent of write and atomic requests generated by the L2 cache that
are routed to the accelerator's local high-bandwidth memory (HBM). This
breakdown does not consider the *size* of the request (meaning that 32B
and 64B requests are both counted as a single request), so this metric
only *approximates* the percent of the L2-Fabric Write and Atomic
bandwidth directed to the local HBM. Note that on current CDNA
accelerators, such as the :ref:`MI2XX <mixxx-note>`, requests are only
considered *atomic* by Infinity Fabric if they are targeted at
:ref:`fine-grained memory <memory-type>` allocations or
:ref:`uncached memory <memory-type>` allocations.
- Percent
* - Remote Write and Atomic Traffic
- The percent of read requests generated by the L2 cache that are routed to
any memory location other than the accelerator's local high-bandwidth
memory (HBM) -- for example, the CPU's DRAM or a remote accelerator's
HBM. This breakdown does not consider the *size* of the request (meaning
that 32B and 64B requests are both counted as a single request), so this
metric only *approximates* the percent of the L2-Fabric Read bandwidth
directed to a remote location. Note that on current CDNA
accelerators, such as the :ref:`MI2XX <mixxx-note>`, requests are only
considered *atomic* by Infinity Fabric if they are targeted at
:ref:`fine-grained memory <memory-type>` allocations or
:ref:`uncached memory <memory-type>` allocations.
- Percent
* - Atomic Traffic
- The percent of write requests generated by the L2 cache that are atomic
requests to *any* memory location. This breakdown does not consider the
*size* of the request (meaning that 32B and 64B requests are both counted
as a single request), so this metric only *approximates* the percent of
the L2-Fabric Read bandwidth directed to a remote location. Note that on
current CDNA accelerators, such as the :ref:`MI2XX <mixxx-note>`,
requests are only considered *atomic* by Infinity Fabric if they are
targeted at :ref:`fine-grained memory <memory-type>` allocations or
:ref:`uncached memory <memory-type>` allocations.
- Percent
* - Uncached Write and Atomic Traffic
- The percent of write and atomic requests generated by the L2 cache that
are targeting :ref:`uncached memory allocations <memory-type>`. This
breakdown does not consider the *size* of the request (meaning that 32B
and 64B requests are both counted as a single request), so this metric
only *approximates* the percent of the L2-Fabric read bandwidth directed
to uncached memory allocations.
- Percent
* - Read Latency
- The time-averaged number of cycles read requests spent in Infinity Fabric
before data was returned to the L2.
- Cycles
* - Write Latency
- The time-averaged number of cycles write requests spent in Infinity
Fabric before a completion acknowledgement was returned to the L2.
- Cycles
* - Atomic Latency
- The time-averaged number of cycles atomic requests spent in Infinity
Fabric before a completion acknowledgement (atomic without return value)
or data (atomic with return value) was returned to the L2.
- Cycles
* - Read Stall
- The ratio of the total number of cycles the L2-Fabric interface was
stalled on a read request to any destination (local HBM, remote PCIe®
connected accelerator or CPU, or remote Infinity Fabric connected
accelerator [#inf]_ or CPU) over the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Write Stall
- The ratio of the total number of cycles the L2-Fabric interface was
stalled on a write or atomic request to any destination (local HBM,
remote accelerator or CPU, PCIe connected accelerator or CPU, or remote
Infinity Fabric connected accelerator [#inf]_ or CPU) over the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
.. _l2-detailed-metrics:
Detailed transaction metrics
----------------------------
The following metrics are available in the detailed L2-Fabric
transaction breakdown table:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - 32B Read Requests
- The total number of L2 requests to Infinity Fabric to read 32B of data
from any memory location, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail. Typically unused on CDNA
accelerators.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Uncached Read Requests
- The total number of L2 requests to Infinity Fabric to read
:ref:`uncached data <memory-type>` from any memory location, per
:ref:`normalization unit <normalization-units>`. 64B requests for
uncached data are counted as two 32B uncached data requests. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - 64B Read Requests
- The total number of L2 requests to Infinity Fabric to read 64B of data
from any memory location, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - HBM Read Requests
- The total number of L2 requests to Infinity Fabric to read 32B or 64B of
data from the accelerator's local HBM, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Remote Read Requests
- The total number of L2 requests to Infinity Fabric to read 32B or 64B of
data from any source other than the accelerator's local HBM, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - 32B Write and Atomic Requests
- The total number of L2 requests to Infinity Fabric to write or atomically
update 32B of data to any memory location, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Uncached Write and Atomic Requests
- The total number of L2 requests to Infinity Fabric to write or atomically
update 32B or 64B of :ref:`uncached data <memory-type>`, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - 64B Write and Atomic Requests
- The total number of L2 requests to Infinity Fabric to write or atomically
update 64B of data in any memory location, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - HBM Write and Atomic Requests
- The total number of L2 requests to Infinity Fabric to write or atomically
update 32B or 64B of data in the accelerator's local HBM, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Remote Write and Atomic Requests
- The total number of L2 requests to Infinity Fabric to write or atomically
update 32B or 64B of data in any memory location other than the
accelerator's local HBM, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Atomic Requests
- The total number of L2 requests to Infinity Fabric to atomically update
32B or 64B of data in any memory location, per
:ref:`normalization unit <normalization-units>`. See
:ref:`l2-request-flow` for more detail. Note that on current CDNA
accelerators, such as the :ref:`MI2XX <mixxx-note>`, requests are only
considered *atomic* by Infinity Fabric if they are targeted at
non-write-cacheable memory, such as
:ref:`fine-grained memory <memory-type>` allocations or
:ref:`uncached memory <memory-type>` allocations on the MI2XX.
- Requests per :ref:`normalization unit <normalization-units>`.
.. _l2-fabric-stalls:
L2-Fabric interface stalls
==========================
When the interface between the L2 cache and Infinity Fabric becomes backed up by
requests, it may stall, preventing the L2 from issuing additional requests to
Infinity Fabric until prior requests complete. This section gives a breakdown of
what types of requests in a kernel caused a stall (like read versus write), and
to which locations -- for instance, to the accelerator’s local memory, or to
remote accelerators or CPUs.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Read - PCIe Stall
- The number of cycles the L2-Fabric interface was stalled on read requests
to remote PCIe connected accelerators [#inf]_ or CPUs as a percent of the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Read - Infinity Fabric Stall
- The number of cycles the L2-Fabric interface was stalled on read requests
to remote Infinity Fabric connected accelerators [#inf]_ or CPUs as a
percent of the :ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Read - HBM Stall
- The number of cycles the L2-Fabric interface was stalled on read requests
to the accelerator's local HBM as a percent of the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Write - PCIe Stall
- The number of cycles the L2-Fabric interface was stalled on write or
atomic requests to remote PCIe connected accelerators [#inf]_ or CPUs as
a percent of the :ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Write - Infinity Fabric Stall
- The number of cycles the L2-Fabric interface was stalled on write or
atomic requests to remote Infinity Fabric connected accelerators [#inf]_
or CPUs as a percent of the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Write - HBM Stall
- The number of cycles the L2-Fabric interface was stalled on write or
atomic requests to accelerator's local HBM as a percent of the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
* - Write - Credit Starvation
- The number of cycles the L2-Fabric interface was stalled on write or
atomic requests to any memory location because too many write/atomic
requests were currently in flight, as a percent of the
:ref:`total active L2 cycles <total-active-l2-cycles>`.
- Percent
.. warning::
On current CDNA accelerators and GCN GPUs, these L2↔Fabric stalls can be undercounted in some circumstances.
.. rubric:: Footnotes
.. [#inf] In addition to being used for on-accelerator data-traffic, AMD
`Infinity Fabric <https://www.amd.com/en/technologies/infinity-architecture>`_
technology can be used to connect multiple accelerators to achieve advanced
peer-to-peer connectivity and enhanced bandwidths over traditional PCIe
connections. Some AMD Instinct MI-series accelerators like the MI250X
`feature coherent CPU↔accelerator connections built using AMD Infinity Fabric <https://www.amd.com/system/files/documents/amd-cdna2-white-paper.pdf>`_.
.. rubric:: Disclaimer
PCIe® is a registered trademark of PCI-SIG Corporation.
docs/conceptual/local-data-share.rst
+183
View File
@@ -0,0 +1,183 @@
.. meta::
:description: Omniperf performance model: Local data share (LDS)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, local, data, share, LDS
**********************
Local data share (LDS)
**********************
.. _lds-sol:
LDS Speed-of-Light
==================
.. warning::
The theoretical maximum throughput for some metrics in this section are
currently computed with the maximum achievable clock frequency, as reported
by ``rocminfo``, for an accelerator. This may not be realistic for all
workloads.
The :ref:`LDS <desc-lds>` speed-of-light chart shows a number of key metrics for
the LDS as a comparison with the peak achievable values of those metrics.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Utilization
- Indicates what percent of the kernel's duration the :ref:`LDS <desc-lds>`
was actively executing instructions (including, but not limited to, load,
store, atomic and HIP's ``__shfl`` operations). Calculated as the ratio
of the total number of cycles LDS was active over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - Access Rate
- Indicates the percentage of SIMDs in the :ref:`VALU <desc-valu>` [#lds-workload]_
actively issuing LDS instructions, averaged over the lifetime of the
kernel. Calculated as the ratio of the total number of cycles spent by
the :ref:`scheduler <desc-scheduler>` issuing :ref:`LDS <desc-lds>`
instructions over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - Theoretical Bandwidth (% of Peak)
- Indicates the maximum amount of bytes that *could* have been loaded from,
stored to, or atomically updated in the LDS in this kernel, as a percent
of the peak LDS bandwidth achievable. See the
:ref:`LDS bandwidth example <lds-bandwidth>` for more detail.
- Percent
* - Bank Conflict Rate
- Indicates the percentage of active LDS cycles that were spent servicing
bank conflicts. Calculated as the ratio of LDS cycles spent servicing
bank conflicts over the number of LDS cycles that would have been
required to move the same amount of data in an uncontended access. [#lds-bank-conflict]_
- Percent
.. rubric:: Footnotes
.. [#lds-workload] Here we assume the typical case where the workload evenly distributes
LDS operations over all SIMDs in a CU (that is, waves on different SIMDs are
executing similar code). For highly unbalanced workloads, where e.g., one
SIMD pair in the CU does not issue LDS instructions at all, this metric is
better interpreted as the percentage of SIMDs issuing LDS instructions on
:ref:`SIMD pairs <desc-lds>` that are actively using the LDS, averaged over
the lifetime of the kernel.
.. [#lds-bank-conflict] The maximum value of the bank conflict rate is less than 100%
(specifically: 96.875%), as the first cycle in the
:ref:`LDS scheduler <desc-lds>` is never considered contended.
.. _lds-stats:
Statistics
==========
The LDS statistics panel gives a more detailed view of the hardware:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - LDS Instructions
- The total number of LDS instructions (including, but not limited to,
read/write/atomics and HIP's ``__shfl`` instructions) executed per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Theoretical Bandwidth
- Indicates the maximum amount of bytes that could have been loaded from,
stored to, or atomically updated in the LDS per
:ref:`normalization unit <normalization-units>`. Does *not* take into
account the execution mask of the wavefront when the instruction was
executed. See the
:ref:`LDS bandwidth example <lds-bandwidth>` for more detail.
- Bytes per :ref:`normalization unit <normalization-units>`
* - LDS Latency
- The average number of round-trip cycles (i.e., from issue to data-return
/ acknowledgment) required for an LDS instruction to complete.
- Cycles
* - Bank Conflicts/Access
- The ratio of the number of cycles spent in the
:ref:`LDS scheduler <desc-lds>` due to bank conflicts (as determined by
the conflict resolution hardware) to the base number of cycles that would
be spent in the LDS scheduler in a completely uncontended case. This is
the unnormalized form of the Bank Conflict Rate.
- Conflicts/Access
* - Index Accesses
- The total number of cycles spent in the :ref:`LDS scheduler <desc-lds>`
over all operations per :ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Atomic Return Cycles
- The total number of cycles spent on LDS atomics with return per
:ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Bank Conflicts
- The total number of cycles spent in the :ref:`LDS scheduler <desc-lds>`
due to bank conflicts (as determined by the conflict resolution hardware)
per :ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Address Conflicts
- The total number of cycles spent in the :ref:`LDS scheduler <desc-lds>`
due to address conflicts (as determined by the conflict resolution
hardware) per :ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Unaligned Stall
- The total number of cycles spent in the :ref:`LDS scheduler <desc-lds>`
due to stalls from non-dword aligned addresses per
:ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Memory Violations
- The total number of out-of-bounds accesses made to the LDS, per
:ref:`normalization unit <normalization-units>`. This is unused and
expected to be zero in most configurations for modern CDNA™ accelerators.
- Accesses per :ref:`normalization unit <normalization-units>`
docs/conceptual/performance-model.rst
+49
View File
@@ -0,0 +1,49 @@
.. meta::
:description: Omniperf performance model
:keywords: Omniperf, ROCm, performance, model, profiler, tool, Instinct,
accelerator, AMD
*****************
Performance model
*****************
Omniperf makes available an extensive list of metrics to better understand
achieved application performance on AMD Instinct™ MI-series accelerators
including Graphics Core Next™ (GCN) GPUs like the AMD Instinct MI50, CDNA™
accelerators like the MI100, and CDNA2 accelerators such as the MI250X, MI250,
and MI210.
To best use profiling data, it's important to understand the role of various
hardware blocks of AMD Instinct accelerators. This section describes each
hardware block on the accelerator as interacted with by a software developer to
give a deeper understanding of the metrics reported by profiling data. Refer to
:doc:`/tutorial/profiling-by-example` for more practical examples and details on how
to use Omniperf to optimize your code.
.. _mixxx-note:
.. note::
In this chapter, **MI2XX** refers to any of the CDNA2 architecture-based AMD
Instinct MI250X, MI250, and MI210 accelerators interchangeably in cases
where the exact product at hand is not relevant.
For a comparison of AMD Instinct accelerator specifications, refer to
:doc:`Hardware specifications <rocm:reference/gpu-arch-specs>`. For product
details, see the :prod-page:`MI250X <mi200/mi250x>`,
:prod-page:`MI250 <mi200/mi250>`, and :prod-page:`MI210 <mi200/mi210>`
product pages.
In this chapter, the AMD Instinct performance model used by Omniperf is divided into a handful of
key hardware blocks, each detailed in the following sections:
* :doc:`compute-unit`
* :doc:`l2-cache`
* :doc:`shader-engine`
* :doc:`command-processor`
* :doc:`system-speed-of-light`
docs/conceptual/pipeline-descriptions.rst
+299
View File
@@ -0,0 +1,299 @@
.. meta::
:description: Omniperf performance model: Shader engine (SE)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, pipeline, VALU, SALU, VMEM, SMEM, LDS, branch,
scheduler, MFMA, AGPRs
*********************
Pipeline descriptions
*********************
This section details the various execution pipelines of the
:doc:`compute unit <compute-unit>`.
.. _desc-valu:
.. _desc-vmem:
Vector arithmetic logic unit (VALU)
-----------------------------------
The vector arithmetic logic unit (VALU) executes vector instructions
over an entire wavefront, each :ref:`work-item <desc-work-item>` (or,
vector-lane) potentially operating on distinct data. The VALU of a CDNA™
accelerator or GCN™ GPU typically consists of:
* Four 16-wide SIMD processors (see :hip-training-pdf:`24` for more details).
* Four 64 or 128 KiB VGPR files (yielding a total of 256-512 KiB total
per CU), see :ref:`AGPRs <desc-agprs>` for more detail.
* An instruction buffer (per-SIMD) that contains execution slots for up
to 8 wavefronts (for 32 total wavefront slots on each CU).
* A vector memory (VMEM) unit which transfers data between VGPRs and
memory; each work-item supplies its own memory address and supplies
or receives unique data.
* CDNA accelerators, such as the MI100 and :ref:`MI2XX <mixxx-note>`, contain
additional
:amd-lab-note:`Matrix Fused Multiply-Add (MFMA) <amd-lab-notes-matrix-cores-readme>`
units.
To support branching and conditionals, each wavefront in the VALU
has a distinct execution mask which determines which work-items in the
wavefront are active for the currently executing instruction. When
executing a VALU instruction, inactive work-items (according to the
current execution mask of the wavefront) do not execute the instruction
and are treated as no-ops.
.. note::
On GCN GPUs and the CDNA MI100 accelerator, there are slots for up to 10
wavefronts in the instruction buffer, but generally occupancy is limited by
other factors to 32 waves per :doc:`compute unit <compute-unit>`.
On the CDNA2 :ref:`MI2XX <mixxx-note>` series accelerators, there are only 8
waveslots per-SIMD.
.. _desc-salu:
.. _desc-smem:
Scalar arithmetic logic unit (SALU)
-----------------------------------
The scalar arithmetic logic unit (SALU) executes instructions that are
shared between all work-items in a wavefront. This includes control flow
such as if/else conditionals, branches and looping pointer arithmetic, loading
common values, and more.
The SALU consists of:
* A scalar processor capable of various arithmetic, conditional, and
comparison (etc.) operations. See
:mi200-isa-pdf:`Chapter 5. Scalar ALU Operations <35>`
of the CDNA2 Instruction Set Architecture (ISA) Reference Guide for more
detail.
* A 12.5 KiB Scalar General Purpose Register (SGPR) file
* A scalar memory (SMEM) unit which transfers data between SGPRs and
memory
Data loaded by the SMEM can be cached in the :ref:`scalar L1 data cache <desc-sl1d>`,
and is typically only used for read-only, uniform accesses such as kernel
arguments, or HIP’s ``__constant__`` memory.
.. _desc-lds:
Local data share (LDS)
----------------------
The local data share (LDS, a.k.a., "shared memory") is fast on-CU scratchpad
that can be explicitly managed by software to effectively share data and to
coordinate between wavefronts in a workgroup.
.. figure:: ../data/performance-model/lds.*
:align: center
:alt: Performance model of the local data share (LDS) on AMD Instinct
accelerators
:width: 800
Performance model of the local data share (LDS) on AMD Instinct MI-series
accelerators.
Above is Omniperf's performance model of the LDS on CDNA accelerators (adapted
from :mantor-gcn-pdf:`20`). The SIMDs in the :ref:`VALU <desc-valu>` are
connected to the LDS in pairs (see above). Only one SIMD per pair may issue an
LDS instruction at a time, but both pairs may issue concurrently.
On CDNA accelerators, the LDS contains 32 banks and each bank is 4B wide.
The LDS is designed such that each bank can be read from, written to, or
atomically updated every cycle, for a total throughput of 128B/clock
(:gcn-crash-course:`40`).
On each of the two ports to the SIMDs, 64B can be sent in each direction per
cycle. So, a single wavefront, coming from one of the 2 SIMDs in a pair, can
only get back 64B/cycle (16 lanes per cycle). The input port is shared between
data and address and this can affect achieved bandwidth for different data
sizes. For example, a 64-wide store where each lane is sending a 4B value takes
8 cycles (50% peak bandwidth) while a 64-wide store where each lane is sending
a 16B value takes 20 cycles (80% peak bandwidth).
In addition, the LDS contains conflict-resolution hardware to detect and handle
bank conflicts. A bank conflict occurs when two (or more)
:ref:`work-items <desc-work-item>` in a :ref:`wavefront <desc-wavefront>` want
to read, write, or atomically update different addresses that map to the same
bank in the same cycle. In this case, the conflict detection hardware will
determine a new schedule such that the access is split into multiple cycles with
no conflicts in any single cycle.
When multiple work-items want to read from the same address within a bank, the
result can be efficiently broadcasted (:gcn-crash-course:`41`). Multiple
work-items writing to the same address within a bank typically results undefined
behavior in HIP and other high-level languages, as the LDS will write the value from the
last work-item as determined by the hardware scheduler (:gcn-crash-course:`41`).
This behavior may be useful in the very specific case of storing a uniform
value.
Relatedly, an address conflict is defined as occurring when two (or more)
work-items in a wavefront want to atomically update the same address on the same
cycle. As in a bank-conflict, this may cause additional cycles of work for the
LDS operation to complete.
.. _desc-branch:
Branch
------
The branch unit is responsible for executing jumps and branches to execute
control flow operations.
Note that Branch operations are not used for execution mask updates, but only
for “whole wavefront” control-flow changes.
.. _desc-scheduler:
Scheduler
---------
The scheduler is responsible for arbitration and issue of instructions for all
the wavefronts currently executing on the :doc:`CU <compute-unit>`. On every
clock cycle, the scheduler:
* Considers waves from one of the SIMD units for execution, selected in a
round-robin fashion between the SIMDs in the compute unit
* Issues up to one instruction per wavefront on the selected SIMD
* Issues up to one instruction per each of the instruction categories among the waves on the selected SIMD:
* :ref:`VALU <desc-valu>`
* :ref:`VMEM <desc-vmem>` operations
* :ref:`SALU <desc-salu>` / SMEM operations
* :ref:`LDS <desc-lds>`
* :ref:`Branch <desc-branch>` operations
This gives a maximum of five issued Instructions Per Cycle (IPC), per-SIMD,
per-CU (:hip-training-pdf:`Introduction to AMD GPU Programming with HIP <>`,
:gcn-crash-course:`The AMD GCN Architecture - A Crash Course <>`). On CDNA
accelerators with :ref:`MFMA <desc-mfma>` instructions, these are issued via the
:ref:`VALU <desc-valu>`. Some of them will execute on a separate functional unit
and typically allow other :ref:`VALU <desc-valu>` operations to execute in their
shadow (see the :ref:`MFMA <desc-mfma>` section for more detail).
.. note::
The IPC model used by Omniperf omits the following two complications for
clarity. First, CDNA accelerators contain other execution units on the CU
that are unused for compute applications. Second, so-called "internal"
instructions (see :gcn-crash-course:`29`) are not issued to a functional
unit, and can technically cause the maximum IPC to *exceed* 5 instructions
per-cycle in special (largely unrealistic) cases. The latter issue is
discussed in more detail in the
:ref:`'internal' IPC <ipc-internal-instructions>` example.
.. _desc-mfma:
Matrix fused multiply-add (MFMA)
--------------------------------
CDNA accelerators, such as the MI100 and :ref:`MI2XX <mixxx-note>`, contain
specialized hardware to accelerate matrix-matrix multiplications, also
known as Matrix Fused Multiply-Add (MFMA) operations. The exact
operation types and supported formats may vary by accelerator. Refer to the
:amd-lab-note:`AMD matrix cores <amd-lab-notes-matrix-cores-readme>`
blog post on GPUOpen for a general discussion of these hardware units.
In addition, to explore the available MFMA instructions in-depth on
various AMD accelerators (including the CDNA line), we recommend the
`AMD Matrix Instruction Calculator <https://github.com/ROCm/amd_matrix_instruction_calculator>`_:
.. code-block:: shell
:caption: Partial snapshot of the AMD Matrix Instruction Calculator Tool
$ ./matrix_calculator.py –architecture cdna2 –instruction v_mfma_f32_4x4x1f32 –detail-instruction
Architecture: CDNA2
Instruction: V_MFMA_F32_4X4X1F32
Encoding: VOP3P-MAI
VOP3P Opcode: 0x42
VOP3P-MAI Opcode: 0x2
Matrix Dimensions:
M: 4
N: 4
K: 1
blocks: 16
Execution statistics:
FLOPs: 512
Execution cycles: 8
FLOPs/CU/cycle: 256
Can co-execute with VALU: True
VALU co-execution cycles possible: 4
Register usage:
GPRs required for A: 1
GPRs required for B: 1
GPRs required for C: 4
GPRs required for D: 4
GPR alignment requirement: 8 bytes
For the purposes of Omniperf, the MFMA unit is typically treated as a separate
pipeline from the :ref:`VALU <desc-valu>`, as other VALU instructions (along
with other execution pipelines such as the :ref:`SALU <desc-salu>`) typically can be
issued during a portion of the total duration of an MFMA operation.
.. note::
The exact details of VALU and MFMA operation co-execution vary by
instruction, and can be explored in more detail via the following fields in
the
`AMD Matrix Instruction Calculator's detailed instruction information <https://github.com/ROCm/amd_matrix_instruction_calculator#example-of-querying-instruction-information>`_:
* ``Can co-execute with VALU``
* ``VALU co-execution cycles possible``
Non-pipeline resources
----------------------
In this section, we describe a few resources that are not standalone
pipelines but are important for understanding performance optimization
on CDNA accelerators.
.. _desc-barrier:
Barrier
^^^^^^^
Barriers are resources on the compute-unit of a CDNA accelerator that
are used to implement synchronization primitives (for example, HIP’s
``__syncthreads``). Barriers are allocated to any workgroup that
consists of more than a single wavefront.
.. _desc-agprs:
Accumulation vector general-purpose registers (AGPRs)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
Accumulation vector general-purpose registers, or AGPRs, are special
resources that are accessible to a subset of instructions focused on
:ref:`MFMA <desc-mfma>` operations. These registers allow the MFMA
unit to access more than the normal maximum of 256 architected
:ref:`vector general-purpose registers (VGPRs) <desc-valu>` by having up to 256
in the architected space and up to 256 in the accumulation space.
Traditional VALU instructions can only use VGPRs in the architected
space, and data can be moved to/from VGPRs↔AGPRs using specialized
instructions (``v_accvgpr_*``). These data movement instructions may be
used by the compiler to implement lower-cost register-spill/fills on
architectures with AGPRs.
AGPRs are not available on all AMD Instinct™ accelerators. GCN GPUs,
such as the AMD Instinct MI50 had a 256 KiB VGPR file. The AMD
Instinct MI100 (CDNA) has a 2x256 KiB register file, where one half
is available as general-purpose VGPRs, and the other half is for matrix
math accumulation VGPRs (AGPRs). The AMD Instinct :ref:`MI2XX <mixxx-note>`
(CDNA2) has a 512 KiB VGPR file per CU, where each wave can dynamically request
up to 256 KiB of VGPRs and an additional 256 KiB of AGPRs. For more information,
refer to `this comment <https://github.com/ROCm/ROCm/issues/1689#issuecomment-1553751913>`_.
docs/conceptual/pipeline-metrics.rst
+909
View File
@@ -0,0 +1,909 @@
.. meta::
:description: Omniperf performance model: Pipeline metrics
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, pipeline, wavefront, metrics, launch, runtime
VALU, MFMA, instruction mix, FLOPs, arithmetic, operations
****************
Pipeline metrics
****************
In this section, we describe the metrics available in Omniperf to analyze the
pipelines discussed in the :doc:`pipeline-descriptions`.
.. _wavefront:
Wavefront
=========
.. _wavefront-launch-stats:
Wavefront launch stats
----------------------
The wavefront launch stats panel gives general information about the
kernel launch:
.. list-table::
:header-rows: 1
:widths: 20 65 15
* - Metric
- Description
- Unit
* - Grid Size
- The total number of work-items (or, threads) launched as a part of
the kernel dispatch. In HIP, this is equivalent to the total grid size
multiplied by the total workgroup (or, block) size.
- :ref:`Work-items <desc-work-item>`
* - Workgroup Size
- The total number of work-items (or, threads) in each workgroup
(or, block) launched as part of the kernel dispatch. In HIP, this is
equivalent to the total block size.
- :ref:`Work-items <desc-work-item>`
* - Total Wavefronts
- The total number of wavefronts launched as part of the kernel dispatch.
On AMD Instinct™ CDNA™ accelerators and GCN™ GPUs, the wavefront size is
always 64 work-items. Thus, the total number of wavefronts should be
equivalent to the ceiling of grid size divided by 64.
- :ref:`Wavefronts <desc-wavefront>`
* - Saved Wavefronts
- The total number of wavefronts saved at a context-save. See
`cwsr_enable <https://docs.kernel.org/gpu/amdgpu/module-parameters.html?highlight=cwsr>`_.
- :ref:`Wavefronts <desc-wavefront>`
* - Restored Wavefronts
- The total number of wavefronts restored from a context-save. See
`cwsr_enable <https://docs.kernel.org/gpu/amdgpu/module-parameters.html?highlight=cwsr>`_.
- :ref:`Wavefronts <desc-wavefront>`
* - VGPRs
- The number of architected vector general-purpose registers allocated for
the kernel, see :ref:`VALU <desc-valu>`. Note: this may not exactly
match the number of VGPRs requested by the compiler due to allocation
granularity.
- :ref:`VGPRs <desc-valu>`
* - AGPRs
- The number of accumulation vector general-purpose registers allocated for
the kernel, see :ref:`AGPRs <desc-agprs>`. Note: this may not exactly
match the number of AGPRs requested by the compiler due to allocation
granularity.
- :ref:`AGPRs <desc-agprs>`
* - SGPRs
- The number of scalar general-purpose registers allocated for the kernel,
see :ref:`SALU <desc-salu>`. Note: this may not exactly match the number
of SGPRs requested by the compiler due to allocation granularity.
- :ref:`SGPRs <desc-salu>`
* - LDS Allocation
- The number of bytes of :doc:`LDS <local-data-share>` memory (or, shared
memory) allocated for this kernel. Note: This may also be larger than
what was requested at compile time due to both allocation granularity and
dynamic per-dispatch LDS allocations.
- Bytes per :ref:`workgroup <desc-workgroup>`
* - Scratch Allocation
- The number of bytes of :ref:`scratch memory <memory-spaces>` requested
per work-item for this kernel. Scratch memory is used for stack memory
on the accelerator, as well as for register spills and restores.
- Bytes per :ref:`work-item <desc-work-item>`
.. _wavefront-runtime-stats:
Wavefront runtime stats
-----------------------
The wavefront runtime statistics gives a high-level overview of the
execution of wavefronts in a kernel:
.. list-table::
:header-rows: 1
:widths: 18 65 17
* - Metric
- Description
- Unit
* - :ref:`Kernel time <kernel-time>`
- The total duration of the executed kernel. Note: this should not be
directly compared to the wavefront cycles / timings below.
- Nanoseconds
* - :ref:`Kernel cycles <kernel-cycles>`
- The total duration of the executed kernel in cycles. Note: this should
not be directly compared to the wavefront cycles / timings below.
- Cycles
* - Instructions per wavefront
- The average number of instructions (of all types) executed per wavefront.
This is averaged over all wavefronts in a kernel dispatch.
- Instructions / wavefront
* - Wave cycles
- The number of cycles a wavefront in the kernel dispatch spent resident on
a compute unit per :ref:`normalization unit <normalization-units>`. This
is averaged over all wavefronts in a kernel dispatch. Note: this should
not be directly compared to the kernel cycles above.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Dependency wait cycles
- The number of cycles a wavefront in the kernel dispatch stalled waiting
on memory of any kind (e.g., instruction fetch, vector or scalar memory,
etc.) per :ref:`normalization unit <normalization-units>`. This counter
is incremented at every cycle by *all* wavefronts on a CU stalled at a
memory operation. As such, it is most useful to get a sense of how waves
were spending their time, rather than identification of a precise limiter
because another wave could be actively executing while a wave is stalled.
The sum of this metric, Issue Wait Cycles and Active Cycles should be
equal to the total Wave Cycles metric.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Issue Wait Cycles
- The number of cycles a wavefront in the kernel dispatch was unable to
issue an instruction for any reason (e.g., execution pipe back-pressure,
arbitration loss, etc.) per
:ref:`normalization unit <normalization-units>`. This counter is
incremented at every cycle by *all* wavefronts on a CU unable to issue an
instruction. As such, it is most useful to get a sense of how waves were
spending their time, rather than identification of a precise limiter
because another wave could be actively executing while a wave is issue
stalled. The sum of this metric, Dependency Wait Cycles and Active
Cycles should be equal to the total Wave Cycles metric.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Active Cycles
- The average number of cycles a wavefront in the kernel dispatch was
actively executing instructions per
:ref:`normalization unit <normalization-units>`. This measurement is made
on a per-wavefront basis, and may include cycles that another wavefront
spent actively executing (on another execution unit, for example) or was
stalled. As such, it is most useful to get a sense of how waves were
spending their time, rather than identification of a precise limiter. The
sum of this metric, Issue Wait Cycles and Active Wait Cycles should be
equal to the total Wave Cycles metric.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Wavefront Occupancy
- The time-averaged number of wavefronts resident on the accelerator over
the lifetime of the kernel. Note: this metric may be inaccurate for
short-running kernels (less than 1ms).
- :ref:`Wavefronts <desc-wavefront>`
.. note::
As mentioned earlier, the measurement of kernel cycles and time typically
cannot be directly compared to, for example, wave cycles. This is due to two factors:
first, the kernel cycles/timings are measured using a counter that is
impacted by scheduling overhead, this is particularly noticeable for
"short-running" kernels (less than 1ms) where scheduling overhead forms a
significant portion of the overall kernel runtime. Secondly, the wave cycles
metric is incremented per-wavefront scheduled to a SIMD every cycle whereas
the kernel cycles counter is incremented only once per-cycle when *any*
wavefront is scheduled.
.. _instruction-mix:
Instruction mix
===============
The instruction mix panel shows a breakdown of the various types of instructions
executed by the user’s kernel, and which pipelines on the
:doc:`CU <compute-unit>` they were executed on. In addition, Omniperf reports
further information about the breakdown of operation types for the
:ref:`VALU <desc-valu>`, vector-memory, and :ref:`MFMA <desc-mfma>`
instructions.
.. note::
All metrics in this section count *instructions issued*, and *not* the total
number of operations executed. The values reported by these metrics will not
change regardless of the execution mask of the wavefront. Note that even if
the execution mask is identically zero (meaning that *no lanes are active*)
the instruction will still be counted, as CDNA accelerators still consider
these instructions *issued*. See
:mi200-isa-pdf:`EXECute Mask, section 3.3 of the CDNA2 ISA guide<19>` for
examples and further details.
Overall instruction mix
-----------------------
This panel shows the total number of each type of instruction issued to
the :doc:`various compute pipelines </conceptual/pipeline-descriptions>` on the
:doc:`CU </conceptual/compute-unit>`. These are:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - :ref:`VALU <desc-valu>` instructions
- The total number of vector arithmetic logic unit (VALU) operations
issued. These are the workhorses of the
:doc:`compute unit <compute-unit>`, and are used to execute a wide range of
instruction types including floating point operations, non-uniform
address calculations, transcendental operations, integer operations,
shifts, conditional evaluation, etc.
- Instructions
* - VMEM instructions
- The total number of vector memory operations issued. These include most
loads, stores and atomic operations and all accesses to
:ref:`generic, global, private and texture <memory-spaces>` memory.
- Instructions
* - :doc:`LDS <local-data-share>` instructions
- The total number of LDS (also known as shared memory) operations issued.
These include loads, stores, atomics, and HIP's ``__shfl`` operations.
- Instructions
* - :ref:`MFMA <desc-mfma>` instructions
- The total number of matrix fused multiply-add instructions issued.
- Instructions
* - :ref:`SALU <desc-salu>` instructions
- The total number of scalar arithmetic logic unit (SALU) operations
issued. Typically these are used for address calculations, literal
constants, and other operations that are *provably* uniform across a
wavefront. Although scalar memory (SMEM) operations are issued by the
SALU, they are counted separately in this section.
- Instructions
* - SMEM instructions
- The total number of scalar memory (SMEM) operations issued. These are
typically used for loading kernel arguments, base-pointers and loads
from HIP's ``__constant__`` memory.
- Instructions
* - :ref:`Branch <desc-branch>` instructions
- The total number of branch operations issued. These typically consist of
jump or branch operations and are used to implement control flow.
- Instructions
.. note::
Note, as mentioned in the :ref:`desc-branch` section: branch
operations are not used for execution mask updates, but only for "whole
wavefront" control flow changes.
.. _valu-arith-instruction-mix:
VALU arithmetic instruction mix
-------------------------------
.. warning::
Not all metrics in this section (for instance, the floating-point instruction
breakdowns) are available on CDNA accelerators older than the
:ref:`MI2XX <mixxx-note>` series.
This panel details the various types of vector instructions that were
issued to the :ref:`VALU <desc-valu>`. The metrics in this section do *not*
include :ref:`MFMA <desc-mfma>` instructions using the same precision; for
instance, the “F16-ADD” metric does not include any 16-bit floating point
additions executed as part of an MFMA instruction using the same precision.
.. list-table::
:header-rows: 1
:widths: 15 65 20
* - Metric
- Description
- Unit
* - INT32
- The total number of instructions operating on 32-bit integer operands
issued to the VALU per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - INT64
- The total number of instructions operating on 64-bit integer operands
issued to the VALU per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F16-ADD
- The total number of addition instructions operating on 16-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F16-MUL
- The total number of multiplication instructions operating on 16-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F16-FMA
- The total number of fused multiply-add instructions operating on 16-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F16-TRANS
- The total number of transcendental instructions (e.g., `sqrt`) operating
on 16-bit floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F32-ADD
- The total number of addition instructions operating on 32-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F32-MUL
- The total number of multiplication instructions operating on 32-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F32-FMA
- The total number of fused multiply-add instructions operating on 32-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F32-TRANS
- The total number of transcendental instructions (such as ``sqrt``)
operating on 32-bit floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F64-ADD
- The total number of addition instructions operating on 64-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F64-MUL
- The total number of multiplication instructions operating on 64-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F64-FMA
- The total number of fused multiply-add instructions operating on 64-bit
floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - F64-TRANS
- The total number of transcendental instructions (such as `sqrt`)
operating on 64-bit floating-point operands issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Conversion
- The total number of type conversion instructions (such as converting data
to or from F32↔F64) issued to the VALU per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
For an example of these counters in action, refer to
:ref:`valu-arith-instruction-mix-ex`.
.. _vmem-instruction-mix:
VMEM instruction mix
--------------------
This section breaks down the types of vector memory (VMEM) instructions
that were issued. Refer to the
:ref:`Instruction Counts metrics section <ta-instruction-counts>` under address
processor front end of the vL1D cache for descriptions of these VMEM
instructions.
.. _mfma-instruction-mix:
MFMA instruction mix
--------------------
.. warning::
The metrics in this section are only available on CDNA2
(:ref:`MI2XX <mixxx-note>`) accelerators and newer.
This section details the types of Matrix Fused Multiply-Add
(:ref:`MFMA <desc-mfma>`) instructions that were issued. Note that
MFMA instructions are classified by the type of input data they operate on, and
*not* the data type the result is accumulated to.
.. list-table::
:header-rows: 1
:widths: 25 60 17
* - Metric
- Description
- Unit
* - MFMA-I8 Instructions
- The total number of 8-bit integer :ref:`MFMA <desc-mfma>` instructions
issued per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - MFMA-F16 Instructions
- The total number of 16-bit floating point :ref:`MFMA <desc-mfma>`
instructions issued per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - MFMA-BF16 Instructions
- The total number of 16-bit brain floating point :ref:`MFMA <desc-mfma>`
instructions issued per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - MFMA-F32 Instructions
- The total number of 32-bit floating-point :ref:`MFMA <desc-mfma>`
instructions issued per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - MFMA-F64 Instructions
- The total number of 64-bit floating-point :ref:`MFMA <desc-mfma>`
instructions issued per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
Compute pipeline
================
.. _metrics-flop-count:
FLOP counting conventions
-------------------------
Omniperf’s conventions for VALU FLOP counting are as follows:
* Addition or multiplication: 1 operation
* Transcendentals: 1 operation
* Fused multiply-add (FMA): 2 operations
Integer operations (IOPs) do not use this convention. They are counted
as a single operation regardless of the instruction type.
.. note::
Packed operations which operate on multiple operands in the same instruction
are counted identically to the underlying instruction type. For example, the
``v_pk_add_f32`` instruction on :ref:`MI2XX <mixxx-note>`, which performs an
add operation on two pairs of aligned 32-bit floating-point operands is
counted only as a single addition -- that is, 1 operation.
As discussed in the :ref:`instruction-mix` section, the FLOP/IOP
metrics in this section do not take into account the execution mask of
the operation, and will report the same value even if the execution mask
is identically zero.
For example, a FMA instruction operating on 32-bit floating-point
operands (such as ``v_fma_f32`` on a :ref:`MI2XX <mixxx-note>` accelerator)
would be counted as 128 total FLOPs: 2 operations (due to the
instruction type) multiplied by 64 operations (because the wavefront is
composed of 64 work-items).
.. _compute-speed-of-light:
Compute Speed-of-Light
----------------------
.. warning::
The theoretical maximum throughput for some metrics in this section are
currently computed with the maximum achievable clock frequency, as reported
by ``rocminfo``, for an accelerator. This may not be realistic for all
workloads.
This section reports the number of floating-point and integer operations
executed on the :ref:`VALU <desc-valu>` and :ref:`MFMA <desc-mfma>` units in
various precisions. We note that unlike the
:ref:`VALU instruction mix <valu-arith-instruction-mix>` and
:ref:`MFMA instruction mix <mfma-instruction-mix>` sections, the metrics here
are reported as FLOPs and IOPs, that is, the total number of operations
executed.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - VALU FLOPs
- The total floating-point operations executed per second on the
:ref:`VALU <desc-valu>`. This is also presented as a percent of the peak
theoretical FLOPs achievable on the specific accelerator. Note: this does
not include any floating-point operations from :ref:`MFMA <desc-mfma>`
instructions.
- GFLOPs
* - VALU IOPs
- The total integer operations executed per second on the
:ref:`VALU <desc-valu>`. This is also presented as a percent of the peak
theoretical IOPs achievable on the specific accelerator. Note: this does
not include any integer operations from :ref:`MFMA <desc-mfma>`
instructions.
- GIOPs
* - MFMA FLOPs (BF16)
- The total number of 16-bit brain floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 16-bit
brain floating point operations from :ref:`VALU <desc-valu>`
instructions. This is also presented as a percent of the peak theoretical
BF16 MFMA operations achievable on the specific accelerator.
- GFLOPs
* - MFMA FLOPs (F16)
- The total number of 16-bit floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 16-bit
floating point operations from :ref:`VALU <desc-valu>` instructions. This
is also presented as a percent of the peak theoretical F16 MFMA
operations achievable on the specific accelerator.
- GFLOPs
* - MFMA FLOPs (F32)
- The total number of 32-bit floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 32-bit
floating point operations from :ref:`VALU <desc-valu>` instructions. This
is also presented as a percent of the peak theoretical F32 MFMA
operations achievable on the specific accelerator.
- GFLOPs
* - MFMA FLOPs (F64)
- The total number of 64-bit floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 64-bit
floating point operations from :ref:`VALU <desc-valu>` instructions. This
is also presented as a percent of the peak theoretical F64 MFMA
operations achievable on the specific accelerator.
- GFLOPs
* - MFMA IOPs (INT8)
- The total number of 8-bit integer :ref:`MFMA <desc-mfma>` operations
executed per second. Note: this does not include any 8-bit integer
operations from :ref:`VALU <desc-valu>` instructions. This is also
presented as a percent of the peak theoretical INT8 MFMA operations
achievable on the specific accelerator.
- GIOPs
.. _pipeline-stats:
Pipeline statistics
-------------------
This section reports a number of key performance characteristics of
various execution units on the :doc:`CU <compute-unit>`. Refer to
:ref:`ipc-example` for a detailed dive into these metrics, and the
:ref:`scheduler <desc-scheduler>` the for a high-level overview of execution
units and instruction issue.
.. list-table::
:header-rows: 1
:widths: 20 65 15
* - Metric
- Description
- Unit
* - IPC
- The ratio of the total number of instructions executed on the
:doc:`CU <compute-unit>` over the
:ref:`total active CU cycles <total-active-cu-cycles>`.
- Instructions per-cycle
* - IPC (Issued)
- The ratio of the total number of
(non-:ref:`internal <ipc-internal-instructions>`) instructions issued over
the number of cycles where the :ref:`scheduler <desc-scheduler>` was
actively working on issuing instructions. Refer to the
:ref:`Issued IPC <issued-ipc>` example for further detail.
- Instructions per-cycle
* - SALU utilization
- Indicates what percent of the kernel's duration the
:ref:`SALU <desc-salu>` was busy executing instructions. Computed as the
ratio of the total number of cycles spent by the
:ref:`scheduler <desc-scheduler>` issuing SALU / :ref:`SMEM <desc-smem>`
instructions over the :ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - VALU utilization
- Indicates what percent of the kernel's duration the
:ref:`VALU <desc-valu>` was busy executing instructions. Does not include
:ref:`VMEM <desc-vmem>` operations. Computed as the ratio of the total
number of cycles spent by the :ref:`scheduler <desc-scheduler>` issuing
VALU instructions over the :ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - VMEM utilization
- Indicates what percent of the kernel's duration the
:ref:`VMEM <desc-vmem>` unit was busy executing instructions, including
both global/generic and spill/scratch operations (see the
:ref:`VMEM instruction count metrics <ta-instruction-counts>` for more
detail). Does not include :ref:`VALU <desc-valu>` operations. Computed
as the ratio of the total number of cycles spent by the
:ref:`scheduler <desc-scheduler>` issuing VMEM instructions over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - Branch utilization
- Indicates what percent of the kernel's duration the
:ref:`branch <desc-branch>` unit was busy executing instructions.
Computed as the ratio of the total number of cycles spent by the
:ref:`scheduler <desc-scheduler>` issuing branch instructions over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - VALU active threads
- Indicates the average level of :ref:`divergence <desc-divergence>` within
a wavefront over the lifetime of the kernel. The number of work-items
that were active in a wavefront during execution of each
:ref:`VALU <desc-valu>` instruction, time-averaged over all VALU
instructions run on all wavefronts in the kernel.
- Work-items
* - MFMA utilization
- Indicates what percent of the kernel's duration the
:ref:`MFMA <desc-mfma>` unit was busy executing instructions. Computed as
the ratio of the total number of cycles spent by the
:ref:`MFMA <desc-salu>` was busy over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - MFMA instruction cycles
- The average duration of :ref:`MFMA <desc-mfma>` instructions in this
kernel in cycles. Computed as the ratio of the total number of cycles the
MFMA unit was busy over the total number of MFMA instructions. Compare
to, for example, the
`AMD Matrix Instruction Calculator <https://github.com/RadeonOpenCompute/amd_matrix_instruction_calculator>`_.
- Cycles per instruction
* - VMEM latency
- The average number of round-trip cycles (that is, from issue to data
return / acknowledgment) required for a VMEM instruction to complete.
- Cycles
* - SMEM latency
- The average number of round-trip cycles (that is, from issue to data
return / acknowledgment) required for a SMEM instruction to complete.
- Cycles
.. note::
The branch utilization reported in this section also includes time spent in
other instruction types (namely: ``s_endpgm``) that are *typically* a very
small percentage of the overall kernel execution. This complication is
omitted for simplicity, but may result in small amounts of branch utilization
(typically less than 1%) for otherwise branch-less kernels.
.. _arithmetic-operations:
Arithmetic operations
---------------------
This section reports the total number of floating-point and integer
operations executed in various precisions. Unlike the
:ref:`compute-speed-of-light` panel, this section reports both
:ref:`VALU <desc-valu>` and :ref:`MFMA <desc-mfma>` operations of the same precision
(e.g., F32) in the same metric. Additionally, this panel lets the user
control how the data is normalized (i.e., control the
:ref:`normalization unit <normalization-units>`), while the speed-of-light panel does
not. For more detail on how operations are counted see the
:ref:`FLOP counting convention <metrics-flop-count>` section.
.. warning::
As discussed in :ref:`instruction-mix`, the metrics in this section do not
take into account the execution mask of the operation, and will report the
same value even if EXEC is identically zero.
.. list-table::
:header-rows: 1
:widths: 18 65 17
* - Metric
- Description
- Unit
* - FLOPs (Total)
- The total number of floating-point operations executed on either the
:ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`.
- FLOP per :ref:`normalization unit <normalization-units>`
* - IOPs (Total)
- The total number of integer operations executed on either the
:ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`.
- IOP per :ref:`normalization unit <normalization-units>`
* - F16 OPs
- The total number of 16-bit floating-point operations executed on either the
:ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`.
- FLOP per :ref:`normalization unit <normalization-units>`
* - BF16 OPs
- The total number of 16-bit brain floating-point operations executed on either the
:ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`. Note: on current CDNA
accelerators, the VALU has no native BF16 instructions.
- FLOP per :ref:`normalization unit <normalization-units>`
* - F32 OPs
- The total number of 32-bit floating-point operations executed on either
the :ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`.
- FLOP per :ref:`normalization unit <normalization-units>`
* - F64 OPs
- The total number of 64-bit floating-point operations executed on either
the :ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`.
- FLOP per :ref:`normalization unit <normalization-units>`
* - INT8 OPs
- The total number of 8-bit integer operations executed on either the
:ref:`VALU <desc-valu>` or :ref:`MFMA <desc-mfma>` units, per
:ref:`normalization unit <normalization-units>`. Note: on current CDNA
accelerators, the VALU has no native INT8 instructions.
- IOPs per :ref:`normalization unit <normalization-units>`
docs/conceptual/references.rst
+26
View File
@@ -0,0 +1,26 @@
.. meta::
:description: Omniperf performance model: References
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, HIP, GCN, LLVM, docs, documentation, training
**********
References
**********
Some sections in :doc:`this chapter </conceptual/performance-model>` cite the
following publicly available documentation.
* :hip-training-pdf:`Introduction to AMD GPU Programming with HIP <>`
* :mi200-isa-pdf:`CDNA2 ISA Reference Guide <>`
* :cdna2-white-paper:`CDNA2 white paper <>`
* :hsa-runtime-pdf:`HSA Runtime Programmer's Reference Manual <>`
* :gcn-crash-course:`The AMD GCN Architecture - A Crash Course (Layla Mah) <>`
* :mantor-gcn-pdf:`AMD Radeon HD7970 with GCN Architecture <>`
* :mantor-vega10-pdf:`AMD Radeon Next Generation GPU Architecture - Vega10 <>`
* :llvm-docs:`LLVM User Guide for AMDGPU Backend <>`
docs/conceptual/shader-engine.rst
+707
View File
@@ -0,0 +1,707 @@
.. meta::
:description: Omniperf performance model: Shader engine (SE)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, shader, engine, sL1D, L1I, workgroup manager, SPI
******************
Shader engine (SE)
******************
The :doc:`compute units <compute-unit>` on a CDNA™ accelerator are grouped
together into a higher-level organizational unit called a shader engine (SE):
.. figure:: ../data/performance-model/selayout.png
:align: center
:alt: Example of CU-grouping into shader engines
:width: 800
Example of CU-grouping into shader engines on AMD Instinct MI-series
accelerators.
The number of CUs on a SE varies from chip to chip -- see for example
:hip-training-pdf:`20`. In addition, newer accelerators such as the AMD
Instinct™ MI 250X have 8 SEs per accelerator.
For the purposes of Omniperf, we consider resources that are shared between
multiple CUs on a single SE as part of the SE's metrics.
These include:
* The :ref:`scalar L1 data cache <desc-sl1d>`
* The :ref:`L1 instruction cache <desc-l1i>`
* The :ref:`workgroup manager <desc-spi>`
.. _desc-sl1d:
Scalar L1 data cache (sL1D)
===========================
The Scalar L1 Data cache (sL1D) can cache data accessed from scalar load
instructions (and scalar store instructions on architectures where they exist)
from wavefronts in the :doc:`CUs <compute-unit>`. The sL1D is shared between
multiple CUs (:gcn-crash-course:`36`) -- the exact number of CUs depends on the
architecture in question (3 CUs in GCN™ GPUs and MI100, 2 CUs in
:ref:`MI2XX <mixxx-note>`) -- and is backed by the :doc:`L2 cache <l2-cache>`.
In typical usage, the data in the sL1D is comprised of:
* Kernel arguments, such as pointers,
`non-populated <https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-sgpr-register-set-up-order-table>`_
grid and block dimensions, and others
* HIP's ``__constant__`` memory, when accessed in a provably uniform manner
[#uniform-access]_
* Other memory, when accessed in a provably uniform manner, *and* the backing
memory is provably constant [#uniform-access]_
.. _desc-sl1d-sol:
Scalar L1D Speed-of-Light
-------------------------
.. warning::
The theoretical maximum throughput for some metrics in this section are
currently computed with the maximum achievable clock frequency, as reported
by ``rocminfo``, for an accelerator. This may not be realistic for all
workloads.
The Scalar L1D speed-of-light chart shows some key metrics of the sL1D
cache as a comparison with the peak achievable values of those metrics:
.. list-table::
:header-rows: 1
:widths: 20 65 15
* - Metric
- Description
- Unit
* - Bandwidth
- The number of bytes looked up in the sL1D cache, as a percent of the peak
theoretical bandwidth. Calculated as the ratio of sL1D requests over the
:ref:`total sL1D cycles <total-sl1d-cycles>`.
- Percent
* - Cache Hit Rate
- The percent of sL1D requests that hit [#sl1d-cache]_ on a previously
loaded line in the cache. Calculated as the ratio of the number of sL1D
requests that hit over the number of all sL1D requests.
- Percent
* - sL1D-L2 BW
- The number of bytes requested by the sL1D from the L2 cache, as a percent
of the peak theoretical sL1D → L2 cache bandwidth. Calculated as the
ratio of the total number of requests from the sL1D to the L2 cache over
the :ref:`total sL1D-L2 interface cycles <total-sl1d-cycles>`.
- Percent
.. _desc-sl1d-stats:
Scalar L1D cache accesses
-------------------------
This panel gives more detail on the types of accesses made to the sL1D,
and the hit/miss statistics.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Requests
- The total number of requests, of any size or type, made to the sL1D per
:ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Hits
- The total number of sL1D requests that hit on a previously loaded cache
line, per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Misses - Non Duplicated
- The total number of sL1D requests that missed on a cache line that *was
not* already pending due to another request, per
:ref:`normalization unit <normalization-units>`. See :ref:`desc-sl1d-sol`
for more detail.
- Requests per :ref:`normalization unit <normalization-units>`
* - Misses - Duplicated
- The total number of sL1D requests that missed on a cache line that *was*
already pending due to another request, per
:ref:`normalization unit <normalization-units>`. See
:ref:`desc-sl1d-sol` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`
* - Cache Hit Rate
- Indicates the percent of sL1D requests that hit on a previously loaded
line the cache. The ratio of the number of sL1D requests that hit
[#sl1d-cache]_ over the number of all sL1D requests.
- Percent
* - Read Requests (Total)
- The total number of sL1D read requests of any size, per
:ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Atomic Requests
- The total number of sL1D atomic requests of any size, per
:ref:`normalization unit <normalization-units>`. Typically unused on CDNA
accelerators.
- Requests per :ref:`normalization unit <normalization-units>`
* - Read Requests (1 DWord)
- The total number of sL1D read requests made for a single dword of data
(4B), per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Read Requests (2 DWord)
- The total number of sL1D read requests made for a two dwords of data
(8B), per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Read Requests (4 DWord)
- The total number of sL1D read requests made for a four dwords of data
(16B), per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Read Requests (8 DWord)
- The total number of sL1D read requests made for a eight dwords of data
(32B), per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Read Requests (16 DWord)
- The total number of sL1D read requests made for a sixteen dwords of data
(64B), per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
.. _desc-sl1d-l2-interface:
sL1D ↔ L2 Interface
-------------------
This panel gives more detail on the data requested across the
sL1D↔
:doc:`L2 <l2-cache>` interface.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - sL1D-L2 BW
- The total number of bytes read from, written to, or atomically updated
across the sL1D↔:doc:`L2 <l2-cache>` interface, per
:ref:`normalization unit <normalization-units>`. Note that sL1D writes
and atomics are typically unused on current CDNA accelerators, so in the
majority of cases this can be interpreted as an sL1D→L2 read bandwidth.
- Bytes per :ref:`normalization unit <normalization-units>`
* - Read Requests
- The total number of read requests from sL1D to the :doc:`L2 <l2-cache>`,
per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Write Requests
- The total number of write requests from sL1D to the :doc:`L2 <l2-cache>`,
per :ref:`normalization unit <normalization-units>`. Typically unused on
current CDNA accelerators.
- Requests per :ref:`normalization unit <normalization-units>`
* - Atomic Requests
- The total number of atomic requests from sL1D to the
:doc:`L2 <l2-cache>`, per
:ref:`normalization unit <normalization-units>`. Typically unused on
current CDNA accelerators.
- Requests per :ref:`normalization unit <normalization-units>`
* - Stall Cycles
- The total number of cycles the sL1D↔
:doc:`L2 <l2-cache>` interface was stalled, per
:ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
.. rubric:: Footnotes
.. [#uniform-access] The scalar data cache is used when the compiler emits
scalar loads to access data. This requires that the data be *provably*
uniformly accesses (that is, the compiler can verify that all work-items in a
wavefront access the same data), *and* that the data can be proven to be
read-only (for instance, HIP's ``__constant__`` memory, or properly
``__restrict__``\ed pointers to avoid write-aliasing). Access of
``__constant__`` memory for example is not guaranteed to go through the sL1D
if the wavefront loads a non-uniform value.
.. [#sl1d-cache] Unlike the :doc:`vL1D <vector-l1-cache>` and
:doc:`L2 <l2-cache>` caches, the sL1D cache on AMD Instinct MI-series CDNA
accelerators does *not* use the "hit-on-miss" approach to reporting cache
hits. That is, if while satisfying a miss, another request comes in that
would hit on the same pending cache line, the subsequent request will be
counted as a *duplicated miss*.
.. _desc-l1i:
L1 Instruction Cache (L1I)
==========================
As with the :ref:`sL1D <desc-sL1D>`, the L1 Instruction (L1I) cache is shared
between multiple CUs on a shader-engine, where the precise number of CUs
sharing a L1I depends on the architecture in question (:gcn-crash-course:`36`)
and is backed by the :doc:`L2 cache <l2-cache>`. Unlike the sL1D, the
instruction cache is read-only.
.. _desc-l1i-sol:
L1I Speed-of-Light
------------------
.. warning::
The theoretical maximum throughput for some metrics in this section are
currently computed with the maximum achievable clock frequency, as reported
by ``rocminfo``, for an accelerator. This may not be realistic for all
workloads.
The L1 Instruction Cache speed-of-light chart shows some key metrics of
the L1I cache as a comparison with the peak achievable values of those
metrics:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Bandwidth
- The number of bytes looked up in the L1I cache, as a percent of the peak
theoretical bandwidth. Calculated as the ratio of L1I requests over the
:ref:`total L1I cycles <total-l1i-cycles>`.
- Percent
* - Cache Hit Rate
- The percent of L1I requests that hit on a previously loaded line the
cache. Calculated as the ratio of the number of L1I requests that hit
[#l1i-cache]_ over the number of all L1I requests.
- Percent
* - L1I-L2 BW
- The percent of the peak theoretical L1I → L2 cache request bandwidth
achieved. Calculated as the ratio of the total number of requests from
the L1I to the L2 cache over the
:ref:`total L1I-L2 interface cycles <total-l1i-cycles>`.
- Percent
* - Instruction Fetch Latency
- The average number of cycles spent to fetch instructions to a
:doc:`CU <compute-unit>`.
- Cycles
.. _desc-l1i-stats:
L1I cache accesses
------------------
This panel gives more detail on the hit/miss statistics of the L1I:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Requests
- The total number of requests made to the L1I per
:ref:`normalization-unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Hits
- The total number of L1I requests that hit on a previously loaded cache
line, per :ref:`normalization-unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Misses - Non Duplicated
- The total number of L1I requests that missed on a cache line that
*were not* already pending due to another request, per
:ref:`normalization-unit <normalization-units>`. See note in
:ref:`desc-l1i-sol` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`.
* - Misses - Duplicated
- The total number of L1I requests that missed on a cache line that *were*
already pending due to another request, per
:ref:`normalization-unit <normalization-units>`. See note in
:ref:`desc-l1i-sol` for more detail.
- Requests per :ref:`normalization unit <normalization-units>`
* - Cache Hit Rate
- The percent of L1I requests that hit [#l1i-cache]_ on a previously loaded
line the cache. Calculated as the ratio of the number of L1I requests
that hit over the number of all L1I requests.
- Percent
L1I - L2 interface
------------------
This panel gives more detail on the data requested across the
L1I-:doc:`L2 <l2-cache>` interface.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - L1I-L2 BW
- The total number of bytes read across the L1I-:doc:`L2 <l2-cache>`
interface, per :ref:`normalization unit <normalization-units>`.
- Bytes per :ref:`normalization unit <normalization-units>`
.. rubric:: Footnotes
.. [#l1i-cache] Unlike the :doc:`vL1D <vector-l1-cache>` and
:doc:`L2 <l2-cache>` caches, the L1I cache on AMD Instinct MI-series CDNA
accelerators does *not* use the "hit-on-miss" approach to reporting cache
hits. That is, if while satisfying a miss, another request comes in that
would hit on the same pending cache line, the subsequent request will be
counted as a *duplicated miss*.
.. _desc-spi:
Workgroup manager (SPI)
=======================
The workgroup manager (SPI) is the bridge between the
:doc:`command processor <command-processor>` and the
:doc:`compute units <compute-unit>`. After the command processor processes a
kernel dispatch, it will then pass the dispatch off to the workgroup manager,
which then schedules :ref:`workgroups <desc-workgroup>` onto the compute units.
As workgroups complete execution and resources become available, the
workgroup manager will schedule new workgroups onto compute units. The workgroup
manager’s metrics therefore are focused on reporting the following:
* Utilizations of various parts of the accelerator that the workgroup
manager interacts with (and the workgroup manager itself)
* How many workgroups were dispatched, their size, and how many
resources they used
* Percent of scheduler opportunities (cycles) where workgroups failed
to dispatch, and
* Percent of scheduler opportunities (cycles) where workgroups failed
to dispatch due to lack of a specific resource on the CUs (for instance, too
many VGPRs allocated)
This gives you an idea of why the workgroup manager couldn’t schedule more
wavefronts onto the device, and is most useful for workloads that you suspect to
be limited by scheduling or launch rate.
As discussed in :doc:`Command processor <command-processor>`, the command
processor on AMD Instinct MI-series architectures contains four hardware
scheduler-pipes, each with eight software threads (:mantor-vega10-pdf:`19`). Each
scheduler-pipe can issue a kernel dispatch to the workgroup manager to schedule
concurrently. Therefore, some workgroup manager metrics are presented relative
to the utilization of these scheduler-pipes (for instance, whether all four are
issuing concurrently).
.. note::
Current versions of the profiling libraries underlying Omniperf attempt to
serialize concurrent kernels running on the accelerator, as the performance
counters on the device are global (that is, shared between concurrent
kernels). This means that these scheduler-pipe utilization metrics are
expected to reach (for example) a maximum of one pipe active -- only 25%.
Workgroup manager utilizations
------------------------------
This section describes the utilization of the workgroup manager, and the
hardware components it interacts with.
.. list-table::
:header-rows: 1
:widths: 20 65 15
* - Metric
- Description
- Unit
* - Accelerator utilization
- The percent of cycles in the kernel where the accelerator was actively
doing any work.
- Percent
* - Scheduler-pipe utilization
- The percent of :ref:`total scheduler-pipe cycles <total-pipe-cycles>` in
the kernel where the scheduler-pipes were actively doing any work. Note:
this value is expected to range between 0% and 25%. See :ref:`desc-spi`.
- Percent
* - Workgroup manager utilization
- The percent of cycles in the kernel where the workgroup manager was
actively doing any work.
- Percent
* - Shader engine utilization
- The percent of :ref:`total shader engine cycles <total-se-cycles>` in the
kernel where any CU in a shader-engine was actively doing any work,
normalized over all shader-engines. Low values (e.g., << 100%) indicate
that the accelerator was not fully saturated by the kernel, or a
potential load-imbalance issue.
- Percent
* - SIMD utilization
- The percent of :ref:`total SIMD cycles <total-simd-cycles>` in the kernel
where any :ref:`SIMD <desc-valu>` on a CU was actively doing any work,
summed over all CUs. Low values (less than 100%) indicate that the
accelerator was not fully saturated by the kernel, or a potential
load-imbalance issue.
- Percent
* - Dispatched workgroups
- The total number of workgroups forming this kernel launch.
- Workgroups
* - Dispatched wavefronts
- The total number of wavefronts, summed over all workgroups, forming this
kernel launch.
- Wavefronts
* - VGPR writes
- The average number of cycles spent initializing :ref:`VGPRs <desc-valu>`
at wave creation.
- Cycles/wave
* - SGPR Writes
- The average number of cycles spent initializing :ref:`SGPRs <desc-salu>`
at wave creation.
- Cycles/wave
Resource allocation
-------------------
This panel gives more detail on how workgroups and wavefronts were scheduled
onto compute units, and what occupancy limiters they hit -- if any. When
analyzing these metrics, you should also take into account their
achieved occupancy -- such as
:ref:`wavefront occupancy <wavefront-runtime-stats>`. A kernel may be occupancy
limited by LDS usage, for example, but may still achieve high occupancy levels
such that improving occupancy further may not improve performance. See
:ref:`occupancy-example` for details.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Not-scheduled rate (Workgroup Manager)
- The percent of :ref:`total scheduler-pipe cycles <total-pipe-cycles>` in
the kernel where a workgroup could not be scheduled to a
:doc:`CU <compute-unit>` due to a bottleneck within the workgroup manager
rather than a lack of a CU or :ref:`SIMD <desc-valu>` with sufficient
resources. Note: this value is expected to range between 0-25%. See note
in :ref:`workgroup manager <desc-spi>` description.
- Percent
* - Not-scheduled rate (Scheduler-Pipe)
- The percent of :ref:`total scheduler-pipe cycles <total-pipe-cycles>` in
the kernel where a workgroup could not be scheduled to a
:doc:`CU <compute-unit>` due to a bottleneck within the scheduler-pipes
rather than a lack of a CU or :ref:`SIMD <desc-valu>` with sufficient
resources. Note: this value is expected to range between 0-25%, see note
in :ref:`workgroup manager <desc-spi>` description.
- Percent
* - Scheduler-Pipe Stall Rate
- The percent of :ref:`total scheduler-pipe cycles <total-pipe-cycles>` in
the kernel where a workgroup could not be scheduled to a
:doc:`CU <compute-unit>` due to occupancy limitations (like a lack of a
CU or :ref:`SIMD <desc-valu>` with sufficient resources). Note: this
value is expected to range between 0-25%, see note in
:ref:`workgroup manager <desc-spi>` description.
- Percent
* - Scratch Stall Rate
- The percent of :ref:`total shader-engine cycles <total-se-cycles>` in the
kernel where a workgroup could not be scheduled to a
:doc:`CU <compute-unit>` due to lack of
:ref:`private (a.k.a., scratch) memory <memory-type>` slots. While this
can reach up to 100%, note that the actual occupancy limitations on a
kernel using private memory are typically quite small (for example, less
than 1% of the total number of waves that can be scheduled to an
accelerator).
- Percent
* - Insufficient SIMD Waveslots
- The percent of :ref:`total SIMD cycles <total-simd-cycles>` in the kernel
where a workgroup could not be scheduled to a :ref:`SIMD <desc-valu>`
due to lack of available :ref:`waveslots <desc-valu>`.
- Percent
* - Insufficient SIMD VGPRs
- The percent of :ref:`total SIMD cycles <total-simd-cycles>` in the kernel
where a workgroup could not be scheduled to a :ref:`SIMD <desc-valu>`
due to lack of available :ref:`VGPRs <desc-valu>`.
- Percent
* - Insufficient SIMD SGPRs
- The percent of :ref:`total SIMD cycles <total-simd-cycles>` in the kernel
where a workgroup could not be scheduled to a :ref:`SIMD <desc-valu>`
due to lack of available :ref:`SGPRs <desc-salu>`.
- Percent
* - Insufficient CU LDS
- The percent of :ref:`total CU cycles <total-cu-cycles>` in the kernel
where a workgroup could not be scheduled to a :doc:`CU <compute-unit>`
due to lack of available :doc:`LDS <local-data-share>`.
- Percent
* - Insufficient CU Barriers
- The percent of :ref:`total CU cycles <total-cu-cycles>` in the kernel
where a workgroup could not be scheduled to a :doc:`CU <compute-unit>`
due to lack of available :ref:`barriers <desc-barrier>`.
- Percent
* - Reached CU Workgroup Limit
- The percent of :ref:`total CU cycles <total-cu-cycles>` in the kernel
where a workgroup could not be scheduled to a :doc:`CU <compute-unit>`
due to limits within the workgroup manager. This is expected to be
always be zero on CDNA2 or newer accelerators (and small for previous
accelerators).
- Percent
* - Reached CU Wavefront Limit
- The percent of :ref:`total CU cycles <total-cu-cycles>` in the kernel
where a wavefront could not be scheduled to a :doc:`CU <compute-unit>`
due to limits within the workgroup manager. This is expected to be
always be zero on CDNA2 or newer accelerators (and small for previous
accelerators).
- Percent
docs/conceptual/system-speed-of-light.rst
+318
View File
@@ -0,0 +1,318 @@
.. meta::
:description: Omniperf performance model: System Speed-of-Light
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, AMD, system, speed of light
*********************
System Speed-of-Light
*********************
System Speed-of-Light summarizes some of the key metrics from various sections
of Omniperf’s profiling report.
.. warning::
The theoretical maximum throughput for some metrics in this section are
currently computed with the maximum achievable clock frequency, as reported
by ``rocminfo``, for an accelerator. This may not be realistic for
all workloads.
Also, not all metrics -- such as FLOP counters -- are available on all AMD
Instinct™ MI-series accelerators. For more detail on how operations are
counted, see the :ref:`metrics-flop-count` section.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - :ref:`VALU <desc-valu>` FLOPs
- The total floating-point operations executed per second on the
:ref:`VALU <desc-valu>`. This is also presented as a percent of the peak
theoretical FLOPs achievable on the specific accelerator. Note: this does
not include any floating-point operations from :ref:`MFMA <desc-mfma>`
instructions.
- GFLOPs
* - :ref:`VALU <desc-valu>` IOPs
- The total integer operations executed per second on the
:ref:`VALU <desc-valu>`. This is also presented as a percent of the peak
theoretical IOPs achievable on the specific accelerator. Note: this does
not include any integer operations from :ref:`MFMA <desc-mfma>`
instructions.
- GIOPs
* - :ref:`MFMA <desc-mfma>` FLOPs (BF16)
- The total number of 16-bit brain floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 16-bit
brain floating point operations from :ref:`VALU <desc-valu>`
instructions. This is also presented as a percent of the peak theoretical
BF16 MFMA operations achievable on the specific accelerator.
- GFLOPs
* - :ref:`MFMA <desc-mfma>` FLOPs (F16)
- The total number of 16-bit floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 16-bit
floating point operations from :ref:`VALU <desc-valu>` instructions. This
is also presented as a percent of the peak theoretical F16 MFMA
operations achievable on the specific accelerator.
- GFLOPs
* - :ref:`MFMA <desc-mfma>` FLOPs (F32)
- The total number of 32-bit floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 32-bit
floating point operations from :ref:`VALU <desc-valu>` instructions. This
is also presented as a percent of the peak theoretical F32 MFMA
operations achievable on the specific accelerator.
- GFLOPs
* - :ref:`MFMA <desc-mfma>` FLOPs (F64)
- The total number of 64-bit floating point :ref:`MFMA <desc-mfma>`
operations executed per second. Note: this does not include any 64-bit
floating point operations from :ref:`VALU <desc-valu>` instructions. This
is also presented as a percent of the peak theoretical F64 MFMA
operations achievable on the specific accelerator.
- GFLOPs
* - :ref:`MFMA <desc-mfma>` IOPs (INT8)
- The total number of 8-bit integer :ref:`MFMA <desc-mfma>` operations
executed per second. Note: this does not include any 8-bit integer
operations from :ref:`VALU <desc-valu>` instructions. This is also
presented as a percent of the peak theoretical INT8 MFMA operations
achievable on the specific accelerator.
- GIOPs
* - :ref:`SALU <desc-salu>` utilization
- Indicates what percent of the kernel's duration the
:ref:`SALU <desc-salu>` was busy executing instructions. Computed as the
ratio of the total number of cycles spent by the
:ref:`scheduler <desc-scheduler>` issuing :ref:`SALU <desc-salu>` or
:ref:`SMEM <desc-salu>` instructions over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - :ref:`VALU <desc-valu>` utilization
- Indicates what percent of the kernel's duration the
:ref:`VALU <desc-valu>` was busy executing instructions. Does not include
:ref:`VMEM <desc-vmem>` operations. Computed as the ratio of the total
number of cycles spent by the :ref:`scheduler <desc-scheduler>` issuing
:ref:`VALU <desc-valu>` instructions over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - :ref:`MFMA <desc-mfma>` utilization
- Indicates what percent of the kernel's duration the
:ref:`MFMA <desc-mfma>` unit was busy executing instructions. Computed as
the ratio of the total number of cycles the MFMA was busy over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - :ref:`VMEM <desc-valu>` utilization
- Indicates what percent of the kernel's duration the
:ref:`VMEM <desc-valu>` unit was busy executing instructions, including
both global/generic and spill/scratch operations (see the
:ref:`VMEM instruction count metrics <ta-instruction-counts>`) for more
detail). Does not include :ref:`VALU <desc-valu>` operations. Computed as
the ratio of the total number of cycles spent by the
:ref:`scheduler <desc-scheduler>` issuing VMEM instructions over the
:ref:`total CU cycles <total-cu-cycles>`.
- Percent
* - :ref:`Branch <desc-branch>` utilization
- Indicates what percent of the kernel's duration the
:ref:`branch <desc-branch>` unit was busy executing instructions.
Computed as the ratio of the total number of cycles spent by the
:ref:`scheduler <desc-scheduler>` issuing :ref:`branch <desc-branch>`
instructions over the :ref:`total CU cycles <total-cu-cycles>`
- Percent
* - :ref:`VALU <desc-valu>` active threads
- Indicates the average level of :ref:`divergence <desc-divergence>` within
a wavefront over the lifetime of the kernel. The number of work-items
that were active in a wavefront during execution of each
:ref:`VALU <desc-valu>` instruction, time-averaged over all VALU
instructions run on all wavefronts in the kernel.
- Work-items
* - IPC
- The ratio of the total number of instructions executed on the
:doc:`CU <compute-unit>` over the
:ref:`total active CU cycles <total-active-cu-cycles>`. This is also
presented as a percent of the peak theoretical bandwidth achievable on
the specific accelerator.
- Instructions per-cycle
* - Wavefront occupancy
- The time-averaged number of wavefronts resident on the accelerator over
the lifetime of the kernel. Note: this metric may be inaccurate for
short-running kernels (less than 1ms). This is also presented as a
percent of the peak theoretical occupancy achievable on the specific
accelerator.
- Wavefronts
* - :doc:`LDS <local-data-share>` theoretical bandwidth
- Indicates the maximum amount of bytes that could have been loaded from,
stored to, or atomically updated in the LDS per unit time (see
:ref:`LDS Bandwidth <lds-bandwidth>` example for more detail). This is
also presented as a percent of the peak theoretical F64 MFMA operations
achievable on the specific accelerator.
- GB/s
* - :doc:`LDS <local-data-share>` bank conflicts/access
- The ratio of the number of cycles spent in the
:doc:`LDS scheduler <local-data-share>` due to bank conflicts (as
determined by the conflict resolution hardware) to the base number of
cycles that would be spent in the LDS scheduler in a completely
uncontended case. This is also presented in normalized form (i.e., the
Bank Conflict Rate).
- Conflicts/Access
* - :doc:`vL1D <vector-l1-cache>` cache hit rate
- The ratio of the number of vL1D cache line requests that hit in vL1D
cache over the total number of cache line requests to the
:ref:`vL1D cache RAM <desc-tc>`.
- Percent
* - :doc:`vL1D <vector-l1-cache>` cache bandwidth
- The number of bytes looked up in the vL1D cache as a result of
:ref:`VMEM <desc-vmem>` instructions per unit time. The number of bytes
is calculated as the number of cache lines requested multiplied by the
cache line size. This value does not consider partial requests, so e.g.,
if only a single value is requested in a cache line, the data movement
will still be counted as a full cache line. This is also presented as a
percent of the peak theoretical bandwidth achievable on the specific
accelerator.
- GB/s
* - :doc:`L2 <l2-cache>` cache hit rate
- The ratio of the number of L2 cache line requests that hit in the L2
cache over the total number of incoming cache line requests to the L2
cache.
- Percent
* - :doc:`L2 <l2-cache>` cache bandwidth
- The number of bytes looked up in the L2 cache per unit time. The number
of bytes is calculated as the number of cache lines requested multiplied
by the cache line size. This value does not consider partial requests, so
e.g., if only a single value is requested in a cache line, the data
movement will still be counted as a full cache line. This is also
presented as a percent of the peak theoretical bandwidth achievable on
the specific accelerator.
- GB/s
* - :doc:`L2 <l2-cache>`-fabric read BW
- The number of bytes read by the L2 over the
:ref:`Infinity Fabric™ interface <l2-fabric>` per unit time. This is also
presented as a percent of the peak theoretical bandwidth achievable on
the specific accelerator.
- GB/s
* - :doc:`L2 <l2-cache>`-fabric write and atomic BW
- The number of bytes sent by the L2 over the
:ref:`Infinity Fabric interface <l2-fabric>` by write and atomic
operations per unit time. This is also presented as a percent of the peak
theoretical bandwidth achievable on the specific accelerator.
- GB/s
* - :doc:`L2 <l2-cache>`-fabric read latency
- The time-averaged number of cycles read requests spent in Infinity Fabric
before data was returned to the L2.
- Cycles
* - :doc:`L2 <l2-cache>`-fabric write latency
- The time-averaged number of cycles write requests spent in Infinity
Fabric before a completion acknowledgement was returned to the L2.
- Cycles
* - :ref:`sL1D <desc-sl1d>` cache hit rate
- The percent of sL1D requests that hit on a previously loaded line the
cache. Calculated as the ratio of the number of sL1D requests that hit
over the number of all sL1D requests.
- Percent
* - :ref:`sL1D <desc-sl1d>` bandwidth
- The number of bytes looked up in the sL1D cache per unit time. This is
also presented as a percent of the peak theoretical bandwidth achievable
on the specific accelerator.
- GB/s
* - :ref:`L1I <desc-l1i>` bandwidth
- The number of bytes looked up in the L1I cache per unit time. This is
also presented as a percent of the peak theoretical bandwidth achievable
on the specific accelerator.
- GB/s
* - :ref:`L1I <desc-l1i>` cache hit rate
- The percent of L1I requests that hit on a previously loaded line the
cache. Calculated as the ratio of the number of L1I requests that hit
over the number of all L1I requests.
- Percent
* - :ref:`L1I <desc-l1i>` fetch latency
- The average number of cycles spent to fetch instructions to a
:doc:`CU <compute-unit>`.
- Cycles
docs/conceptual/vector-l1-cache.rst
+767
View File
@@ -0,0 +1,767 @@
.. meta::
:description: Omniperf performance model: Vector L1 cache (vL1D)
:keywords: Omniperf, ROCm, profiler, tool, Instinct, accelerator, AMD, vector, l1, cache, vl1d
**********************
Vector L1 cache (vL1D)
**********************
The vector L1 data (vL1D) cache is local to each
:doc:`compute unit <compute-unit>` on the accelerator, and handles vector memory
operations issued by a wavefront. The vL1D cache consists of several components:
* An address processing unit, also known as the
:ref:`texture addresser <desc-ta>` which receives commands (instructions) and
write/atomic data from the :doc:`compute unit <compute-unit>`, and coalesces
them into fewer requests for the cache to process.
* An address translation unit, also known as the
:ref:`L1 Unified Translation Cache (UTCL1) <desc-utcl1>`, that translates
requests from virtual to physical addresses for lookup in the cache. The
translation unit has an L1 translation lookaside buffer (L1TLB) to reduce the
cost of repeated translations.
* A Tag RAM that looks up whether a requested cache line is already
present in the :ref:`cache <desc-tc>`.
* The result of the Tag RAM lookup is placed in the L1 cache controller
for routing to the correct location; for instance, the
:ref:`L2 Memory Interface <vl1d-l2-transaction-detail>` for misses or the
:ref:`cache RAM <desc-tc>` for hits.
* The cache RAM, also known as the :ref:`texture cache (TC) <desc-tc>`, stores
requested data for potential reuse. Data returned from the
:doc:`L2 cache <l2-cache>` is placed into the cache RAM before going down the
:ref:`data-return path <desc-td>`.
* A backend data processing unit, also known as the
:ref:`texture data (TD) <desc-td>` that routes data back to the requesting
:doc:`compute unit <compute-unit>`.
Together, this complex is known as the vL1D, or Texture Cache per Pipe
(TCP). A simplified diagram of the vL1D is presented below:
.. figure:: ../data/performance-model/l1perf_model.png
:align: center
:alt: Performance model of the vL1D Cache on AMD Instinct
:width: 800
Performance model of the vL1D Cache on AMD Instinct MI-series accelerators.
.. _vl1d-sol:
vL1D Speed-of-Light
===================
.. warning::
The theoretical maximum throughput for some metrics in this section are
currently computed with the maximum achievable clock frequency, as reported
by ``rocminfo``, for an accelerator. This may not be realistic for all
workloads.
The vL1D’s speed-of-light chart shows several key metrics for the vL1D
as a comparison with the peak achievable values of those metrics.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Hit Rate
- The ratio of the number of vL1D cache line requests that hit [#vl1d-hit]_
in vL1D cache over the total number of cache line requests to the
:ref:`vL1D Cache RAM <desc-tc>`.
- Percent
* - Bandwidth
- The number of bytes looked up in the vL1D cache as a result of
:ref:`VMEM <desc-vmem>` instructions, as a percent of the peak
theoretical bandwidth achievable on the specific accelerator. The number
of bytes is calculated as the number of cache lines requested multiplied
by the cache line size. This value does not consider partial requests, so
for instance, if only a single value is requested in a cache line, the
data movement will still be counted as a full cache line.
- Percent
* - Utilization
- Indicates how busy the :ref:`vL1D Cache RAM <desc-tc>` was during the
kernel execution. The number of cycles where the vL1D Cache RAM is
actively processing any request divided by the number of cycles where the
vL1D is active [#vl1d-activity]_.
- Percent
* - Coalescing
- Indicates how well memory instructions were coalesced by the
:ref:`address processing unit <desc-ta>`, ranging from uncoalesced (25%)
to fully coalesced (100%). Calculated as the average number of
:ref:`thread-requests <thread-requests>` generated per instruction
divided by the ideal number of thread-requests per instruction.
- Percent
.. _desc-ta:
Address processing unit or Texture Addresser (TA)
=================================================
The :doc:`vL1D <vector-l1-cache>`’s address processing unit receives vector
memory instructions (commands) along with write/atomic data from a
:doc:`compute unit <compute-unit>` and is responsible for coalescing these into
requests for lookup in the :ref:`vL1D RAM <desc-tc>`. The address processor
passes information about the commands (coalescing state, destination SIMD,
etc.) to the :ref:`data processing unit <desc-td>` for use after the requested
data has been retrieved.
Omniperf reports several metrics to indicate performance bottlenecks in
the address processing unit, which are broken down into a few
categories:
- :ref:`ta-busy-stall`
- :ref:`ta-instruction-counts`
- :ref:`ta-spill-stack`
.. _ta-busy-stall:
Busy / stall metrics
--------------------
When executing vector memory instructions, the compute unit must send an
address (and in the case of writes/atomics, data) to the address
processing unit. When the front-end cannot accept any more addresses, it
must backpressure the wave-issue logic for the VMEM pipe and prevent the
issue of further vector memory instructions.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Busy
- Percent of the :ref:`total CU cycles <total-cu-cycles>` the address
processor was busy
- Percent
* - Address Stall
- Percent of the :ref:`total CU cycles <total-cu-cycles>` the address
processor was stalled from sending address requests further into the vL1D
pipeline
- Percent
* - Data Stall
- Percent of the :ref:`total CU cycles <total-cu-cycles>` the address
processor was stalled from sending write/atomic data further into the
vL1D pipeline
- Percent
* - Data-Processor → Address Stall
- Percent of :ref:`total CU cycles <total-cu-cycles>` the address processor
was stalled waiting to send command data to the
:ref:`data processor <desc-td>`
- Percent
.. _ta-instruction-counts:
Instruction counts
------------------
The address processor also counts instruction types to give the user
information on what sorts of memory instructions were executed by the
kernel. These are broken down into a few major categories:
.. list-table::
:header-rows: 1
* - Memory type
- Usage
- Description
* - Global
- Global memory
- Global memory can be seen by all threads from a process. This includes
the local accelerator's DRAM, remote accelerator's DRAM, and the host's
DRAM.
* - Generic
- Dynamic address spaces
- Generic memory, or "flat" memory, is used when the compiler cannot
statically prove that a pointer is to memory in one or the other address
spaces. The pointer could dynamically point into global, local, constant,
or private memory.
* - Private Memory
- Register spills / Stack memory
- Private memory, or "scratch" memory, is only visible to a particular
:ref:`work-item <desc-work-item>` in a particular
:ref:`workgroup <desc-workgroup>`. On AMD Instinct™ MI-series
accelerators, private memory is used to implement both register spills
and stack memory accesses.
The address processor counts these instruction types as follows:
.. list-table::
:header-rows: 1
* - Type
- Description
- Unit
* - Global/Generic
- The total number of global & generic memory instructions executed on all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Global/Generic Read
- The total number of global & generic memory read instructions executed on
all :doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Global/Generic Write
- The total number of global & generic memory write instructions executed
on all :doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Global/Generic Atomic
- The total number of global & generic memory atomic (with and without
return) instructions executed on all :doc:`compute units <compute-unit>`
on the accelerator, per :ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Spill/Stack
- The total number of spill/stack memory instructions executed on all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Spill/Stack Read
- The total number of spill/stack memory read instructions executed on all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Spill/Stack Write
- The total number of spill/stack memory write instructions executed on all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`.
- Instruction per :ref:`normalization unit <normalization-units>`
* - Spill/Stack Atomic
- The total number of spill/stack memory atomic (with and without return)
instructions executed on all :doc:`compute units <compute-unit>` on the
accelerator, per :ref:`normalization unit <normalization-units>`.
Typically unused as these memory operations are typically used to
implement thread-local storage.
- Instructions per :ref:`normalization unit <normalization-units>`
.. note::
The above is a simplified model specifically for the HIP programming language
that does not consider inline assembly usage, constant memory usage or
texture memory.
These categories correspond to:
* Global/Generic: global and flat memory operations, that are used for global
and generic memory access.
* Spill/Stack: buffer instructions which are used on the MI50, MI100, and
:ref:`MI2XX <mixxx-note>` accelerators for register spills / stack memory.
These concepts are described in more detail in the :ref:`memory-spaces`,
while generic memory access is explored in the
:ref:`generic memory benchmark <flat-memory-ex>` section.
.. _ta-spill-stack:
Spill / stack metrics
---------------------
Finally, the address processing unit contains a separate coalescing
stage for spill/stack memory, and thus reports:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Spill/Stack Total Cycles
- The number of cycles the address processing unit spent working on
spill/stack instructions, per
:ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Spill/Stack Coalesced Read Cycles
- The number of cycles the address processing unit spent working on
coalesced spill/stack read instructions, per
:ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
* - Spill/Stack Coalesced Write Cycles
- The number of cycles the address processing unit spent working on
coalesced spill/stack write instructions, per
:ref:`normalization unit <normalization-units>`.
- Cycles per :ref:`normalization unit <normalization-units>`
.. _desc-utcl1:
L1 Unified Translation Cache (UTCL1)
====================================
After a vector memory instruction has been processed/coalesced by the
address processing unit of the vL1D, it must be translated from a
virtual to physical address. This process is handled by the L1 Unified
Translation Cache (UTCL1). This cache contains a L1 Translation
Lookaside Buffer (TLB) which stores recently translated addresses to
reduce the cost of subsequent re-translations.
Omniperf reports the following L1 TLB metrics:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Requests
- The number of translation requests made to the UTCL1 per
:ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Hits
- The number of translation requests that hit in the UTCL1, and could be
reused, per :ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Hit Ratio
- The ratio of the number of translation requests that hit in the UTCL1
divided by the total number of translation requests made to the UTCL1.
- Percent
* - Translation Misses
- The total number of translation requests that missed in the UTCL1 due to
translation not being present in the cache, per
:ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - Permission Misses
- The total number of translation requests that missed in the UTCL1 due to
a permission error, per :ref:`normalization unit <normalization-units>`.
This is unused and expected to be zero in most configurations for modern
CDNA™ accelerators.
- Requests per :ref:`normalization unit <normalization-units>`
.. note::
On current CDNA accelerators, such as the :ref:`MI2XX <mixxx-note>`, the
UTCL1 does *not* count hit-on-miss requests.
.. _desc-tc:
Vector L1 Cache RAM or Texture Cache (TC)
=========================================
After coalescing in the :ref:`address processing unit <desc-ta>` of the v1LD,
and address translation in the :ref:`L1 TLB <desc-utcl1>` the request proceeds
to the Cache RAM stage of the pipeline. Incoming requests are looked up
in the cache RAMs using parts of the physical address as a tag. Hits
will be returned through the :ref:`data-return path <desc-td>`, while misses
will routed out to the :doc:`L2 Cache <l2-cache>` for servicing.
The metrics tracked by the vL1D RAM include:
- :ref:`Stall metrics <vl1d-cache-stall-metrics>`
- :ref:`Cache access metrics <vl1d-cache-access-metrics>`
- :ref:`vL1D-L2 transaction detail metrics <vl1d-l2-transaction-detail>`
.. _vl1d-cache-stall-metrics:
vL1D cache stall metrics
------------------------
The vL1D also reports where it is stalled in the pipeline, which may
indicate performance limiters of the cache. A stall in the pipeline may
result in backpressuring earlier parts of the pipeline, e.g., a stall on
L2 requests may backpressure the wave-issue logic of the :ref:`VMEM <desc-vmem>`
pipe and prevent it from issuing more vector memory instructions until
the vL1D’s outstanding requests are completed.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Stalled on L2 Data
- The ratio of the number of cycles where the vL1D is stalled waiting for
requested data to return from the :doc:`L2 cache <l2-cache>` divided by
the number of cycles where the vL1D is active [#vl1d-activity]_.
- Percent
* - Stalled on L2 Requests
- The ratio of the number of cycles where the vL1D is stalled waiting to
issue a request for data to the :doc:`L2 cache <l2-cache>` divided by the
number of cycles where the vL1D is active [#vl1d-activity]_.
- Percent
* - Tag RAM Stall (Read/Write/Atomic)
- The ratio of the number of cycles where the vL1D is stalled due to
Read/Write/Atomic requests with conflicting tags being looked up
concurrently, divided by the number of cycles where the
vL1D is active [#vl1d-activity]_.
- Percent
.. _vl1d-cache-access-metrics:
vL1D cache access metrics
-------------------------
The vL1D cache access metrics broadly indicate the type of requests
incoming from the :ref:`cache front-end <desc-ta>`, the number of requests that
were serviced by the vL1D, and the number & type of outgoing requests to
the :doc:`L2 cache <l2-cache>`. In addition, this section includes the
approximate latencies of accesses to the cache itself, along with
latencies of read/write memory operations to the :doc:`L2 cache <l2-cache>`.
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Total Requests
- The total number of incoming requests from the
:ref:`address processing unit <desc-ta>` after coalescing.
- Requests
* - Total read/write/atomic requests
- The total number of incoming read/write/atomic requests from the
:ref:`address processing unit <desc-ta>` after coalescing per
:ref:`normalization unit <normalization-units>`
- Requests per :ref:`normalization unit <normalization-units>`
* - Cache Bandwidth
- The number of bytes looked up in the vL1D cache as a result of
:ref:`VMEM <desc-vmem>` instructions per
:ref:`normalization unit <normalization-units>`. The number of bytes is
calculated as the number of cache lines requested multiplied by the cache
line size. This value does not consider partial requests, so for
instance, if only a single value is requested in a cache line, the data
movement will still be counted as a full cache line.
- Bytes per :ref:`normalization unit <normalization-units>`
* - Cache Hit Rate [#vl1d-hit]_
- The ratio of the number of vL1D cache line requests that hit in vL1D
cache over the total number of cache line requests to the
:ref:`vL1D Cache RAM <desc-tc>`.
- Percent
* - Cache Accesses
- The total number of cache line lookups in the vL1D.
- Cache lines
* - Cache Hits [#vl1d-hit]_
- The number of cache accesses minus the number of outgoing requests to the
:doc:`L2 cache <l2-cache>`, that is, the number of cache line requests
serviced by the :ref:`vL1D Cache RAM <desc-tc>` per
:ref:`normalization unit <normalization-units>`.
- Cache lines per :ref:`normalization unit <normalization-units>`
* - Invalidations
- The number of times the vL1D was issued a write-back invalidate command
during the kernel's execution per
:ref:`normalization unit <normalization-units>`. This may be triggered
by, for instance, the ``buffer_wbinvl1`` instruction.
- Invalidations per :ref:`normalization unit <normalization-units>`
* - L1-L2 Bandwidth
- The number of bytes transferred across the vL1D-L2 interface as a result
of :ref:`VMEM <desc-vmem>` instructions, per
:ref:`normalization unit <normalization-units>`. The number of bytes is
calculated as the number of cache lines requested multiplied by the cache
line size. This value does not consider partial requests, so for
instance, if only a single value is requested in a cache line, the data
movement will still be counted as a full cache line.
- Bytes per :ref:`normalization unit <normalization-units>`
* - L1-L2 Reads
- The number of read requests for a vL1D cache line that were not satisfied
by the vL1D and must be retrieved from the to the
:doc:`L2 Cache <l2-cache>` per
:ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - L1-L2 Writes
- The number of write requests to a vL1D cache line that were sent through
the vL1D to the :doc:`L2 cache <l2-cache>`, per
:ref:`normalization unit <normalization-units>`.
- Requests per :ref:`normalization unit <normalization-units>`
* - L1-L2 Atomics
- The number of atomic requests that are sent through the vL1D to the
:doc:`L2 cache <l2-cache>`, per
:ref:`normalization unit <normalization-units>`. This includes requests
for atomics with, and without return.
- Requests per :ref:`normalization unit <normalization-units>`
* - L1 Access Latency
- Calculated as the average number of cycles that a vL1D cache line request
spent in the vL1D cache pipeline.
- Cycles
* - L1-L2 Read Access Latency
- Calculated as the average number of cycles that the vL1D cache took to
issue and receive read requests from the :doc:`L2 Cache <l2-cache>`. This
number also includes requests for atomics with return values.
- Cycles
* - L1-L2 Write Access Latency
- Calculated as the average number of cycles that the vL1D cache took to
issue and receive acknowledgement of a write request to the
:doc:`L2 Cache <l2-cache>`. This number also includes requests for
atomics without return values.
- Cycles
.. note::
All cache accesses in vL1D are for a single cache line's worth of data.
The size of a cache line may vary, however on current AMD Instinct MI CDNA
accelerators and GCN™ GPUs the L1 cache line size is 64B.
.. rubric :: Footnotes
.. [#vl1d-hit] The vL1D cache on AMD Instinct MI-series CDNA accelerators
uses a "hit-on-miss" approach to reporting cache hits. That is, if while
satisfying a miss, another request comes in that would hit on the same
pending cache line, the subsequent request will be counted as a "hit".
Therefore, it is also important to consider the access latency metric in the
:ref:`Cache access metrics <vl1d-cache-stall-metrics>` section when
evaluating the vL1D hit rate.
.. [#vl1d-activity] Omniperf considers the vL1D to be active when any part of
the vL1D (excluding the :ref:`address processor <desc-ta>` and
:ref:`data return <desc-td>` units) are active, for example, when performing
a translation, waiting for data, accessing the Tag or Cache RAMs, etc.
.. _vl1d-l2-transaction-detail:
vL1D - L2 Transaction Detail
----------------------------
This section provides a more granular look at the types of requests made
to the :doc:`L2 cache <l2-cache>`. These are broken down by the operation type
(read / write / atomic, with, or without return), and the
:ref:`memory type <memory-type>`.
.. _desc-td:
Vector L1 data-return path or Texture Data (TD)
===============================================
The data-return path of the vL1D cache, also known as the Texture Data
(TD) unit, is responsible for routing data returned from the
:ref:`vL1D cache RAM <desc-tc>` back to a wavefront on a SIMD. As described in
the :ref:`vL1D cache front-end <desc-ta>` section, the data-return path is passed
information about the space requirements and routing for data requests
from the :ref:`VALU <desc-valu>`. When data is returned from the
:ref:`vL1D cache RAM <desc-tc>`, it is matched to this previously stored request
data, and returned to the appropriate SIMD.
Omniperf reports the following vL1D data-return path metrics:
.. list-table::
:header-rows: 1
* - Metric
- Description
- Unit
* - Data-return Busy
- Percent of the :ref:`total CU cycles <total-cu-cycles>` the data-return
unit was busy processing or waiting on data to return to the
:doc:`CU <compute-unit>`.
- Percent
* - Cache RAM → Data-return Stall
- Percent of the :ref:`total CU cycles <total-cu-cycles>` the data-return
unit was stalled on data to be returned from the
:ref:`vL1D Cache RAM <desc-tc>`.
- Percent
* - Workgroup manager → Data-return Stall
- Percent of the :ref:`total CU cycles <total-cu-cycles>` the data-return
unit was stalled by the :ref:`workgroup manager <desc-spi>` due to
initialization of registers as a part of launching new workgroups.
- Percent
* - Coalescable Instructions
- The number of instructions submitted to the
:ref:`data-return unit <desc-td>` by the
:ref:`address processor <desc-ta>` that were found to be coalescable, per
:ref:`normalization unit <normalization-units>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Read Instructions
- The number of read instructions submitted to the
:ref:`data-return unit <desc-td>` by the
:ref:`address processor <desc-ta>` summed over all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`. This is expected to be
the sum of global/generic and spill/stack reads in the
:ref:`address processor <desc-ta>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Write Instructions
- The number of store instructions submitted to the
:ref:`data-return unit <desc-td>` by the
:ref:`address processor <desc-ta>` summed over all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`. This is expected to be
the sum of global/generic and spill/stack stores counted by the
:ref:`vL1D cache-front-end <ta-instruction-counts>`.
- Instructions per :ref:`normalization unit <normalization-units>`
* - Atomic Instructions
- The number of atomic instructions submitted to the
:ref:`data-return unit <desc-td>` by the
:ref:`address processor <desc-ta>` summed over all
:doc:`compute units <compute-unit>` on the accelerator, per
:ref:`normalization unit <normalization-units>`. This is expected to be
the sum of global/generic and spill/stack atomics in the
:ref:`address processor <desc-ta>`.
- Instructions per :ref:`normalization unit <normalization-units>`