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.
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
Try “vector”, “memory” or “DMA”.
How work is expressed and scheduled
| Concept & meaning | NVIDIA GPUsHopper · CUDA | AMD Instinct GPUsCDNA 3 · HIP | Google TPUv6e · Pallas | AWS 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.
|
Kernel programTile operations use VMEM/SMEM buffers.
| ||
| Execution Group The hardware grouping below a program; these are different abstractions. |
Warp32 threads.
|
Vector lanes / sublanesElements of a VREG, not GPU threads.
| ||
| 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.
|
Physical compute blocks and engines
| Concept & meaning | NVIDIA GPUsHopper · CUDA | AMD Instinct GPUsCDNA 3 · HIP | Google TPUv6e · Pallas | AWS 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, Scalar, GPSIMDThe Scalar Engine also operates across many elements.
|
Where values live and how they move
| Concept & meaning | NVIDIA GPUsHopper · CUDA | AMD Instinct GPUsCDNA 3 · HIP | Google TPUv6e · Pallas | AWS TrainiumTrainium2 · NKI |
|---|---|---|---|---|
| Device Memory Large accelerator-attached storage for tensors. |
Global memory / HBM3For the MI300 example.
|
HBMExplicit transfers to and from SBUF.
|
||
| Operand Scratchpad Software-managed on-chip storage for reusable input and intermediate tiles. |
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.
| |
| Matrix Accumulators Where partial matrix results are kept between multiply-accumulate steps. |
Register accumulatorsFor Hopper mma / wgmma.
|
VGPRs / AccVGPRsDestination depends on the MFMA form.
| ||
| 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. |
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.
|
Ordering work and exchanging data
| Concept & meaning | NVIDIA GPUsHopper · CUDA | AMD Instinct GPUsCDNA 3 · HIP | Google TPUv6e · Pallas | AWS TrainiumTrainium2 · NKI |
|---|---|---|---|---|
| Local Ordering Ensure producers finish before dependent local consumers proceed. |
__syncthreads()A converged block barrier; not a grid barrier.
|
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.
|
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.
|
NeuronLink-v3Trainium2 device-to-device communication.
|
Precision, utilization and data access
| Concept & meaning | NVIDIA GPUsHopper · CUDA | AMD Instinct GPUsCDNA 3 · HIP | Google TPUv6e · Pallas | AWS 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.
|
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. |
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.
| |
| Branching and Selection Conditional work can waste parallel arithmetic; control and selection differ by model. |
Scalar control / vector selectNo GPU thread-warp divergence model.
|
Compiled scalar loopsPython branches are resolved during specialization.
|
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.
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.
Distinguish operand storage, multiply precision and accumulation. FP32 output does not guarantee full FP32 multiplication. Format support and throughput vary by instruction and generation.
Select “Show citations” to see the sources for each vendor term. Definitions summarize the common concept.
CUDA grids, blocks, warps, memory spaces, barriers, streams and matrix operations.
Hopper SMs, TMA, shared memory, caches, HBM, NVLink and occupancy.
mma and wgmma instructions, register fragments, matrix types and asynchronous completion.
Latency hiding, occupancy, coalescing, banking and register spilling.
CUs, Matrix Cores, XCD caches, HBM and Infinity Fabric links.
Thread/block/grid hierarchy, shared memory and SIMT execution.
Wave64, VALU/SALU, SGPR/VGPR/AccVGPR, LDS, MFMA, EXEC masking and wait counters.
Occupancy, divergence, memory coalescing and shared-memory banking.
One TensorCore per chip, MXUs, vector/scalar units, HBM and ICI.
Sequential grid semantics, vector registers, scalar control, block layout, spills and matrix precision.
VMEM/SMEM, VREG/SREG, scratch buffers, DMA, semaphores and pipelined transfers.
Remote DMA, send/receive semaphores and inter-device synchronization.
NeuronCore-v3, matrix formats, SBUF/PSUM, DMA, NeuronLink-v3 and engine overlap.
Sequential tensor programming, dependency-preserving scheduling, scalar values and control flow.
Software-managed HBM, SBUF and PSUM; explicit transfers and absence of a transparent hardware cache hierarchy.
Tensor memory placement and partition/free-dimension organization.
Stationary/moving tiles, nc_matmul, FP32 PSUM results and accumulation.
Tensor Engine/DMA utilization, blocking, reuse and overlapping data with compute.
Tensor, Vector, Scalar and GPSIMD engines; their distinct functions.
SPMD program coordinates within a launch grid.
SM CUDA cores and Tensor Cores, their matrix formats, register files and memory hierarchy.