Library Secrets at a Glance

The five pillars that take vendor BLAS to 90% of peak

Your tiled GEMM hits ~5% of peak. cuBLAS hits 80%+. The gap is built on five engineering pillars, accumulated over decades. Read them as a checklist; if any pillar cracks, vendor performance craters too.

Pillar	What it does	Where you observe it
Blocking / tiling	Keeps hot data in shared / L2 — same idea as your Track 6 tiled GEMM, but multi-level	Nsight 'shared mem reads', Xcode 'L2 hit ratio'
Micro-kernel (Tensor / matrix unit)	Fused matrix-multiply ops doing 32–128 FLOPs per instruction (not per FMA)	SASS `mma.*` instructions, Apple Instruments matrix-unit counters
Split-K & Streaming-K	Parallelize the K-dimension across blocks, then reduce partial sums	`cublasLtMatmulAlgoGetHeuristic`, MLX's split-K paths
Heuristic algo picker	Auto-searches tile sizes / epilogue / splitK for your specific shape	Nsight 'Algorithm ID' panel, cuBLASLt logs
Epilogue fusion	Folds bias add + activation + residual add into the kernel's final loop	cuBLASLt `activationType`, `β·C` handling

Read the table once. Then notice: your tiled GEMM has only the first pillar. That's why it's at 5%. Each remaining pillar is independently worth ~2× — multiplied together, you understand the gap.

What breaks each pillar in production

Tiny m or k — micro-kernel can't fill its expected shape; library falls back to GEMV-grade kernel.
Hidden size not a multiple of 16 — Tensor Core paths require this; off-multiples force a slower fallback.
L2 thrashing — concurrent kernels evict your tiles; performance halves with no obvious cause.
Bank conflicts in shared memory — your tile loader stalls; profiler shows high 'shared bank conflicts'.

Code

Reading cuBLASLt heuristic output — what algo got picked?·cuda

// cuBLASLt is the way to make the algo picker visible.
// (Pseudocode — the real API is verbose; consult the docs.)

cublasLtMatmulHeuristicResult_t heuristic[10];
int returnedResults = 0;
cublasLtMatmulAlgoGetHeuristic(
    ltHandle,                 // lib handle
    matmulDesc,               // operation description
    Adesc, Bdesc, Cdesc, Cdesc,
    preference,
    /*requestedAlgoCount=*/ 10,
    heuristic,
    &returnedResults);

for (int i = 0; i < returnedResults; ++i) {
    printf("Algo %d: workspace=%zu, waves=%f\n",
           i, heuristic[i].workspaceSize, heuristic[i].wavesCount);
}
// First entry is the library's best pick for your specific (m, n, k, dtype).

Force vs auto — when to override the library's pick·cuda

// Sometimes the heuristic is wrong (especially for unusual shapes).
// You can sweep the available algos and pick the actual winner:

// 1. Get all algorithm IDs supported for this op
int algos[32]; int n_algos;
cublasLtMatmulAlgoGetIds(handle, ..., 32, algos, &n_algos);

// 2. Time each one over a few warmup runs
float best_ms = INFINITY; int best = -1;
for (int i = 0; i < n_algos; ++i) {
    float ms = bench_algo(algos[i], A, B, C, m, n, k);
    if (ms < best_ms) { best_ms = ms; best = i; }
}

// 3. Cache the winner — algo selection is shape-specific
shape_to_algo[std::make_tuple(m, n, k)] = algos[best];

Exercise

On a CUDA box, run a 4096³ FP16 GEMM through cuBLAS once with default settings, then explore which Tensor Core algo it picked (use Nsight Compute or read cuBLASLt logs). Then deliberately use a non-Tensor-Core algo and re-run. The 2–3× gap you observe is the value of pillar 2 (micro-kernel) alone — and it's selected automatically by the library because of pillar 4 (heuristic picker).

The five pillars that take vendor BLAS to 90% of peak

What breaks each pillar in production

Code

External links

Exercise

Progress

Comments 0