Tensor Programs → GPU Kernels, Explained. Part 2

Part 1: Semantic Graphs traced how Inductor, MAX, TinyGrad, and Luminal represent y = torch.matmul(x, weight.T) + b. Here we follow one selected implementation from each compiler. The same linear layer becomes one to five GPU kernels, depending on how the compiler handles the matrix product, intermediate storage, and bias addition.

The workload is unchanged: M = N = 1024, K = 768, with FP16 inputs and output, FP32 matrix arithmetic, and an FP16 projection rounded before bias addition. Keep two questions in mind: where is that projection stored, and which implementation computes it?

We use the pinned October 18, 2026 captures: expanded autotuning for Inductor, MAX's captured library-dispatch configuration (tile_based_fusion=False), a completed BEAM=2 search for TinyGrad, and expanded plan search for Luminal. MAX has no separate tuning experiment here. These are concrete selected implementations, not evidence of a global optimum or a performance ranking.

From Graph Operations to GPU Execution

Five interacting decisions turn a semantic graph into GPU launches. They need not happen in this order:

Autotuning measures candidates before selecting one. Those candidates can include precompiled library calls as well as generated kernels.

Inside a kernel, a thread block contains 32-thread warps and processes a tile of tensor elements. A schedule assigns work to threads and decides how to load operands into registers, optionally staging them in shared memory. Tensor Cores perform matrix multiply-accumulate operations.

Generated CUDA is compiled through PTX, NVIDIA's virtual instruction set, to native SASS; a compiler can also emit PTX directly. On H100, HMMA implements warp-level matrix operations, while HGMMA implements operations across a warpgroup of four warps.

Inductor Fuses Matrix and Bias in a Triton Template

Inductor lowers the FX graph from Part 1 into a loop-level IR, where Python functions describe how to compute individual output elements. Operator lowerings can compose into generated kernels, while specialized operations can use external implementations.

With autotuning enabled, the selected implementation is a single Triton matrix-and-bias template. Its indexing reads the transposed weight without making a copy. Each block computes a 64×128 output tile with 4 warps, processing the reduction in chunks of BLOCK_K = 64. Five software pipeline stages overlap operand loads with matrix arithmetic. These selected lines show the schedule and projection rounding (generated caller):

Scroll code horizontally to read long lines.

num_stages=5,
num_warps=4,
    ACC_TYPE : tl.constexpr = tl.float32
    BLOCK_M : tl.constexpr = 64
    BLOCK_N : tl.constexpr = 128
    BLOCK_K : tl.constexpr = 64
        acc += tl.dot(a, b, allow_tf32=ALLOW_TF32, out_dtype=ACC_TYPE)
    tmp0 = acc.to(tl.float16)
    tmp1 = tmp0.to(tl.float32)

tmp0 rounds the FP32 accumulator to FP16 and tmp1 widens it for bias addition. This preserves Part 1's projection boundary inside the fused kernel, without a separate global projection buffer.

The selected PTX uses wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 for FP32 accumulation. It then executes cvt.rn.f16x2.f32 before add.f16x2. The packed FP16 add implements the widened addition and final narrowing on already-rounded operands. For these finite FP16 values, it produces the same rounded result.

Implementation selection therefore determines more than the matrix routine: selecting a template with an epilogue lets bias share the matrix kernel.

MAX Dispatches Two Kernels from One Graph Routine

MAX replaces the explicit transpose with an NK layout annotation on weight (%arg1). This records that the weight is stored as [N, K], so the matrix kernel can consume it without a transposed copy. The graph also groups matrix multiplication and bias inside one mogg.routine (routine excerpt):

Scroll code horizontally to read long lines.

%0 = mo.rebind(%arg1) : !mo.tensor<[1024, 768], f16, gpu:0> -> !mo.tensor<[1024, 768], f16,
    gpu:0, {layout = #mo.layout<NK>}> loc(#loc)
context (%arg7) attributes {device = #M.device_ref<"gpu", 0>, fusionInfo =
    #matmul_static_broadcast_to_add, hasReadEffect = false, hasWriteEffect = false} {
mogg.call @__matmul_static_broadcast_to_add %arg3, %arg4, %arg5, %arg6, %arg7 :
    !mogg.pointer, !mogg.pointer, !mogg.pointer, !mogg.pointer, !mogg.context loc(#loc)

The next lowering stage extracts input buffers and allocates the 2,097,152-byte (2 MiB, 256-byte aligned) FP16 output buffer %8 before asynchronous dispatch (buffer excerpt):

Scroll code horizontally to read long lines.

%idx2097152 = index.constant 2097152 loc(#loc)
%8 = mgp.buffer.alloc %0, %idx2097152, %arg4, align 256 : !mgp.buffer<<"gpu", 0>>,
    !mgp.device_context<<"gpu", 0>> loc(#loc)

Under tile_based_fusion=False, the device timeline records two GPU launches: a cuBLAS matrix kernel followed by MAX's elementwise bias kernel, with no transpose launch. The dispatch log identifies the vendor BLAS fallback, and the machine-code capture contains HGMMA.64x64x16.F32.

MAX: One Graph Routine, Two GPU Kernels

Captured library path · tile_based_fusion=False

1. Semantic graph (rmo)Transpose → MatMul → add; explicit [768, 1024] view.
2. Graph routine (mogg)NK weight layout; MatMul and bias grouped together.
3. Buffers (mgp)Allocate 2 MiB FP16 storage, then dispatch.
4. GPU launchescuBLAS matrix → bias: two kernels, no transpose kernel.

Grouping operations in a graph routine does not guarantee a single kernel. MAX also lets authors write kernels in Mojo and integrate them into a graph, placing tiling and memory-staging decisions in kernel code. That is an extension point for human or agent authors; the path captured here uses the library matrix implementation.

TinyGrad Builds a Fused Tensor Core Kernel from Primitives

TinyGrad starts with the multiply-and-reduce primitives from Part 1. It must reconstruct the matrix multiplication pattern from those operations and their indexing before selecting Tensor Core instructions. Its lowering pipeline makes iteration ranges and kernel boundaries explicit; BEAM search explores schedules for the resulting kernel.

The selected program from a completed BEAM=2 search fuses matrix multiplication and bias. It launches a 32×8 grid of 32×2 thread blocks: 256 blocks, each with 64 threads. Each block computes a 128×32 output tile. Address calculations read x[m*768+k], weight[n*768+k], and b[n], so no transposed weight buffer is needed.

The generated CUDA helper issues inline PTX. Its name lists the fragment shape in TinyGrad's N, M, K order, so 8_16_16 matches the instruction's m16n8k16:

Scroll code horizontally to read long lines.

__device__ float4 __WMMA_8_16_16_f16_f32(half8 a, half4 b, float4 c){
  int *a_pk = (int *)(&a), *b_pk = (int *)(&b), *c_pk = (int *)(&c);
  asm("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32"
      "{%0, %1, %2, %3}, {%4, %5, %6, %7},"
      "{%8, %9}, {%0, %1, %2, %3};"
    : "+r"(c_pk[0]), "+r"(c_pk[1]), "+r"(c_pk[2]), "+r"(c_pk[3])
    : "r"(a_pk[0]), "r"(a_pk[1]), "r"(a_pk[2]), "r"(a_pk[3]), "r"(b_pk[0]), "r"(b_pk[1]));
  return c;
}

The instruction computes a 16×8 matrix fragment cooperatively across a warp, with a K = 16 step, FP16 inputs, and FP32 accumulation (instruction semantics). Multiple instructions cover the block's larger tile and reduction chunk.

The generated CUDA loops over 24 chunks of 32 reduction elements (768 = 24 × 32). Threads load pairs of FP16 values directly from global memory, with no shared-memory staging. Each thread holds 64 FP32 accumulator values: 64 threads × 64 values = 4,096 outputs per block. The epilogue casts each accumulator to FP16 before adding bias; the PTX confirms cvt.rn.f16.f32 before add.f16.

TinyGrad: Matrix and Bias in One Kernel

Selected BEAM=2 run · no shared-memory staging

1. Map the output tileGrid 32×8; block 32×2 threads; 128×32 outputs per block.
2. Reduce over K24 chunks × 32 elements; 64 FP32 accumulator values per thread.
3. Round and add biasFP32 → FP16 → add bias; write the final FP16 output.

Luminal Measures Alternative Execution Plans

Luminal's pinned CUDA backend (77d3060) has matrix-site rules for FP32 and BF16 inputs, but none that read FP16 inputs and accumulate in FP32. The explicit FP32 casts from Part 1 therefore become conversion kernels that materialize FP32 copies of x and weight before cuBLASLt runs.

Extraction selects compatible expressions, layouts, and backend operations from the e-graph to form an executable plan. Luminal measures complete plans, allowing search to compare storage and layout choices along with computation. The expanded search measures 107 plans and selects five GPU kernels: cast x, cast and transpose weight, multiply, round the projection to FP16, and add bias. The matrix layout produces the projection in the required orientation, so no projection transpose copy is needed.

The matrix call selects a cuBLAS FP32 kernel whose sampled instructions are scalar FFMA, rather than Tensor Core instructions. A generated kernel then rounds the 4 MiB FP32 projection to 2 MiB of FP16 values:

Scroll code horizontally to read long lines.

extern "C" __global__ void k(const float* a, __half* out, const long long* params) {
    const unsigned long long n = 1048576LL;
    unsigned long long i = (unsigned long long)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = (__half)a[i];
}

The bias kernel reads that rounded projection and b[n], then writes the final output. The plan marks its four temporaries free after their last uses. The device trace records the five kernels within a CUDA Graph, so five GPU kernels do not imply five host submission calls.

Compare the Selected Execution Paths

The four paths preserve the same projection rounding at different places: inside a fused kernel, in a buffer shared by two kernels, or in a separate conversion kernel.

One Selected Implementation per Compiler

One step = one GPU kernel. Steps show execution order, not duration.

InductorExpanded autotune · 1 kernel

1. Triton matrix + biasProjection rounded within the fused kernel.

MAXLibrary dispatch · 2 kernels

1. cuBLAS matrixWrite the FP16 projection.
2. BiasRead the projection and add bias.

TinyGradBEAM=2 · 1 kernel

1. Generated matrix + biasProjection rounded within the fused kernel.

LuminalExpanded search · 5 kernels

1. Cast xFP16 → FP32.
2. Cast weightFP32 conversion and transpose.
3. cuBLASLt matrixFP32 arithmetic with scalar FFMA.
4. Round projectionFP32 → FP16.
5. BiasAdd bias and write the final FP16 output.

Reason about Hardware Limits

A roofline model compares compute throughput for the instruction path in use with memory bandwidth multiplied by arithmetic intensity (operations per byte moved). Bandwidth and byte counts must refer to the same memory level, such as HBM or L2 cache. The lower ceiling bounds throughput in that model. Luminal's scalar FP32 matrix path has a different compute ceiling from the FP16-input Tensor Core paths. Dependency stalls and insufficient parallel work can keep any kernel below its ceiling; host submission overhead adds a separate cost.

This matrix multiply performs about 2MNK, or 1.6 billion, floating-point operations. Reusing operands across a tile can amortize their loads. A separate bias kernel does just one addition per output plus loads and stores, giving it much lower arithmetic intensity (NVIDIA's matrix multiplication guide). The selected structures suggest these tradeoffs:

Intermediate storage makes one cost concrete. Counting one write and one read per materialized intermediate gives MAX 4 MiB for its FP16 projection. Luminal totals 24 MiB: FP32 input copies (12 MiB), the FP32 projection (8 MiB), and the FP16 projection (4 MiB). Inductor and TinyGrad have no separate projection buffer. This accounting excludes original input reads, bias, final output, repeated tile loads, and library workspace; it is not measured HBM traffic. Caches affect how much reaches HBM.

These are potential limits, not measured bottlenecks. In particular, the selected Inductor and TinyGrad runs have no matching Nsight resource capture, so we cannot infer their achieved occupancy or bandwidth from source alone.

What the Four Compilers Teach Us

The numerical contract fixes the required behavior while leaving room for different implementations. Our FP16 projection must round before bias addition, but that rounding can happen inside a fused kernel or between separate kernels. A semantic boundary does not require a memory buffer or a kernel launch.

Compiler design shapes the available choices. Templates, rewrite rules, library interfaces, and authored kernels determine which implementations can be constructed; heuristics and measured search select among them. The pinned Luminal backend makes this distinction concrete: searching more plans cannot select an FP16-input matrix implementation absent from its backend rules. Improving selection and expanding the available implementations are separate engineering tasks.

The practical question is which complete implementation preserves the contract while balancing data movement, matrix throughput, and parallelism. Fusion can remove intermediate storage and launches, but its value depends on the resulting kernel's schedule. Agentic kernel generation could expand the available implementations; integrating a candidate still requires preserving its surrounding dependencies and evaluating the complete execution path.

The experiment report retains the pinned versions, configurations, and evidence for this linear layer on H100.