← BACK TO LESSONPARALLEL HORIZONSREAD // 04

LESSON READING // 04

TENSOR COMPUTE

HISTORICAL ANCHOR2017 · VOLTA TENSOR CORE / WMMAREAD TIME · ABOUT 60 MIN

Re-express scalar multiply-add loops as warp-cooperative matrix tiles, then choose appropriate precision for inputs, multiplication, and accumulation.

Write down answers as you go; the estimate includes code tracing and practice.

PLAY THE INTERACTIVE LESSON
COURSE PACINGABOUT 60 MINUTES
  1. 01ORIENT + OBJECTIVES05 MIN
  2. 02MENTAL MODEL08 MIN
  3. 03DEEP DIVES15 MIN
  4. 04CODE TRACE12 MIN
  5. 05HISTORY + ECOSYSTEM08 MIN
  6. 06PRACTICE + REVIEW12 MIN
ACCELERATED COMPUTING CAPABILITY MAPINDEPENDENT LEARNING COMPANION · NOT AN NVIDIA COURSE OR CERTIFICATION
PATH // 04
GPU LIBRARIESPROFILINGMIXED PRECISION
THE OFFICIAL PATH EMPHASIZES

Prefer mature accelerated libraries and tools while selecting algorithms and data types through validated application evidence.

HOW THIS LESSON CONNECTS

The game explains Tensor Cores, warp tiles, and mixed precision; the lab compares a scalar GEMM against a library baseline.

SUGGESTED PREREQUISITES
  • Lessons 00–01
  • Basic matrix multiplication
  • FP16 / FP32 representation
PRACTICE TOOLS
  • cuBLAS
  • Nsight Compute
  • CUTLASS Profiler (optional)
01 · 05 MIN

ORIENT + OBJECTIVES

REMEMBER THIS FIRST

Specialized compute activates only when software expresses work in a shape the unit understands.

BY THE END, YOU SHOULD BE ABLE TO

  1. 01

    Distinguish scalar-thread matrix multiplication from warp-level matrix tiles.

  2. 02

    Explain cooperative fragment ownership and opaque lane mapping in WMMA.

  3. 03

    Analyze lower-precision inputs versus higher-precision accumulation.

  4. 04

    Decide among libraries, template libraries, and handwritten WMMA.

02 · 08 MIN

BUILD THE MENTAL MODEL

  1. 01
    TILE

    Take a fixed-shape piece from a large matrix as one warp-level matrix-operation unit.

  2. 02
    FRAGMENT

    The warp collectively owns a fragment; its element-to-lane mapping is an opaque implementation detail.

  3. 03
    ACCUMULATE

    Lower-precision inputs reduce compute and movement cost, while higher-precision accumulation can protect summation.

03 · 15 MIN

CONCEPT DEEP DIVES

01

WMMA IS WARP-SCOPE

The whole warp participates consistently in load, MMA, and store; this is not a matrix instruction owned by one thread.

02

MIXED PRECISION IS A PATH

It is not a global FP16 switch. Inputs, accumulators, master weights, and loss scaling have different jobs.

03

LIBRARIES ARE PART OF THE ECOSYSTEM

cuBLAS, cuDNN, and CUTLASS hide tiling, layout, and generational differences behind more stable interfaces.

PART 01

FROM OUTPUT PARALLELISM TO MATRIX-OP PARALLELISM

A conventional kernel can assign one C[row][col] to each thread while every thread still loops over K with scalar operations. That is parallel, but it does not express a Tensor Core matrix shape. WMMA makes a warp collectively own and compute a tile.

A large GEMM still needs hierarchical tiling: a block owns a larger tile, warps own subtiles, and data moves from global memory through shared memory into fragments. Tensor Cores accelerate core MMA; loading, layout, edges, and stores still control end-to-end efficiency.

PAUSE AND REASONWhy do 256 output threads not prove Tensor Core use?

REFERENCE ANSWERThey may execute ordinary scalar FMAs. A supported library or WMMA/MMA path must map operation, type, and shape onto Tensor Cores.

PART 02

FRAGMENT LANE MAPPING IS INTENTIONALLY OPAQUE

wmma::fragment represents matrix state distributed across a warp. Which lane stores which element is implementation-dependent and not a stable contract. Use load_matrix_sync, mma_sync, and store_matrix_sync to manipulate it.

Every participating lane must reach WMMA operations with consistent template arguments and control flow. Diverging half a warp around the operation can hang or corrupt results. Use only documented fragment element access when modification is necessary.

PAUSE AND REASONCan lane 0 be assumed to own the top-left element?

REFERENCE ANSWERNo. Element-to-lane mapping is opaque and may change across architectures.

PART 03

MIXED PRECISION PROTECTS AN ERROR PATH

Low-precision multiplicands reduce storage and compute cost, but summing many terms can lose small increments near a large partial sum. FP32 accumulation reduces that loss, yet cannot recover information already removed when inputs were quantized.

Training may also retain FP32 master weights, use loss scaling, or keep selected operations in high precision. Production code usually starts with cuBLAS, cuDNN, or framework AMP, then reaches for CUTLASS or WMMA only when customization is justified and numerically tested.

PAUSE AND REASONDoes FP16 input with FP32 accumulation guarantee full-FP32 results?

REFERENCE ANSWERNo. Input quantization has already lost information; FP32 accumulation only protects later summation.

04 · 12 MIN

PUT IT BACK INTO CODE

PROGRAM MODELA warp loads fragments and performs matrix multiply-accumulate
01wmma::fragment<matrix_a, 16, 16, 16, half, row_major> a;02wmma::fragment<matrix_b, 16, 16, 16, half, col_major> b;03wmma::fragment<accumulator, 16, 16, 16, float> c;04wmma::load_matrix_sync(a, A, lda);05wmma::load_matrix_sync(b, B, ldb);06wmma::mma_sync(c, a, b, c);

TRACE THE CODE

  1. 01
    DECLARETYPE AND SHAPE ARE PART OF THE CONTRACT

    A/B layout, input types, and tile shape must be supported by the target architecture.

  2. 02
    LOADTHE WARP LOADS CONSISTENTLY

    Base alignment and leading dimension must meet requirements; load converts storage into an opaque fragment.

  3. 03
    MMAEXPRESS D = A×B + C

    The warp cooperates in mma_sync; accumulator type changes summation, not already-quantized inputs.

  4. 04
    STORERETURN FRAGMENTS TO NORMAL LAYOUT

    store_matrix_sync writes distributed state to row- or column-major storage for normal access.

REF · 04

KNOWLEDGE ATLAS

OPTIONAL REFERENCE · REVISIT AS NEEDED · OUTSIDE THE CORE 60 MINUTES

CORE VOCABULARY

GEMM
General matrix multiplication C = αAB + βC, the compute core of many scientific and deep-learning operators.
Tile
A matrix subproblem shaped for a thread group, on-chip storage, and matrix instructions.
Tensor Core
Specialized matrix multiply-accumulate hardware requiring compatible type, layout, shape, and software path.
Accumulator
The intermediate precision used to sum products; it can exceed input or output precision.
Mixed Precision
Combining types to balance throughput, capacity, bandwidth, and numerical stability.
EXAMPLE

WORKED EXAMPLE

  1. 01
    COUNT WORK

    2 × 1,024³ ≈ 2.15 billion floating-point operations.

  2. 02
    COUNT MINIMUM BYTES

    A and B are about 2 MiB each and C is 4 MiB, for roughly 8 MiB total.

  3. 03
    COMPUTE INTENSITY

    About 2.15 billion FLOPs ÷ 8 MiB ≈ 256 FLOP/byte.

RESULT

High arithmetic intensity comes from repeated reuse of A/B tiles, not from tiny inputs.

WHY IT MATTERS

Real comparisons still fix shape, layout, types, and error bounds, and verify the library selected matrix hardware.

DIAGNOSTIC PLAYBOOK

01SYMPTOM

A GEMM call does not use matrix instructions

INSPECT FIRST
Types, compute type, layout/alignment, shape, and selected library algorithm
EVIDENCE
Actual profiler instructions and library logs—not the API name
02SYMPTOM

Tensor path is fast but inaccurate

INSPECT FIRST
Input range, accumulation, reduction length, and error definition
EVIDENCE
Absolute/relative error distributions against a high-precision reference
03SYMPTOM

Small-matrix throughput is poor

INSPECT FIRST
Launch cost, batch count, edge tiles, and parallelism
EVIDENCE
Shape-swept latency, batched-GEMM comparison, and occupancy

HARDWARE → ECOSYSTEM

  1. 01CUDA CORE FMA

    Threads explicitly implemented multiply-accumulate and tiling.

    Layout and register/shared-memory reuse set efficiency.
  2. 02VOLTA TENSOR CORE

    Warp-level matrix multiply-accumulate gained a specialized path.

    Software had to express matrix shape, type, and accumulation.
  3. 03LIBRARIES + TRANSFORMER ENGINE

    cuBLAS, CUTLASS, and frameworks select tiles and recipes per hardware.

    Optimization shifts toward library choice, numerical validation, and path profiling.
05 · 08 MIN

HARDWARE + ECOSYSTEM COORDINATE

2017 · VOLTA TENSOR CORES

Volta first added Tensor Cores to the SM, and CUDA 9 exposed warp-level matrix multiply-accumulate through WMMA. The software stack then mapped more data types, layouts, and model operators onto tensor paths. CUDA cores did not disappear.

06 · 12 MIN

PRACTICE + REVIEW

Who owns a 16×16×16 matrix tile in WMMA?

  1. A. One thread
  2. B. One cooperating warp
  3. C. The entire CPU process
REVEAL ANSWER
B

WMMA load, MMA, and store are warp-level operations. Software must not assume how fragment elements map to individual lanes.

02CLASSIFY

One output per thread, one tile per warp, one large tile per block: what levels are these?

HINT

Separate scalar ownership, warp cooperation, and block partitioning.

REFERENCE ANSWER

Scalar-thread parallelism, WMMA/MMA warp scope, and higher-level block tiling containing multiple warp tiles.

03NUMERICAL

Why can repeatedly adding 0.5 to 4096 disappear in FP16 accumulation?

HINT

Consider representable spacing near 4096.

REFERENCE ANSWER

FP16 spacing there exceeds 0.5, so each small increment rounds away. FP32 accumulation provides finer mantissa resolution.

04DEBUG

Only lanes threadIdx.x % 32 < 16 call mma_sync. What is wrong?

HINT

WMMA cooperation scope.

REFERENCE ANSWER

The warp does not participate uniformly, causing undefined behavior. The full warp must execute on consistent control flow.

05DECIDE

For a standard large GEMM, start with handwritten WMMA or cuBLAS?

HINT

Compare mature libraries with custom-kernel cost.

REFERENCE ANSWER

Usually cuBLAS first, with correct algorithms, layouts, and precision. Use CUTLASS/WMMA for justified fusion or special layouts.

LAB · 40–60 MIN

OPTIONAL HANDS-ON LAB

OUTSIDE THE CORE 60 MINUTES · REQUIRES A SUITABLE CUDA / GPU ENVIRONMENT

LAB GOAL

MOVE FROM SCALAR GEMM TO A LIBRARY + TENSOR PATH

Build a correct, repeatable GEMM baseline across scalar code, cuBLAS, and supported Tensor Core types.

PROCEDURE

  1. 01

    Create deterministic matrices and explicit error thresholds.

  2. 02

    Run the scalar kernel and record shape/layout/end-to-end time.

  3. 03

    Implement the same operation in cuBLAS with explicit input/accumulator/output types.

  4. 04

    Use Nsight Compute to confirm matrix instructions, memory, and occupancy.

EVIDENCE OF COMPLETION

Submit error, time, and profiler-path evidence for each implementation, plus why production usually starts with a library.

STRETCH CHALLENGE

Compare two CUTLASS Profiler tile/stage configurations under identical shape, type, and correctness.

OFFICIAL LEARNING ENTRY POINTScuBLAS DocumentationCUTLASS
REF

GO DEEPER

01
CUDA Programming Guide · Warp Matrix FunctionsWMMA fragments, shapes, and cooperation rules
OFFICIAL SOURCE · OPENS IN A NEW TAB02
Programming Tensor Cores in CUDA 9Historical context for Volta, libraries, and WMMA
OFFICIAL SOURCE · OPENS IN A NEW TAB