← Learning Notes

AI Accelerator Terminology Cheat Sheet

Download PDF

Shared roles, different architectures. Read each column within its stated scope.

Examples: Hopper · CDNA 3 · TPU v6e · Trainium2. APIs: CUDA · HIP · Pallas · NKI. Sources & Scope

01Execution

How work is expressed and scheduled

Concept & meaningNVIDIA GPUsHopper · CUDAAMD Instinct GPUsCDNA 3 · HIPGoogle TPUv6e · PallasAWS TrainiumTrainium2 · NKI
Execution Model

How a program expresses parallel work; the execution models differ.

SIMTA kernel runs as many threads.
SIMTHIP threads map to wavefronts.
Vector/tile programA sequential instruction stream with asynchronous units.
Tensor programSequential semantics; independent engine work may overlap.
Work Grid

The index space used to divide an operation into program instances.

Grid of thread blocksblockIdx + threadIdx
Grid of thread blocksblockIdx + threadIdx
Pallas gridGenerally iterated sequentially on a core.
SPMD launch gridnl.program_id identifies an instance.
Cooperation Scope

A scope for sharing temporary data within local work.

Thread blockThreads share block shared memory.
Thread block / work-groupThreads share LDS.
Kernel programTile operations use VMEM/SMEM buffers.
Kernel instanceEngine operations use SBUF/PSUM tiles.
Execution Group

The hardware grouping below a program; these are different abstractions.

Warp32 threads.
Wavefront64 work-items on CDNA 3; HIP calls it a warp.
Vector lanes / sublanesElements of a VREG, not GPU threads.
Engine-wide tile operationNo CUDA-style warp in the NKI model.
Matrix Operation

An instruction or kernel API that requests matrix multiply-accumulate work.

mma / wgmmaWarp / asynchronous warp-group operations.
MFMAMatrix fused multiply-add instructions.
lax.dot / lax.dot_generalPallas lowers matrix operations to MXUs.
nisa.nc_matmulStationary and moving operand tiles.

02Hardware

Physical compute blocks and engines

Concept & meaningNVIDIA GPUsHopper · CUDAAMD Instinct GPUsCDNA 3 · HIPGoogle TPUv6e · PallasAWS TrainiumTrainium2 · NKI
Compute Block

A repeated block containing arithmetic and local storage; sizes and capabilities differ.

SMStreaming Multiprocessor.
CUCompute Unit; CUs are grouped in XCDs.
TensorCoreOne per v6e chip; includes matrix, vector and scalar units.
NeuronCore-v3Eight physical cores per Trainium2 chip.
Matrix Engine

Specialized hardware for dense matrix arithmetic.

Tensor CoreInside an SM.
Matrix CoreInside a CU.
MXUMatrix-multiply unit; two per v6e TensorCore.
Tensor EngineSystolic matrix engine in a NeuronCore.
General Arithmetic

Hardware for work around matrix multiplication, such as elementwise operations and control.

CUDA cores / SM pipelinesThread arithmetic and other instructions.
VALU + SALUVector work and wave-uniform scalar work.
Vector unit + scalar unitArray arithmetic and scalar control.
Vector, Scalar, GPSIMDThe Scalar Engine also operates across many elements.

03Memory

Where values live and how they move

Concept & meaningNVIDIA GPUsHopper · CUDAAMD Instinct GPUsCDNA 3 · HIPGoogle TPUv6e · PallasAWS TrainiumTrainium2 · NKI
Device Memory

Large accelerator-attached storage for tensors.

Global memory / HBMFor the H100 example.
Global memory / HBM3For the MI300 example.
HBMStaged into on-chip memory for computation.
HBMExplicit transfers to and from SBUF.
Operand Scratchpad

Software-managed on-chip storage for reusable input and intermediate tiles.

Shared memoryBlock-scoped; Hopper also supports cluster access.
LDSLocal Data Share; work-group allocation.
VMEM / SMEMVector SRAM / scalar SRAM.
SBUFState Buffer shared by the core’s compute engines.
Register Storage

Storage closest to arithmetic; ownership and allocation differ.

RegistersPer-thread architectural values.
VGPRs / SGPRsPer-lane vector / wave-uniform scalar values.
VREGs / SREGsVector / scalar registers.
Scalar VirtualRegisterNKI tensor values live in SBUF/PSUM.
Matrix Accumulators

Where partial matrix results are kept between multiply-accumulate steps.

Register accumulatorsFor Hopper mma / wgmma.
VGPRs / AccVGPRsDestination depends on the MFMA form.
FP32 vector valuesFor floating-point Pallas dot; VREGs and VMEM scratch.
PSUMPartial Sum Buffer; dedicated matrix-result memory.
Tensor Layout

How logical tensor elements map to the storage consumed by an instruction.

Register fragments / shared descriptorsMMA operands use instruction-specific register or shared-memory layouts.
MFMA operand layoutInstruction-specific lane/VGPR mapping.
VREG lanes / sublanesTiling and padding follow vector dimensions.
Partition + free dimensionsSBUF/PSUM use partitioned 2D storage.
Bulk Data Movement

A mechanism to stage large tiles while arithmetic can proceed.

TMATensor Memory Accelerator; global to/from shared copies.
Global-to-LDS loadsISA loads can target LDS; completion needs a wait.
DMAHBM to/from VMEM/SMEM copies.
DMA enginesHBM to/from SBUF; independent of compute engines.
Cache or Explicit Buffer

Whether fast-memory contents are filled transparently or by program-visible transfers.

L1 / L2 cachesShared memory is explicitly managed.
L1 / L2 cachesL2 is shared by CUs within an XCD.
Explicit VMEM/SMEM stagingPallas exposes SRAM buffers and DMA.
Software-managed hierarchySBUF/PSUM are not transparent hardware caches.
Private Temporaries

Storage for values local to execution; “local” does not imply on-chip SRAM.

Local memoryThread-private device-memory space, including spills.
Scratch memoryPrivate per-work-item storage in memory.
VREG spills to VMEMProgram scratch is not per-thread local memory.
SBUF intermediatesKernel-local tensor storage in the NKI model.

04Coordination

Ordering work and exchanging data

Concept & meaningNVIDIA GPUsHopper · CUDAAMD Instinct GPUsCDNA 3 · HIPGoogle TPUv6e · PallasAWS TrainiumTrainium2 · NKI
Local Ordering

Ensure producers finish before dependent local consumers proceed.

__syncthreads()A converged block barrier; not a grid barrier.
__syncthreads()HIP block barrier. ISA s_barrier alone does not complete memory operations.
Program dependenciesSequential core program; asynchronous work needs completion tracking.
Program dependenciesNKI preserves sequential semantics when scheduling engines.
Transfer Completion

Wait for an in-flight data transfer before reading or reusing its buffers.

cuda::barrier / pipelineTrack asynchronous shared-memory copies.
s_waitcntISA counters track relevant outstanding memory operations.
DMA semaphore + waitCopy descriptors track completion.
Compiler dependenciesNKI schedules required ordering between DMA and consumers.
Chip-to-Chip Exchange

Dedicated links carry data between accelerator chips; topology and scope vary.

NVLinkGPU-to-GPU links; separate from the SM hierarchy.
Infinity Fabric linksConnect MI300X GPUs within an eight-GPU node.
ICIInter-chip interconnect; Pallas exposes remote DMA.
NeuronLink-v3Trainium2 device-to-device communication.

05Performance

Precision, utilization and data access

Concept & meaningNVIDIA GPUsHopper · CUDAAMD Instinct GPUsCDNA 3 · HIPGoogle TPUv6e · PallasAWS TrainiumTrainium2 · NKI
Matrix Input Precision

Example arithmetic formats; a tensor’s storage type alone does not select the math.

FP8 / FP16 / BF16 / TF32Examples of Hopper Tensor Core modes.
FP8 / FP16 / BF16 / INT8Examples of CDNA 3 Matrix Core modes.
BF16 / INT8Documented v6e matrix compute examples.
FP8 / FP16 / BF16 / TF32Examples of Trainium2 Tensor Engine modes.
Accumulation Precision

The running sum can be wider than the multiply inputs; mode and instruction matter.

FP32 accumulationAvailable for BF16/FP16 matrix operations.
MFMA_F32 formsBF16/FP16 inputs can accumulate into FP32.
FP32 floating-point dot resultsFP32 inputs may be rounded to BF16 by default.
FP32 PSUMNKI matmul writes FP32 partial results.
Latency Hiding

Overlap useful work with data waits. Occupancy counts resident state, not utilization.

Warp occupancy + pipelinesResident warps and asynchronous tile staging.
Wave occupancy + pipeliningResident waves and staged LDS data.
Software pipeliningOverlap DMA with a core’s computation.
Engine/DMA overlapPipeline tiles; inspect each engine’s utilization.
Memory Access Efficiency

Match data access patterns to the memory system to reduce wasted transfers.

Coalescing / bank conflictsWarp global accesses / shared-memory banks.
Coalescing / LDS bankingWave accesses / shared-memory banks.
Block shape / paddingTile layout must fit the vector/matrix units.
DMA efficiency / partition layoutBlock transfers and reuse SBUF tiles.
Branching and Selection

Conditional work can waste parallel arithmetic; control and selection differ by model.

Warp divergence / predicationDifferent thread paths can reduce efficiency.
Wave divergence / EXEC maskInactive lanes do not contribute useful work.
Scalar control / vector selectNo GPU thread-warp divergence model.
Compiled scalar loopsPython branches are resolved during specialization.

Notes

Same Role, Different Unit

An SM, CU, Google TensorCore and NeuronCore have different sizes and capabilities. Google TensorCore names a compute block; NVIDIA Tensor Core names a matrix engine.

Scratchpads Have Different Scopes

CUDA shared memory and AMD LDS support thread cooperation. TPU VMEM and Trainium SBUF serve tile programs and engines. These are role comparisons, not interchangeable memory models.

Precision Has Three Parts

Distinguish operand storage, multiply precision and accumulation. FP32 output does not guarantee full FP32 multiplication. Format support and throughput vary by instruction and generation.

Sources & Scope 21 references

Select “Show citations” to see the sources for each vendor term. Definitions summarize the common concept.

  • Hardware examples are explicitly bounded to Hopper, CDNA 3, TPU v6e and Trainium2; this is not a claim about every NVIDIA, AMD, Google or AWS chip.
  • CUDA and HIP use a thread hierarchy. Pallas TPU and NKI express tile/tensor programs. A Pallas grid is generally sequential on a core; an NKI program_id is an SPMD instance coordinate, not a thread ID.
  • Pallas and NKI entries are software interfaces to hardware concepts, identified in column subtitles. Pallas TPU is experimental; these APIs can evolve. The matrix-format row describes hardware examples, not a promise that every type is accepted by every kernel API.
  • Precision rows show selected matrix modes, not all supported storage types or instructions. TF32 is reduced-precision multiplication; FP8 variants are not automatically numerically interchangeable.
  • The original publication date is retained; this expanded four-platform edition was reviewed on October 21, 2026. No cross-platform speed or popularity ranking is implied.
  1. 01
    NVIDIA · CUDA C++ Programming Guide 13.0.2

    CUDA grids, blocks, warps, memory spaces, barriers, streams and matrix operations.

  2. 02
    NVIDIA · Hopper Tuning Guide

    Hopper SMs, TMA, shared memory, caches, HBM, NVLink and occupancy.

  3. 03
    NVIDIA · PTX ISA 9.0

    mma and wgmma instructions, register fragments, matrix types and asynchronous completion.

  4. 04
    NVIDIA · CUDA Best Practices 13.0.2

    Latency hiding, occupancy, coalescing, banking and register spilling.

  5. 05
    AMD · MI300 / CDNA 3 Microarchitecture

    CUs, Matrix Cores, XCD caches, HBM and Infinity Fabric links.

  6. 06
    AMD · HIP Programming Model 6.3.1

    Thread/block/grid hierarchy, shared memory and SIMT execution.

  7. 07
    AMD · MI300 / CDNA 3 ISA

    Wave64, VALU/SALU, SGPR/VGPR/AccVGPR, LDS, MFMA, EXEC masking and wait counters.

  8. 08
    AMD · HIP Performance Guidelines

    Occupancy, divergence, memory coalescing and shared-memory banking.

  9. 09
    Google · TPU v6e Architecture

    One TensorCore per chip, MXUs, vector/scalar units, HBM and ICI.

  10. 10
    JAX · Writing TPU Kernels with Pallas

    Sequential grid semantics, vector registers, scalar control, block layout, spills and matrix precision.

  11. 11
    JAX · TPU Pipelining

    VMEM/SMEM, VREG/SREG, scratch buffers, DMA, semaphores and pipelined transfers.

  12. 12
    JAX · Distributed Pallas on TPUs

    Remote DMA, send/receive semaphores and inter-device synchronization.

  13. 13
    AWS · Trainium2 NKI Architecture · Neuron 2.33.0

    NeuronCore-v3, matrix formats, SBUF/PSUM, DMA, NeuronLink-v3 and engine overlap.

  14. 14
    AWS · NKI Language Guide · Neuron 2.33.0

    Sequential tensor programming, dependency-preserving scheduling, scalar values and control flow.

  15. 15
    AWS · Trainium Memory Hierarchy · Neuron 2.33.0

    Software-managed HBM, SBUF and PSUM; explicit transfers and absence of a transparent hardware cache hierarchy.

  16. 16
    AWS · NKI Data Representation · Neuron 2.33.0

    Tensor memory placement and partition/free-dimension organization.

  17. 17
    AWS · NKI Matrix Multiplication · Neuron 2.33.0

    Stationary/moving tiles, nc_matmul, FP32 PSUM results and accumulation.

  18. 18
    AWS · NKI Kernel Optimization · Neuron 2.33.0

    Tensor Engine/DMA utilization, blocking, reuse and overlapping data with compute.

  19. 19
    AWS · NeuronCore-v3 Architecture · Neuron 2.33.0

    Tensor, Vector, Scalar and GPSIMD engines; their distinct functions.

  20. 20
    AWS · NKI program_id · Neuron 2.33.0

    SPMD program coordinates within a launch grid.

  21. 21
    NVIDIA · Hopper Architecture In-Depth

    SM CUDA cores and Tensor Cores, their matrix formats, register files and memory hierarchy.