TL;DR: Tensor cores accelerate supported matrix-multiply-accumulate operations. A kernel must use a compatible instruction, operand format and fragment layout. Libraries can tile, pad or handle tails for awkward matrix shapes; an odd dimension does not automatically disable tensor cores. Confirm the executed path and measure its time. Tensor-core use does not guarantee compute-bound performance.
How to approach it
Define the operation as D = AB + C, then distinguish a matrix from the instruction tiles used to compute it. Name one concrete supported form. Separate numerical settings, memory delivery and hardware evidence before proposing a performance change.
A strong answer
A linear layer delivers a small fraction of the advertised BF16 peak. That result could come from memory traffic, a small problem, launch overhead or a different arithmetic path. Low throughput alone cannot identify which one.
Tensor Cores and Matrix Units perform matrix operations on structured operand fragments. A fragment is the portion of a tile held or described by participating threads. A complete general matrix multiplication (GEMM) combines many such operations with data movement, address calculations and output handling.
For a concrete form, the PTX instruction reference (PTX is NVIDIA's low-level virtual instruction set that CUDA code compiles to) includes mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32. It describes BF16 operands with a 16×8 output tile and K=16, accumulating into FP32. This is an instruction identifier to inspect, not a standalone runnable program. Hopper also provides warpgroup operations; later architectures have other interfaces. Do not apply one form's layout rules to every tensor operation.
The worked count is 16×8×16 = 2,048 multiply-adds, or 4,096 FLOPs using two operations per multiply-add. This count describes work, not instruction latency or the number of cycles needed to retire it.
| Requirement | What to verify |
|---|---|
| Supported operation | The selected kernel actually issues matrix instructions |
| Input and accumulator formats | Match the exact instruction and required accuracy |
| Operand layout | Fragments, alignment and synchronization follow that interface |
| Whole-problem shape | Library handling of edges, padding and tiles is appropriate |
| Supply of work and data | Enough useful tiles and reuse to justify the chosen path |
An instruction has a fixed tile shape; a user's entire matrix need not have that shape or be divisible by it. A library can choose edge handling, padding or another kernel. That handling may cost time, so benchmark it, but “dimension 32,001 means scalar fallback” is not a valid diagnosis. NVIDIA's matrix guide discusses alignment and tiling efficiency.
Likewise, FP32 input tensors do not establish full FP32 multiplication precision: a configured TF32 path may use reduced input precision with a wider exponent range. Some devices also support FP64 matrix instructions. Check the requested numerical mode and generated path instead of treating “tensor core” as one dtype. Wider accumulation helps preserve sums; it does not make input rounding disappear or make all accumulated arithmetic exact.
Here is a traffic-bound example, independent of which suitable kernel is chosen. Let M=8, N=K=8,192; read BF16 A and B once and write a BF16 result once. Exclude a separate C read by taking its coefficient as zero:
F = 2 × 8 × 8192 × 8192 = 1,073,741,824 FLOPs
D = 2 × (8×8192 + 8192×8192 + 8×8192)
= 134,479,872 bytes
intensity = F / D ≈ 7.984 FLOPs/byte
HBM floor at 3.35 TB/s = 40.143 microseconds
compute floor at 989 TFLOP/s = 1.086 microseconds
sanity: the weight matrix alone is 134,217,728 bytes.
Those H100 SXM peak specifications give ideal resource bounds. A compatible tensor-core kernel can still be constrained by data delivery at this small M. The calculation does not measure 40 microseconds of runtime or prove an idle percentage. Caching, extra traffic and overhead can change the observed result. The Roofline Model explains the distinction.
Use a profiler to find the actual kernel and inspect its instruction/tensor pipeline evidence. Nsight Compute's Compute Workload Analysis and instruction views are more informative than inferring implementation from a kernel name. Metric availability varies by device and tool version. Low tensor activity can mean limited matrix work or poor data supply, not necessarily absence of matrix instructions.
I would first compare the supported library path with an equal-accuracy reference. If the trace shows substantial tail or layout overhead, test padding or a better layout and measure the complete operation. Reverse that change if conversion, extra arithmetic or added parameters outweigh the gain. A smaller non-tensor kernel can be the fastest correct choice for a tiny problem.
What interviewers probe next
- “Can softmax use tensor cores?” Softmax itself is not a matrix product. A fused attention kernel can combine matrix instructions with reductions and elementwise work; inspect each portion.
- “Can I quote the sparse peak?” Only for the supported sparsity format and operation. A dense matrix does not automatically earn that peak.
- “Does BF16 storage guarantee BF16 execution everywhere?” No. Conversions, accumulator formats and non-matrix operations can differ.
- “Is a tensor kernel always preferable?” No. Compare correct end-to-end time at the actual shape, including any setup.
Common mistakes
- Turning fixed instruction tiles into a blanket divisibility rule for user matrices.
- Calling an ideal bandwidth floor achieved performance.
- Treating a kernel name or low activity percentage as proof of fallback.
- Assuming FP32 accumulation restores information lost in lower-precision inputs.
Key takeaways
- The example m16n8k16 operation performs 4,096 FLOPs on one tile.
- The small-M example has about 7.984 FLOPs/byte; tensor-core use alone cannot overcome its traffic bound.
- Verify the instruction path, numerical mode and measured runtime separately.