one line to machine code

I wrote the most boring line of tensor code I could and asked PyTorch’s compiler to show its work:

def f(x):
    return torch.relu(x @ x.T + 1.0)

A matrix times its own transpose, then one is added to every result, then every negative result is replaced by zero. That last operation is a rectified linear unit, written relu. In ordinary PyTorch, Python asks for those operations in sequence. I passed the function to torch.compile, ran it once on a square matrix of 64 rows and 64 columns, and opened the generated wrapper:

def call(self, args):
    arg0_1, = args
    assert_size_stride(arg0_1, (64, 64), (64, 1))
    buf0 = empty_strided_cpu((64, 64), (64, 1), torch.float32)
    extern_kernels.mm(arg0_1, reinterpret_tensor(arg0_1, (64, 64), (1, 64), 0), out=buf0)
    buf1 = buf0; del buf0  # reuse
    cpp_fused_add_relu_0(buf1)
    return (buf1, )

Four details did not fit the simple word “compile.” The matrix multiplication became extern_kernels.mm, which is a call through an Inductor lowering rather than a loop visible in this wrapper. The transpose became a different set of strides passed to reinterpret_tensor; no copy appears. The add and relu became generated C++. The wrapper also rejects any input that is not exactly $64 \times 64$ with row-major strides.

The first correction arrived before the first benchmark. I initially described the matrix multiplication as a cut in the graph because its implementation was external. The trace contains one graph and zero breaks. The library call and the generated loop are two implementations scheduled inside one compiled region. “External” identifies where the operation runs. It does not mean Python regained control.

The measurements use Python 3.10.14 and the CPU build of PyTorch 2.9.1 at git revision 5811a8d7da873dd699ff6687092c225caffcf1bb on arm64 macOS 15.7.7. PyTorch reports Apple clang 15 in its build configuration. The local clang --version reports the same major version. CUDA is absent, MPS is present, and this investigation deliberately stays on the CPU. Timed runs set both OMP_NUM_THREADS=1 and torch.set_num_threads(1) because the generated four-thread loop crashes in this environment. That failure appears later, with the one-thread run beside it.

The git revision matters more than the package label. 2.9.1 identifies a release, but implementation claims in this investigation name source functions, default flags, and generated formats that can change between patches. The probe asserts both:

assert torch.__version__.split("+")[0] == "2.9.1"
assert torch.version.git_version == (
    "5811a8d7da873dd699ff6687092c225caffcf1bb"
)

It then reads the installed source through Python’s own module paths. That avoids accidentally citing a current web page while executing an older local package. TensorVariable.method_item, torch._dynamo.config, torch._inductor.kernel.mm, CppScheduling, and compile_fx.py all come from the imported installation.

The generated cache is isolated for probes that inspect artifacts:

export TORCHINDUCTOR_CACHE_DIR="$(mktemp -d /private/tmp/asquare-compile.XXXXXX)"
export OMP_NUM_THREADS=1
python pipeline_probe.py

A fresh directory separates cold generation from artifacts left by another function, shape, or compiler setting. The scripts locate files by content as well as suffix. A random .cpp file in an Inductor cache might be a support extension rather than the add and relu loop. The pipeline probe requires clamp_min, the vector load, and the matching wrapper call before treating a file as evidence.

Random seeds make tensor values repeatable, but they do not make timings deterministic. The operating system can schedule other work, frequency can change, and cache state can vary. The benchmark therefore pins input generation and trial order while reporting a distribution. It does not claim bit-for-bit repeatable microseconds.

Correctness checks precede timing:

torch.testing.assert_close(eager(x), compiled(x))

This is essential because a crashed process, a stale cache artifact, a dtype mismatch, or an invalid optimization can be fast. Performance is measured only after both paths agree under the comparison contract. Exceptional floating-point values receive a separate check because random normal inputs rarely contain NaN or infinity.

The source files and outputs live under writing-notes/probes/compile. The article shortens generated code for readability, while the probes assert the omitted context. This division keeps a claim reproducible without turning the public page into a dump of cache paths and framework logging.

three calls that look like one line

Default PyTorch uses eager execution: an operation is dispatched when Python reaches it. x @ x.T produces a tensor. + 1.0 consumes that tensor and produces another logical result. relu consumes the result of the add. On this CPU build these are operator calls, not GPU kernel launches. The distinction matters because a CPU library call can execute synchronously on the calling thread, invoke a thread pool, or call another library. Calling every operation a launch would import a GPU cost model that this experiment did not measure.

The separate add and relu require separate operator boundaries and a logical intermediate between them. Their algorithmic tensor traffic is easy to count. For $N$ float32 elements, a simple out-of-place add reads $4N$ bytes and writes $4N$ bytes. A following out-of-place relu does the same. That gives $16N$ bytes across the operator interfaces. A fused in-place loop can read and write the matrix once, $8N$ bytes. These are useful lower bounds on traffic implied by the algorithms. They are not measurements of DRAM traffic. A small tensor can remain in cache, stores can be combined, and the hardware can move cache lines for reasons the source does not show.

The arithmetic ratio suggests a prediction. Adding one and comparing against zero require only a few operations for each element moved. Matrix multiplication reuses rows and columns to perform many multiply-add operations per output element. I therefore predicted that fusion would matter more for the elementwise tail than for the complete expression. Calling either path memory-bound or compute-bound would require counters or a roofline measurement. The operation counts make the prediction; the timings later test it.

torch.compile can see the add feeding relu inside one captured region and emit one loop for both. It cannot make the matrix product disappear, and it does not promise that the result avoids every cache or memory access. The first useful question is narrower: what program did the compiler capture?

python stops at a frame boundary

Dynamo receives a Python frame before the interpreter executes that frame normally. A frame holds one invocation of a Python function: its bytecode, local variables, instruction position, and links to the surrounding execution state. CPython exposes a hook for replacing frame evaluation through PEP 523. The first line of the pinned torch._dynamo package says that Dynamo uses this API. The PyTorch 2.9 compiler overview describes the same entry point.

The distinction fixed another bad sentence in my first explanation. Dynamo does not watch tensor calls while throwing away the Python around them. It interprets Python bytecode far enough to decide which values can be represented in a graph, records a region, installs guards for assumptions made while tracing, and leaves residual Python to coordinate compiled regions when capture cannot continue.

The bytecode for the function is small enough to inspect with the standard library:

import dis

dis.dis(f)

The exact instruction names depend on the Python version. Python 3.10.14 shows the same logical sequence as the source: load torch.relu, load x twice, fetch attribute T, perform matrix multiplication, load constant 1.0, add, call relu, and return. None of those instructions says ATen, vector width, or C++. They describe operations in the Python virtual machine.

Dynamo’s first graph stays close to those Python operations:

graph():
    %x       = placeholder[target=L_x_]
    %t       = call_function[target=builtins.getattr](args=(%x, T))
    %matmul  = call_function[target=operator.matmul](args=(%x, %t))
    %add     = call_function[target=operator.add](args=(%matmul, 1.0))
    %relu    = call_function[target=torch.relu](args=(%add,))
    return (relu,)

This is an FX graph. Each line is a node. The percent name is a label for the value produced by that node. The arguments name earlier nodes, so the edges are implicit in the text. matmul consumes x and t; add consumes matmul; relu consumes add. There is no loop over matrix elements here. The graph describes relationships between tensor operations.

The output reports one graph and zero graph breaks. That count matters because it settles the question raised by extern_kernels.mm: Dynamo captured the complete expression before Inductor chose different implementations for its parts.

Another trace, performed with make_fx at the ATen level, produced a graph farther from Python:

=== traced ATen graph ===
  permute   = aten.permute.default
  mm        = aten.mm.default
  add       = aten.add.Tensor
  relu      = aten.relu.default

The two graphs must not be confused. The Dynamo graph retained operator.matmul and the Python attribute access for .T. The ATen trace contains aten.permute.default and aten.mm.default. ATen is PyTorch’s operator library and dispatch layer. The friendly expression has now been expressed in operator forms that later compiler stages know how to lower.

x.T became permute. The @ operator became mm. Adding the Python number 1.0 became aten.add.Tensor. torch.relu became aten.relu.default. Decomposition is the process of replacing a larger or friendlier operation with lower-level operators. It does not happen once at a universal boundary. Different tracing modes, decompositions, devices, and compiler passes can expose different amounts of an operation.

the operator name still does not name an implementation

aten.mm.default looks specific, but it identifies an operator schema and overload rather than one final body of machine code. The schema describes arguments and result behavior. PyTorch’s dispatcher uses properties such as device and dispatch keys to select an implementation for an eager call. CPU, CUDA, MPS, autograd, functionalization, and other modes can route the same named operation differently.

Compilation introduces another mapping. Inductor registers a lowering for the ATen operator. That lowering can retain a call to an external implementation, generate code, decompose the operation further, or build a choice set. The clean expression therefore crosses two distinct selection questions:

What tensor operation is this?
    aten.mm.default

How will this compiled backend implement it here?
    extern_kernels.mm through an ATen GEMM choice

The word default is the overload name. It does not mean a generic slow fallback. Operators can have overloads because Python-facing syntax accepts different combinations of tensors, scalars, dimensions, and keyword arguments. aten.add.Tensor distinguishes tensor addition from other add schemas. Those overload identities let tracing record a precise contract after Python’s flexible call rules have been resolved.

The scalar 1.0 illustrates type promotion. The source mixes a float32 tensor with a Python floating-point value. Eager PyTorch applies its promotion rules and produces a float32 result for this case. The ATen graph records a tensor add with the scalar argument. The generated C++ constructs Vectorized<float>(1.0), so the embedded constant has become the loop’s element type.

Changing the input to float64 changed generated vector types and produced a separate guarded graph. Changing it to float16 produced widening conversions in the CPU artifact. The same Python literal participates in different generated instructions because dtype resolution precedes code generation.

Layout is also part of implementation selection. aten.permute.default describes a logical permutation. The downstream mm lowering observes strides and can pass a transposed view to a library. If a selected library path required contiguous input, the compiler could insert a copy or choose a different implementation. No copy appears in this artifact, so the chosen path accepted the strided operand.

Dispatch continues inside external libraries. An ATen CPU matrix multiplication may reach an optimized BLAS routine, and that library can select microkernels based on shape, dtype, and CPU features. The Python wrapper does not expose that internal selection. Naming extern_kernels.mm is therefore one boundary reached, not the final instruction sequence of the matrix multiplication.

This layered dispatch makes broad source claims fragile. “PyTorch matrix multiplication uses algorithm X” omits device, dtype, layout, version, build configuration, and shape. A defensible claim names the observed path: in the pinned CPU build, for a contiguous $64 \times 64$ float32 input and its transposed view, Inductor emitted extern_kernels.mm with those exact size and stride arguments.

The distinction also explains why two operators can remain in one graph while using different execution engines. A graph is a dependency representation. It need not mean one generated function. The matrix multiplication can call a library, the pointwise tail can call generated C++, and the wrapper can connect their buffers. Fusion decides which compatible operations share generated iteration. It does not require every operation in the captured region to have the same origin.

The reverse mistake is possible too. One ATen node can lower into several loops or calls. _softmax remained one node in the earlier trace but needed several phases of reduction and normalization. Counting graph nodes is therefore not a performance model. To estimate work, the node must be followed into its lowering, iteration domains, buffers, and library calls.

I wanted to see that decomposition happen to something bigger than my toy, so I traced a handful of the operations a real network is actually made of and counted the ATen nodes each one collapsed to:

relu(x@x.T+1):  4 nodes:  permute, mm, add, relu
softmax:        1 node:   _softmax
gelu:           1 node:   gelu
layer_norm:     4 nodes:  native_layer_norm, getitem x3
sdpa (attention): 6 nodes: mul x2, transpose, mm x2, _safe_softmax

This result contradicted my first description of softmax. The arithmetic for softmax needs a row maximum, subtraction, exponentiation, a row sum, and division. The ATen graph above still represents it as _softmax. Those arithmetic steps appear only after lowering. An operator can therefore be primitive for one stage and composite for the next.

I compiled torch.softmax separately and opened cpp_fused__softmax_0. The maximum, exponentiation, and accumulation appeared in the generated C++:

tmp_acc0_vec = at::vec::maximum(tmp_acc0_vec, tmp0);   // the max
// ... reduce the max across the vector ...
auto tmp4 = tmp3.exp();                                // the exp (after subtract)
tmp_acc0 = tmp_acc0 + at::vec::vec_reduce_all(..., tmp_acc0_vec);  // the sum

The single generated function does not prove a single pass over the row. A stable softmax must know the row maximum before it can accumulate exponentials shifted by that maximum, and it must know the exponential sum before it can normalize outputs. The generated function can contain several loop nests and temporary values while still being one callable unit. Function count, operator count, loop count, and memory-pass count are four different measurements. The code established the arithmetic and fusion boundary. I did not reduce it to a one-pass claim.

The transpose can be settled directly from the wrapper. The input has shape (64, 64) and strides (64, 1). A stride says how many elements the address advances when an index increases by one. Moving to the next row advances 64 floats; moving to the next column advances one. The transposed view has strides (1, 64). Its first index walks through adjacent elements and its second index jumps by 64. Both descriptions refer to the same storage and the same starting offset.

reinterpret_tensor passes that changed description to mm. No transpose allocation or copy appears between the input and the external call. “Free” would still be too broad because the matrix multiplication must consume a strided operand, and that layout can affect the library path. The precise observation is that .T did not materialize a second tensor in this generated region.

the buffer name was not an intermediate representation

My first reading treated the generated wrapper as if it were Inductor’s internal representation. That was attractive because the names look compiler-like:

buf0 = empty_strided_cpu((64, 64), (64, 1), torch.float32)
extern_kernels.mm(arg0_1, ..., out=buf0)
buf1 = buf0; del buf0  # reuse
cpp_fused_add_relu_0(buf1)

buf0 is assigned once. buf1 is assigned once. That resemblance is not enough to call the wrapper static single assignment form. Static single assignment, usually shortened to SSA, means every value in an intermediate representation has exactly one definition. It also has rules for values that meet after control-flow branches. A Python wrapper whose local variables happen to be assigned once does not expose those rules or establish the form of an earlier scheduler IR.

The wrapper proves something more concrete. empty_strided_cpu obtains storage for the matrix product. extern_kernels.mm writes into it. buf1 = buf0 gives the same Python object another local name. del buf0 removes the earlier reference. The generated function cpp_fused_add_relu_0 receives buf1 as an in-place argument. The output returned to the caller is that same buffer. No separate output allocation for the add and no separate output allocation for relu appear in the wrapper.

That is a lifetime decision. Once the matrix product has been consumed by the fused tail, the unfused value is no longer observable. Overwriting its storage therefore preserves the function’s result. If some second consumer needed the raw matrix product, this reuse might become illegal. The compiler needs use information and alias information to make that decision, but the wrapper alone does not say which internal pass proved it.

I made the second consumer observable by returning both values:

def keep_product(x):
    product = x @ x.T
    return product, torch.relu(product + 1.0)

The original single-output function generated:

buf0 = empty_strided_cpu((64, 64), (64, 1), torch.float32)
extern_kernels.mm(..., out=buf0)
buf1 = buf0; del buf0  # reuse
cpp_fused_add_relu_0(buf1)
return (buf1,)

The two-output function generated a different contract:

buf0 = empty_strided_cpu((64, 64), (64, 1), torch.float32)
extern_kernels.mm(..., out=buf0)
buf1 = empty_strided_cpu((64, 64), (64, 1), torch.float32)
cpp_fused_add_relu_0(buf0, buf1)
return (buf0, buf1)

The fused function signature changed from one float* to const float* plus float*. buf0 remains readable and unmodified because the caller can observe it as the first returned tensor. buf1 receives the transformed result. One extra consumer changed allocation, pointer constness, and store destination without changing the arithmetic of the second output.

This is liveness in a directly visible form. A value is live while some future operation or returned result may need it. In the first function, the product dies when the fused tail begins because nothing can observe the original afterward. In the second, it remains live through function return. Reusing the buffer would replace the first returned value with the transformed one and violate Python semantics.

Aliasing makes liveness harder. Two tensor objects can name overlapping storage through views. Writing through one can change what another observes. The compiler must distinguish a new allocation, a view with different shape or stride, an in-place mutation, and a copy. reinterpret_tensor for the transpose is a view description. empty_strided_cpu creates output storage. buf1 = buf0 aliases the same tensor object at the wrapper level. cpp_fused_add_relu_0(buf1) mutates that storage.

The original input does not alias the product output in this expression. Matrix multiplication writes a new matrix. That fact makes it safe to return a result whose storage has overwritten the product. A different expression involving a view of the input would need more care. Functionalization and alias analysis exist because source syntax alone does not reveal every shared-storage relationship.

The extra allocation also changes the traffic model. In the one-output function, the fused tail loads and stores buf0. In the two-output function, it loads buf0 and stores buf1, preserving the source. Both perform one explicit load and one store per element inside the fused loop. The two-output version additionally retains two full matrices at return instead of one. Peak live memory changes even when the fused arithmetic traffic does not.

That distinction matters in training graphs, where a forward value may appear dead to the Python return but still be required by backward. AOTAutograd can turn such a value into an additional hidden forward output saved for the gradient program. The later backward probe shows exactly that transformation.

The path between the two printed graphs and this wrapper has several layers. They matter whenever a claim is attached to “the graph” without naming which one:

CPython frame and bytecode
        |
        v
Dynamo FX graph, plus guards and residual Python
        |
        v
AOTAutograd and functionalization
        |
        v
ATen operations and decompositions
        |
        v
Inductor graph lowering and scheduler decisions
        |
        +---- external operation choices such as aten.mm
        |
        +---- generated loop code such as add plus relu
        |
        v
Python wrapper, generated C++, shared library
        |
        v
clang and LLVM machine code

AOTAutograd sits in the middle even when the original function only performs a forward calculation. Its name becomes literal when gradients are needed: it captures and partitions forward and backward work so a backend can compile both. Functionalization rewrites mutations and aliasing into forms that later stages can reason about. Inductor then lowers the resulting graph to its own dependency and scheduling structures. The pinned compile_fx.py imports aot_autograd, defines both forward and backward compilers, and sends graph modules into Inductor. The PyTorch 2.9 overview describes Dynamo, AOTAutograd, and Inductor as distinct components rather than one tracer.

The distinction predicts where an observation belongs. A graph break belongs to Dynamo capture. A saved tensor needed by a backward formula belongs to the AOTAutograd boundary. Selecting an ATen matrix multiplication choice belongs to an Inductor lowering. Splitting a loop for vector width belongs to CPU code generation. A guard failure happens before cached code can be reused. Calling all of these compilation loses the causal path.

The source for the pinned revision makes the matrix multiplication decision inspectable. In torch/_inductor/kernel/mm.py, tuned_mm is registered as the lowering for aten.mm. It builds a list of choices. When ATen GEMM kernels are enabled, an aten_mm choice is added. Triton and other choices are conditional on device and configuration. On this CPU run, the emitted wrapper invokes the ATen external kernel.

This does not mean mm escaped compilation. The lowering was selected while compiling the same region. The wrapper sequences the library operation and the generated loop without returning to the original Python function. A graph break would produce separate captured regions with Python between them. The clean trace reports none.

the two operations enter one loop

The generated C++ for cpp_fused_add_relu_0, trimmed to its arithmetic, is:

extern "C" void kernel(float* in_out_ptr0) {
    #pragma omp parallel num_threads(4)
    for (int64_t x0 = 0; x0 < 4096; x0 += 4) {
        auto tmp0 = at::vec::Vectorized<float>::loadu(in_out_ptr0 + x0, 4);
        auto tmp2 = at::vec::Vectorized<float>(1.0);
        auto tmp3 = tmp0 + tmp2;                              // the + 1.0
        auto tmp4 = at::vec::clamp_min(tmp3, decltype(tmp3)(0));  // the relu
        tmp4.store(in_out_ptr0 + x0);
    }
}

The loaded vector becomes tmp0. The constant vector is tmp2. Addition produces tmp3. clamp_min produces tmp4, which is stored through the same pointer. No store of tmp3 appears between the add and clamp. This is the evidence for fusion: both operations occur inside one loop iteration, and their intermediate exists as a C++ local value rather than as a separately allocated tensor.

For 4,096 float32 elements, the generated loop performs 1,024 vector iterations. At the operator-interface level, the unfused lower bound was 65,536 bytes for two reads and two writes. The fused loop explicitly loads and stores 16,384 bytes each, for 32,768 bytes at this level. Cache lines and write policy can change physical traffic. The generated program still establishes why the fused path has half the explicit element traffic.

The four-thread version shown above was generated with OMP_NUM_THREADS=4. The controlled timing version generated with one thread omits that parallel team. This matters because thread creation, wakeup, scheduling, and work partitioning can exceed the arithmetic cost for 4,096 elements. It also prevents the crash from contaminating the benchmark.

four lanes are visible, four fast cores are not

The loop increments by four elements and asks Vectorized<float>::loadu for four floats. Four float32 values occupy 128 bits. Arm64 NEON vector registers are 128 bits wide, so the source is consistent with one vector register per Vectorized<float> value. Generated C++ is still above final assembly. The compiler may schedule several vector instructions, unroll, or use scalar cleanup. The statement supported here is that Inductor chose four float lanes for this artifact.

The number in num_threads(4) supports less than I first claimed. The environment reported torch.get_num_threads() == 4, and the generated OpenMP pragma requested four threads. I found no trace proving why PyTorch had that default, which cores the operating system assigned, or whether the threads ran only on performance cores. Heterogeneous core scheduling is an operating system and runtime decision unless affinity is explicitly controlled. The probe did not set affinity, and macOS does not expose Linux taskset or NUMA controls.

I wanted to know how load-bearing that width-4 choice really was, so I changed the dtype and re-read the kernel. At float16 and bfloat16, the loop steps by 8, not 4, and the load pulls 8 elements: a 16-bit type is half the size, so twice as many fit in the same 128-bit register. But then something I didn’t expect showed up in the half-precision kernel:

auto tmp0 = at::vec::Vectorized<at::Half>::loadu(in_out_ptr0 + x0, 8);
auto tmp1 = at::vec::convert<float,2,at::Half,1>(tmp0);   // widen 8 halfs -> two float vectors
// ... the +1.0 and clamp happen in float32 ...
auto tmp6 = at::vec::convert<at::Half,1,float,2>(tmp5);   // narrow back to half

For this generated CPU artifact, eight half values are loaded, converted to float32 vectors, processed, and converted back. That does not establish a rule for every arm64 PyTorch operation, every compiler version, or MPS execution. It establishes the chosen path for this expression in the pinned CPU build. bfloat16 generated the same widening shape. The float64 artifact used VectorizedN<double,2>, grouping two native-width vectors into one logical object.

Storage width, arithmetic width, and logical loop width can therefore differ. That difference becomes a possible cost rather than an automatic advantage. Narrow storage halves the bytes per element, but conversions add instructions and wider temporaries consume registers.

the fast loop still owes the same answer

Fusion is legal only if the generated operation preserves the observable numerical contract closely enough for the selected compiler settings. Ordinary finite values are the easy part. Infinity, NaN, signed zero, overflow, and rounding expose shortcuts that can change meaning.

I isolated the generated tail and supplied eight float32 values:

values = torch.tensor([
    -math.inf,
    -2.0,
    -1.0,
    -0.0,
    0.0,
    1.0,
    math.inf,
    math.nan,
])

Both eager and compiled execution returned:

input:    -inf  -2.0  -1.0  -0.0  0.0  1.0  inf  nan
output:    0.0   0.0   0.0   1.0  1.0  2.0  inf  nan

The result follows the order of operations. Addition happens before relu. Negative infinity plus one remains negative infinity, then clamps to zero. Negative one plus one becomes zero. Negative zero plus one becomes positive one. Positive infinity remains infinite. NaN remains NaN. torch.testing.assert_close checked the arrays with equal_nan=True.

NaN is a useful control because different maximum instructions have different rules. Some return the numeric operand when the other operand is NaN; others propagate NaN. A compiler that replaced PyTorch relu with an incompatible maximum could turn an invalid value into zero and hide the upstream fault. The generated C++ calls at::vec::clamp_min, and this artifact matches eager output on the tested NaN.

This eight-element test does not exhaust floating-point behavior. It omits signaling NaNs, different NaN payloads, subnormal values, every rounding boundary, and dtype-specific behavior. It verifies the exceptional cases most likely to be broken by a careless max rewrite.

Reassociation creates another hazard. In exact arithmetic:

$$(a + b) + c = a + (b + c)$$

In floating-point arithmetic the equality can fail because each addition rounds. Let $a = 10^{20}$, $b = -10^{20}$, and $c = 3.14$ in a format where $10^{20}$ is representable but small additions near it are lost. The left grouping first produces zero, then $3.14$. The right grouping first rounds $-10^{20} + 3.14$ back near $-10^{20}$, then produces zero. A fusion pass that changes association can change the answer.

The current expression has a fixed add followed by a clamp, so it offers no legal reassociation among several additions. Matrix multiplication reductions do contain long sums, and different kernels can accumulate them in different orders. That is one reason correctness comparisons use tolerances instead of bit equality for many tensor operations.

Compiler modes and flags govern how aggressive those transformations may be. Fast-math settings can permit assumptions about NaNs, infinities, signed zeros, and association that ordinary IEEE reasoning would reject. I did not enable or inspect a fast-math mode for the generated extension. The edge-case probe checks output behavior for the selected defaults rather than making a claim about every Inductor option.

Integer operations have a different set of contracts. Overflow can wrap, saturate, raise, or be undefined depending on language and operation. Division by zero and shift widths also matter. Shape arithmetic often uses signed 64-bit integers, as const int64_t ks0 shows. Proving that ks0 * ks0 cannot overflow belongs to guard and shape constraints, not to vector floating-point code.

Correctness also includes layout and mutation. Equal numeric values in a new tensor are insufficient if eager execution returns a view, preserves aliasing, updates an input, or raises on an invalid overlap. A tensor compiler must preserve those observable properties or document a mode that changes them. The one-line expression is mercifully simple: it returns a new result and performs no user-visible input mutation.

The probe therefore establishes a bounded contract: for the finite values and exceptional float32 values shown, the one-thread generated CPU tail agrees with eager PyTorch 2.9.1. That evidence earns the timing comparison. A fast answer that changed NaN behavior would be a different program.

the first number was already wrong

I first timed eager and compiled execution in two large uninterrupted batches. That made the result depend on order, thermal state, and whatever changed between batches. The replacement benchmark alternates labels in a seeded random order. Each label receives 31 samples after ten warmup calls. A sample contains repeated inner calls, with fewer repetitions for larger matrices. Garbage collection is disabled during the timed region. Every compiled output is checked against eager output before timing.

The method does not control CPU affinity, frequency, or background macOS activity. The intervals therefore report the 10th percentile, median, and 90th percentile instead of pretending the median is exact:

full_n=64
  eager:    p10 5.69 us   median 5.96 us   p90 6.15 us
  compiled: p10 21.02 us  median 23.59 us  p90 24.16 us
  speedup at medians: 0.253x

full_n=512
  eager:    p10 437.73 us  median 493.24 us  p90 516.09 us
  compiled: p10 421.86 us  median 447.01 us  p90 457.35 us
  speedup at medians: 1.103x

full_n=1024
  eager:    p10 3092.88 us  median 3238.39 us  p90 3730.76 us
  compiled: p10 2736.86 us  median 2857.88 us  p90 3111.42 us
  speedup at medians: 1.133x

At 64 rows, compilation loses by almost a factor of four. The generated loop has fewer explicit element accesses, yet the complete call still checks guards, enters a wrapper, prepares the external matrix multiplication, and invokes a generated function. Fixed costs do not shrink with the tensor. At 4,096 elements, there is too little avoided work to pay those costs.

At 512 and 1,024 rows, the compiled expression wins by 10 and 13 percent at the medians. The matrix multiplication performs work proportional to $n^3$. The elementwise tail touches $n^2$ values. As $n$ grows, the matrix multiplication occupies more of the full expression, and both paths rely on a library for that operation. The compiler can improve the tail without changing the dominant cubic term.

I isolated the part that code generation actually changed:

def tail(x):
    return torch.relu(x + 1.0)

The same randomized method produced:

tail_n=512
  eager:    median 52.31 us
  compiled: median 40.74 us
  speedup: 1.284x

tail_n=1024
  eager:    median 461.11 us
  compiled: median 245.46 us
  speedup: 1.879x

tail_n=2048
  eager:    median 2548.48 us
  compiled: median 1291.60 us
  speedup: 1.973x

The approach to 2x matches the explicit traffic ratio: two operator-level reads and writes compared with one generated load and store. It is supporting evidence, not a bandwidth measurement. I did not collect hardware counters for cache misses, memory-controller traffic, or achieved bandwidth. The result is also specific to one thread. The safe conclusion is that removing the intermediate explains the scale and direction of the timing. Calling the kernel DRAM-bound would go farther than the evidence.

A later run of the complete verifier produced different medians:

full_n=64:    eager 6.05 us     compiled 21.00 us    speedup 0.288x
full_n=512:   eager 459.89 us   compiled 417.82 us   speedup 1.101x
full_n=1024:  eager 2725.11 us  compiled 2439.61 us  speedup 1.117x

tail_n=512:   eager 59.65 us    compiled 39.57 us    speedup 1.507x
tail_n=1024:  eager 417.32 us   compiled 232.26 us   speedup 1.797x
tail_n=2048:  eager 2113.60 us  compiled 1063.85 us  speedup 1.987x

The absolute times and the 512-row tail ratio moved. The three conclusions did not: compiled loses badly at 64 rows, wins modestly for the full larger expressions, and approaches 2x for the largest isolated tail. The change between runs is evidence that one invocation of this laptop benchmark cannot support a tight performance promise. The byte calculation below uses the first recorded run and says so through its input medians.

the byte model changes when the cache changes

The traffic lower bound can be divided by time to form an algorithmic effective bandwidth. This is a rate implied by the bytes visible in the operator model. It is not physical bus bandwidth.

For the eager tail with $N$ float32 elements:

$$B_{\text{eager}} = 4\ \text{interfaces} \times 4N\ \text{bytes} = 16N\ \text{bytes}$$

The four interfaces are the add read, add write, relu read, and relu write. For the generated tail:

$$B_{\text{compiled}} = 2\ \text{interfaces} \times 4N\ \text{bytes} = 8N\ \text{bytes}$$

At $n=2048$, $N=n^2=4{,}194{,}304$ elements. The eager model contains 64 MiB and the compiled model contains 32 MiB. Dividing by the medians:

$$R_{\text{eager}} = \frac{64\ \text{MiB}}{2.54848\ \text{ms}} = 24.56\ \text{GiB/s}$$ $$R_{\text{compiled}} = \frac{32\ \text{MiB}}{1.29160\ \text{ms}} = 24.20\ \text{GiB/s}$$

The two algorithmic rates differ by about 1.5 percent. That agreement is what the near-2x speedup means in this model. Each path appears limited by a similar effective supply rate while the compiled path asks for half the explicit bytes.

At $n=1024$, the same calculation gives about 33.9 GiB/s for eager and 31.8 GiB/s for compiled. They remain close, though less so. At $n=512$, eager implies about 74.7 GiB/s while compiled implies about 47.9 GiB/s. The simple constant-rate model breaks.

That break is expected from a memory hierarchy. A $512 \times 512$ float32 tensor contains 1 MiB. The working sets and intermediates can interact with private and shared caches differently from a 16 MiB tensor at 2,048 rows. The rate of moving bytes through a nearby cache can exceed the rate sustained by a larger working set. Fixed dispatch and loop overhead also occupy a larger fraction of the 40 to 52 microsecond samples.

The model can include a fixed term:

$$T(N) = F + \frac{B(N)}{R(N)}$$

$F$ is call, guard, wrapper, and loop setup measured in seconds. $B(N)$ is algorithmic bytes. $R(N)$ is an effective rate that can change when the working set crosses cache levels. With only three sizes, fitting both $F$ and a size-dependent $R$ would produce parameters rather than understanding. Hardware counters and a denser size sweep are needed to locate cache transitions.

The full expression adds another term:

$$T_{\text{full}}(n) = T_{\text{mm}}(n) + T_{\text{tail}}(n) + T_{\text{boundary}}(n)$$

This equation is conceptual. The terms can overlap through caches and library behavior, and separate medians cannot be added as if they came from paired calls. It still identifies why the full speedup shrinks: $T_{\text{mm}}$ grows with roughly cubic arithmetic while the tail’s element count grows quadratically.

For a square matrix product, every one of $n^2$ outputs accumulates $n$ products. Counting a multiply and an add as two floating-point operations gives approximately:

$$O_{\text{mm}} = 2n^3\ \text{operations}$$

At $n=1024$, that is about $2.15 \times 10^9$ operations. The tail performs only a few million elementwise operations. Saving one tail intermediate can matter, but it cannot transform the dominant operation count.

Arithmetic intensity divides operations by bytes. A well-tiled matrix multiplication reuses input blocks so each loaded value participates in many operations. The tail performs one add and one comparison for each load and store. That contrast predicts why a specialized GEMM library owns the matrix product while fusion is valuable around it.

A proper roofline placement would need achieved floating-point rate and measured memory traffic or at least a defensible bandwidth ceiling for the same core configuration. I collected neither. The calculations above remain source-level traffic and operation models. They explain the direction of results while preserving the difference between a model and a counter.

The distinction prevents an easy but invalid conclusion. Two implied rates near 24 GiB/s do not prove the physical memory system delivered 24 GiB/s. Cache-line fills, write allocation, writeback, prefetching, and reuse can make physical traffic larger or smaller than the explicit element interfaces. The rates are valuable because both programs are evaluated with the same accounting rule, not because the number names a hardware link.

the first call is three different events

A cold first call, a warm call in the same process, and a first call from a second process using the disk cache are different experiments. I gave a child process a new TORCHINDUCTOR_CACHE_DIR, timed its first and second call, then started a fresh child with the same cache:

cold process:
  first call       7988.02 ms
  second call        64.17 us
  cache files             9

fresh process, same disk cache:
  first call       4498.84 ms
  second call        57.25 us
  cache files             9

The first number includes Python startup work after timing begins inside the child only where it affects lazy imports, Dynamo capture, guard construction, AOTAutograd work, Inductor lowering, C++ generation, clang, loading, and execution. It is not “compiler time” in isolation. The second number is a same-process guarded call to already loaded code. The fresh-process number shows that finding the same nine files on disk does not turn the first invocation into the 57 microsecond warm path. Python still reconstructs compiler state, checks cache entries, loads code, and runs initialization.

One run does not provide a distribution for compile latency. These values are labels on observable paths, not stable performance estimates. They correct the earlier claim that a warm disk cache makes compilation a dictionary lookup and permanently removes first-call cost. Several caches participate, and process-local state still matters.

An application that serves a request immediately after startup has to account for the cold path. One that compiles representative shapes before accepting traffic can put that cost outside request latency. A service that creates new worker processes frequently can pay a middle path even when artifacts persist on disk. “Compilation overhead is amortized” only becomes true after naming the number of calls, process lifetime, shape set, and cache state.

four threads end the process

The larger shape initially failed before it could be timed. I reduced the failure to a subprocess so a crash could not kill the verifier. The control and failing cases use separate fresh caches:

OMP_NUM_THREADS=1
torch threads=1
first run OK, shape (128, 128)
second run OK

OMP_NUM_THREADS=4
torch threads=4
about to trigger compile+run
Fatal Python error: Segmentation fault
four_thread_exit=139

Exit 139 is the shell representation of termination by SIGSEGV. The failure reproduces at the first compiled call with four threads. The one-thread process compiles, checks the output shape, and executes the already compiled function a second time.

The generated four-thread C++ contains:

#pragma omp parallel num_threads(4)
{
    #pragma omp for
    for (int64_t x0 = 0; x0 < 16384; x0 += 4) {
        // vector add and clamp
    }
}

otool -L on a generated shared object lists @rpath/libomp.dylib, libc++, and the system library. That list does not prove that only one OpenMP runtime exists anywhere in the full process, nor does it locate the invalid access. Python’s fault handler showed an asynchronous compilation worker thread, but a Python traceback cannot unwind arbitrary native worker state into a root cause.

I tested KMP_DUPLICATE_LIB_OK=TRUE during the earlier investigation and it did not repair the failure. That negative result rejects one workaround; it does not reject every duplicate-runtime mechanism. AddressSanitizer, an LLDB native backtrace, a minimized generated extension independent of PyTorch, and runtime version inspection would be reasonable next probes. They were not completed. The cause remains unknown.

Every performance number above therefore forces one thread. It measures fusion and generated vector code without measuring the intended OpenMP parallel execution. A working four-thread build could change both eager and compiled results because library matrix multiplication also has threading choices. The failure is part of the environment, not evidence that Inductor’s CPU threading generally crashes.

the matrix multiplication stayed inside the region

Now the other half of the hook, the matmul that got shipped out:

extern_kernels.mm(arg0_1, reinterpret_tensor(...), out=buf0)

The matrix multiplication did not become a generated pointwise loop. It became an extern_kernels.mm call with an output buffer supplied by the wrapper. GEMM means general matrix multiplication. Optimized GEMM implementations usually contain packing, cache tiling, register blocking, vector instructions, and shape-dependent choices. Reusing one can be better than emitting a direct triple loop.

The pinned lowering does more than blindly refuse code generation. tuned_mm creates candidate choices according to device and configuration. The installed source describes ATen, Triton, and CUTLASS possibilities, though this CPU run emitted only the external ATen path. With default settings, no autotuning table appeared. That absence does not prove a one-candidate race because logging conditions and choice filtering also affect the table. It only records that this run provided no measured comparison among GEMM candidates.

The full benchmark supports the division of labor. The isolated tail reaches nearly 2x at 2,048 rows, while the full function improves by 13 percent at 1,024 rows. The full expression contains matrix multiplication work that the tail experiment removes. I did not time mm alone in this probe, so I cannot subtract the two medians and assign every remaining microsecond to GEMM. Medians of separate distributions are not additive. The scaling and generated wrapper still explain why a large tail speedup becomes a modest whole-function speedup.

the first guard knows 4,096 elements

One line in the generated call I haven’t explained is the first real statement:

assert_size_stride(arg0_1, (64, 64), (64, 1))

This assertion checks the shape and stride expected by this wrapper. Dynamo guards cover more than this one line: tensor type, device, dtype, Python object identities where relevant, global state, and relationships among symbolic values can all participate. A compiled entry is eligible only when its guards match the current call.

Exact sizes let the generated loop contain x0 < 4096. They can also enable layout and divisibility simplifications. The wrapper does not prove that the shape chose the thread count or vector width, so those claims remain separate.

I counted backend invocations with a backend that returned the captured graph module’s forward method. That avoids the crashing CPU code generator and measures Dynamo’s compiled-region cache behavior:

after 64x64:        compiles = 1
after 64x64 again:  compiles = 1   (guard passed, reused)
after 128x128:      compiles = 2   (guard failed -> recompiled)

The repeated 64-row input reused the first captured entry. The 128-row input did not. “Guard failure” names the mismatch. “Retracing” means Dynamo interprets the frame again. “Recompilation” means the backend receives another graph. “Code generation” means a backend such as Inductor emits or retrieves executable artifacts. A real torch.compile call can pass through all four, but the counting backend stops before code generation. Conflating them makes cache behavior impossible to diagnose.

I then fed distinct square shapes to the same function. I expected one specialization per size:

after  16x16 : compiles = 1
after  32x32 : compiles = 2
after  48x48 : compiles = 2
after  64x64 : compiles = 2
after  80x80 : compiles = 2
...
after 192x192: compiles = 2

It compiled twice and stopped. The installed configuration has automatic_dynamic_shapes = True. After the first shape mismatch, Dynamo generalized the dimensions that changed. The generated relationship stays square because x @ x.T and the input layout let the same symbol describe both dimensions. The pinned torch._dynamo.config source sets the per-code-object recompile limit to eight, but this input never approached it. The dynamic-shapes documentation describes this static-first, generalize-after-mismatch behavior.

Passing dynamic=True asks Dynamo to attempt a dynamic graph from the first capture:

dynamic=True:  64x64 -> 1,  128x128 -> 1,  257x257 -> 1

One backend invocation served all three inputs. The separate Inductor run generated a wrapper containing:

assert_size_stride(arg1_1, (s77, s77), (s77, 1))
extern_kernels.mm(
    arg1_1,
    reinterpret_tensor(arg1_1, (s77, s77), (1, s77), 0),
    out=buf0,
)
cpp_fused_add_relu_0(buf1, s77)

s77 represents the runtime dimension. The two shape positions share it, encoding equality. The row stride also equals it, while the column stride remains one. The generated C++ receives const int64_t ks0 and loops until ks0 * ks0.

The static input has 4,096 elements, divisible by four. A symbolic square can contain a remainder, as the 257-row control demonstrates. The generated loop includes a full-width path and a masked tail:

if (C10_LIKELY(x0 < 4*(ks0*ks0 / 4))) { /* full 4-wide vector */ }
if (C10_UNLIKELY(x0 >= 4*(ks0*ks0 / 4) && x0 < ks0*ks0)) { /* masked leftover */ }

A C10_LIKELY condition handles complete vectors. A C10_UNLIKELY condition handles up to three leftover elements with a partial load and store. The symbolic size also passes through the wrapper and participates in arithmetic at runtime. These are visible costs of generality.

The code does not establish that the static version is faster overall. A static version can remove this tail logic while requiring more recompilations across a variable workload. A dynamic version can do extra integer work while avoiding those compilations. The correct comparison depends on the distribution of shapes and the number of calls per shape.

The torch.compile API documentation also distinguishes dynamic=None, dynamic=False, dynamic=True, and fullgraph. A dynamic request can still specialize where the compiler cannot create a symbolic representation. “Dynamic” means an attempt to generalize supported dimensions, not an unconditional promise that one binary accepts every tensor.

equal shapes can still miss the guard

Shape is only one property of a tensor. I kept the shape at $32 \times 32$ while changing stride, dtype, and gradient state. A counting backend printed the example tensor every time Dynamo produced a new graph:

contiguous_f32
  shape (32, 32), stride (32, 1), float32, requires_grad false
  compiles 1

contiguous_f32_repeat
  shape (32, 32), stride (32, 1), float32, requires_grad false
  compiles 1

transposed_f32
  shape (32, 32), stride (1, 32), float32, requires_grad false
  compiles 2

contiguous_f64
  shape (32, 32), stride (32, 1), float64, requires_grad false
  compiles 3

contiguous_f64_grad
  shape (32, 32), stride (32, 1), float64, requires_grad true
  compiles 4

shape_change_f32
  shape (48, 48), stride (48, 1), float32, requires_grad false
  compiles 5

The repeated contiguous float32 call reused the entry. Transposing another square tensor preserved shape while changing strides, and that required a different capture. The element at logical coordinate $(i,j)$ now lives at storage offset $i + 32j$ rather than $32i + j$. The same loop and matrix multiplication arguments cannot blindly assume both layouts.

Changing float32 to float64 also produced a new graph. Element size doubled, vector width changed, constants needed a different type, and the matrix multiplication dispatcher needed a double-precision implementation. A shape-only cache key would call code with the wrong pointer interpretation.

Setting requires_grad=True on the float64 input produced another graph even though this probe only called forward. Gradient state changes whether autograd records work and whether AOTAutograd must prepare a backward-capable path. A forward result produced under inference assumptions cannot always be reused when gradient recording is required.

The final float32 shape change produced a fifth graph rather than reusing the earlier generalized path. The sequence had already varied layout, dtype, and gradient state, so it is not the clean two-shape experiment used to demonstrate automatic dynamic shapes. Compiler cache entries attach to a combination of guards, and changes along several dimensions can prevent a symbolic shape entry from matching.

Device is absent from the output because this build used CPU tensors only. Moving an input to MPS or CUDA would require a device-specific graph and backend path. The pointer address itself is usually allowed to vary; otherwise every new allocation would recompile. The guard concerns tensor properties needed by code, not the identity of every data buffer.

Python values can create guards too. Consider:

def scaled(x, factor):
    return torch.relu(x @ x.T + factor)

If factor is an ordinary Python float, Dynamo may specialize on its value because the value becomes a constant during tracing. Passing a tensor scalar can make it a graph input instead. The two signatures express different change contracts to the compiler. One offers a constant that can be folded or embedded. The other offers variability without retracing for each scalar value, subject to dtype and device guards.

Global state participates for the same reason. Training mode, default dtype, autocast state, deterministic-algorithm settings, and grad mode can change operation semantics or dispatch. A guard that ignores relevant state risks silently executing code compiled under incompatible assumptions. A guard that watches too much misses often and destroys reuse. Guard design is therefore part correctness proof and part performance policy.

The torch.compile documentation says cache entries are associated with a code object and checked by guards. It also documents a recompile limit of eight before further compilation is skipped for that code object under the default behavior. Hitting the limit is not an optimization result. It can mean the program falls back to eager execution for unmatched calls, changing performance without changing the Python source.

For production measurement, a warm median from one shape is incomplete unless the workload’s guard-hit rate is known. A service can have excellent compiled execution on its common shape and terrible tail latency when rare strides, dtypes, or sequence lengths trigger compilation. Shape frequency and compile events belong beside request latency in the trace.

one scalar sends python back into the path

The clean function never asks Python to inspect a tensor value. I changed that:

def g(x):
    y = x @ x.T
    if y.sum().item() > 0:
        y = y + 1.0
    else:
        y = y - 1.0
    return torch.relu(y)

y.sum() is still a tensor, with shape (), meaning it contains one scalar. .item() converts that tensor value to an ordinary Python number. The if then chooses Python control flow using data produced by the matrix multiplication.

At the pinned revision, torch._dynamo.config.capture_scalar_outputs is false. The source for TensorVariable.method_item explicitly raises an unsupported case with the text Tensor.item() call with capture_scalar_outputs=False. torch._dynamo.explain reported:

clean f:   graphs=1  graph_breaks=0
branchy g: graphs=2  graph_breaks=1
  break 1: Unsupported Tensor.item() call with capture_scalar_outputs=False

This is a graph break. Dynamo finishes one captured region, lets Python obtain the scalar and select the branch, then captures work after the decision as another region. The PyTorch graph-break documentation recommends using fullgraph=True when a break must become an error rather than a silent split. That mode is useful in verification because a successful call then proves that one graph captured the frame. It is less forgiving when partial compilation is still valuable.

Setting scalar-output capture can move .item() into a graph, but it does not make arbitrary data-dependent Python control flow a static graph. A captured symbolic scalar still has to govern an operation the backend can represent. The exact unsupported boundary depends on the surrounding code and compiler version. The pinned default and the observed break are narrower facts.

To measure a seam without changing arithmetic or branch direction, I inserted the explicit operation torch._dynamo.graph_break() between addition and relu:

def clean(x):
    y = x @ x.T
    y = y + 1.0
    return torch.relu(y)

def broken(x):
    y = x @ x.T
    y = y + 1.0
    torch._dynamo.graph_break()
    return torch.relu(y)

explain reported one graph and zero breaks for clean, then two graphs and one break for broken. Both compiled functions produced equal outputs. After ten warmup calls, each received 41 timed calls at $1024 \times 1024$. A seeded random choice decided which function was measured first:

clean:  median 3066.71 us   min 2901.38 us   max 3610.42 us
broken: median 3490.54 us   min 3291.17 us   max 3875.00 us
median difference: 423.83 us

The measurement has two weaknesses. The functions were not alternated for every sample, so slow drift can still bias one set. The min and max are not confidence intervals. A stronger run would interleave paired samples and bootstrap the paired differences.

The complete verifier later repeated the probe and measured medians of 2,414.25 microseconds and 2,635.08 microseconds, a difference of 220.83 microseconds. The mechanism and ordering survived; the estimated penalty nearly halved. That spread prevents using 424 microseconds as a stable graph-break cost. It remains one observation from one run.

The generated-region boundary still explains a real restriction. An operation before the break cannot fuse with one after it because no single Inductor graph contains both. The value crossing the boundary must be represented as an output of the first region and an input of the second. That is materialization at the graph interface.

Materialization does not prove a DRAM round trip. The tensor can remain in shared cache, and fixed dispatch costs can contribute to the 424 microseconds. No cache-miss or memory-controller counter was collected. The byte count gives a possible cost model: a $1024 \times 1024$ float32 tensor contains 4 MiB, so exposing and consuming one extra tensor interface concerns at least that logical payload. The clock cannot say which cache level served it.

The clean function’s add and relu can inhabit one generated loop. The explicit break guarantees they cannot. That is enough to predict a penalty that grows with the boundary value, plus a fixed region transition. Testing several sizes with paired samples would separate those terms. The current probe establishes only one point and the mechanism visible in the graphs.

a python loop can disappear by becoming larger

The scalar branch shows one kind of Python that cannot stay inside the captured region under the current settings. A fixed Python loop behaves differently:

def repeat(x, steps):
    for _ in range(steps):
        x = torch.relu(x + 1.0)
    return x

With steps=1, Dynamo captured one add and one relu. With steps=3, it captured six call nodes:

add -> relu -> add -> relu -> add -> relu

No loop node appeared. Python executed the range protocol while tracing, and Dynamo unrolled the known number of iterations into repeated graph operations. The compiled program has no loop counter for steps; it has a larger straight-line graph specialized to the observed Python integer.

The counting backend confirmed that specialization:

steps 1: compiles 1
steps 1: compiles 1
steps 3: compiles 2
steps 5: compiles 3

Repeating the integer one reused the first graph. Three and five produced new graphs. This is profitable when the trip count belongs to a small stable set and unrolling exposes fusion. It is costly when the trip count varies widely because graph size, compilation time, and cache entries grow with it.

The six-node graph does not imply six generated loops. Inductor can fuse the chain. Algebra also matters. Repeated relu(x + 1) is not generally equivalent to replacing an arbitrary repeated operation with one operation. A compiler must prove any rewrite under the operation’s numeric rules. Seeing repetition does not authorize collapsing it.

A runtime loop over tensor data needs another representation. Higher-order operators can express control flow as graph operations with subgraphs. A symbolic loop construct gives the backend a body and runtime condition rather than forcing Python to choose every iteration during capture. Support varies by operation and backend. The key distinction is whether control flow is data represented in the graph or control exercised by the Python interpreter.

The data-dependent branch can also be rewritten with a tensor selection:

def tensor_select(x):
    condition = x.sum() > 0
    positive = torch.relu(x + 1.0)
    negative = torch.relu(x - 1.0)
    return torch.where(condition, positive, negative)

Dynamo captured this as one graph with zero breaks:

sum
greater-than
add
relu
subtract
relu
where

This removes .item() and keeps the condition as a tensor. It does not preserve the cost behavior of the Python if. torch.where receives both candidate tensors, so the graph computes both positive and negative expressions before selection. The Python branch computes only the selected side after obtaining the scalar. One version has a graph seam and conditional work; the other has a single region and unconditional work.

For cheap branches, computing both sides can be faster because fusion avoids a host decision. For expensive matrix multiplications, doubling branch work can be much worse. The compiler cannot recover laziness from where because the program explicitly constructed both arguments. Avoiding a graph break is not an optimization by itself; the replacement must preserve both values and cost structure that matter.

The difference is especially sharp on an asynchronous GPU. A Python .item() often requires the host to wait until the device produces the scalar, creating synchronization before the branch. A device-side conditional or masked tensor expression can avoid the host round trip, but may create divergence or execute extra work. I did not measure this on a GPU. The CPU graph only establishes the representational trade.

Graph capture therefore depends on more than whether source contains if or for. A Python loop with a known trip count can become repeated nodes. A Python branch on a tensor scalar breaks here. A tensor where remains one graph while evaluating both candidates. A higher-order control-flow operator can carry subgraphs. The execution consequence follows from the captured representation, not the punctuation in the source.

This also explains why fullgraph=True is a diagnostic rather than a universal performance switch. It requires one captured graph or raises. It cannot make unsupported Python semantics disappear. The code may need to express control flow differently, and a different expression can have a different computational cost.

a five-line compiler makes the missing form visible

The incorrect SSA claim left an unresolved question. PyTorch’s internal representations are large and version-sensitive. A smaller compiler can expose every transformation without pretending its structures are Inductor’s.

I restricted the input language to one variable named x, numeric constants, addition, multiplication, and a function named relu. The test expression is:

relu(x * (2 + 1) + 1)

Python’s ast.parse performs the language front end for the probe. A real parser receives tokens produced by a lexer. Tokens identify units such as the name relu, the name x, parentheses, *, +, and numbers. The parser applies grammar rules to form a tree:

Call relu
  Add
    Multiply
      Name x
      Add
        Constant 2
        Constant 1
    Constant 1

Tree structure removes the ambiguity of the flat text. Multiplication consumes x and the inner addition. The outer addition consumes that product and the last constant. The call consumes everything beneath it.

The lowerer visits the tree from the leaves upward. Each nonconstant operation creates a new temporary. When both operands are constants, it computes the result immediately rather than emitting a runtime operation. The result is:

t0 = x * 3.0
t1 = t0 + 1.0
t2 = max(t1, 0.0)

The inner 2 + 1 is absent because it became 3.0 during compilation. This is constant folding. It is valid because the operands are known and because the restricted language gives the operation ordinary numeric meaning. Constant folding in a full compiler has more hazards: integer overflow rules, floating-point rounding, exceptions, mutable state, and observable evaluation order can make an apparently obvious rewrite invalid.

This three-instruction sequence really is SSA. Each temporary has one definition. t1 refers to the value produced by t0, and t2 refers to t1. The straight-line expression contains no branch that rejoins, so it needs no merge operation. A language with:

if condition:
    value = a
else:
    value = b
use(value)

needs a way to name the selected value after the branches meet. Classic SSA uses a phi operation conceptually written value2 = phi(value_from_then, value_from_else). The selected input depends on the predecessor edge. Modern compiler IRs can represent the same idea through block arguments or other region constructs. The single-assignment property is about definitions in the IR, not spelling temporary variables once in generated source.

The probe next emits a C loop:

#include <stddef.h>

void run(const float* input, float* output, size_t n) {
    for (size_t i = 0; i < n; ++i) {
        float x = input[i];
        float t0 = x * 3.0f;
        float t1 = t0 + 1.0f;
        float t2 = t1 > 0.0f ? t1 : 0.0f;
        output[i] = t2;
    }
}

This is code generation. The SSA sequence no longer describes one abstract scalar invocation. It has been placed inside a loop over buffers, its values assigned C types, and its result stored through an output pointer. The loop establishes an iteration space. The pointers establish a calling convention between generated code and its caller. The size_t n argument makes the loop shape dynamic.

I passed that file to:

clang -O3 -dynamiclib expr.c -o expr.dylib
clang -O3 -S expr.c -o expr.s

The first command produces a loadable Mach-O dynamic library. The second stops after assembly text. Python loads the library with ctypes, assigns a signature to run, and passes five floats:

input:  -2.0  -0.25  0.0  0.5  4.0
output:  0.0   0.25  1.0  2.5  13.0

These can be checked by hand. For $-2$, multiplication and addition produce $-5$, so relu returns zero. For $-0.25$, the result before relu is $0.25$. For $4$, it is $13$. The probe compares every result to the Python calculation and fails on a mismatch.

Clang’s arm64 assembly contains a vector loop:

fmov.4s v0, #3.00000000
fmov.4s v1, #1.00000000
movi.2d v2, #0

ldp      q3, q4, [x9, #-32]
fmul.4s  v3, v3, v0
fadd.4s  v3, v3, v1
fmaxnm.4s v3, v3, v2
stp      q3, q4, [x10, #-32]

The actual loop handles 16 floats per iteration with four 128-bit vector registers, then runs a scalar remainder loop. fmul.4s multiplies four single-precision lanes. fadd.4s adds four lanes. fmaxnm.4s chooses the numeric maximum with zero for four lanes. The source contained no vector type. LLVM recognized independent loop iterations and vectorized them.

The compiler also emitted alias protection before that vector path. It compares the distance between input and output pointers and falls back when the buffers are too close for the vectorized assumptions. That check is a consequence of the C function accepting arbitrary pointers. PyTorch’s compiler has richer tensor alias information and can make different decisions.

This tiny path exposes three optimization levels that were easy to mix together in the main artifact. The custom lowerer folded a constant before C existed. Clang vectorized the generated C loop. The hardware executes the final instructions. If performance changes, the responsible transformation must be located at one of those boundaries.

It also shows why “machine code” is not the generated C++ text. C++ is another source language. Clang parses it, lowers it into LLVM IR, optimizes it, selects arm64 instructions, allocates physical registers, emits relocatable code, and links a shared object. The earlier toolchain investigation follows those steps for a standalone program. Inductor invokes a comparable path during a running Python process.

the shared object still needs a caller

The fresh PyTorch cache contained a Python wrapper, generated .cpp files, and loadable .so files. On macOS, otool -L showed the generated extension’s dynamic dependencies:

@rpath/libomp.dylib
@rpath/libc++.1.dylib
/usr/lib/libSystem.B.dylib

The .so contains machine code and relocation records in a Mach-O container despite the Linux-style suffix. Loading resolves its dynamic dependencies and makes exported entry points callable from Python. The wrapper supplies tensor data pointers, sizes, and other scalar arguments in the order expected by the generated binding.

This boundary has failure modes separate from graph capture. A graph can be valid while clang is absent. C++ can compile while the loader cannot find libomp. A library can load while a wrong pointer or lifetime causes a crash. Guards can pass while native code contains a bug. The four-thread failure occurs after capture and code generation have succeeded, which narrows the fault but does not locate it.

The cache names are content-derived hashes, but a filename alone does not prove every dependency of correctness is in that hash. PyTorch’s cache key and surrounding metadata decide which compiler settings, source fragments, architecture properties, and runtime versions participate. Version compatibility is therefore operational, not merely a disk-space concern. Reusing an artifact across an incompatible runtime would be dangerous; refusing reuse makes startup slower.

The cache experiment found nine files after the first process and still nine after the second. That proves artifact reuse in this controlled directory. It does not prove which individual files were read, memory-mapped, or rebuilt. File-system tracing could answer that. The first-process and second-process times already show why file count is an insufficient cache metric.

backward creates another compiled program

The original f returns a tensor but computes no gradient. A training call changes what compilation must produce. I wrapped the expression in a scalar loss:

def loss(x):
    y = torch.relu(x @ x.T + 1.0)
    return y.square().mean()

The input has requires_grad=True. Calling the compiled loss creates the forward value; calling backward() asks for the gradient of that scalar with respect to every element of x.

The forward cannot discard everything after producing the scalar. The backward formula for relu needs to know which preactivation values were positive. The gradient of the square needs the forward value. Matrix multiplication backward needs operands or equivalent saved information. AOTAutograd analyzes those needs while partitioning forward and backward work.

The run produced two Python wrappers and three C++ files. The AOTAutograd counters recorded one total compilation, one cache miss, one successful result, and one saved cache entry. The loss was 109.42401123046875; the gradient sum was -5.621288299560547; every gradient value was finite.

The forward wrapper contains:

extern_kernels.mm(
    primals_1,
    reinterpret_tensor(primals_1, (64, 64), (1, 64), 0),
    out=buf0,
)
cpp_fused_add_mean_pow_relu_0(buf2, buf0)
return (buf2, primals_1, buf0)

buf2 is the scalar loss. The wrapper also returns primals_1 and buf0 to the autograd machinery. They are saved values for the backward. The inference wrapper returned only its final tensor. Compilation did not erase the mathematical need for intermediate state; it made that state explicit at the forward and backward boundary.

The backward wrapper receives the original input, the saved matrix product, and the incoming scalar tangent:

assert_size_stride(primals_1, (64, 64), (64, 1))
assert_size_stride(mm, (64, 64), (64, 1))
assert_size_stride(tangents_1, (), ())

cpp_fused_add_div_expand_mul_pow_relu_threshold_backward_0(
    buf0, tangents_1
)
extern_kernels.mm(buf0, primals_1, out=buf1)
extern_kernels.addmm(
    buf1,
    reinterpret_tensor(buf0, (64, 64), (1, 64), 0),
    primals_1,
    alpha=1,
    beta=1,
    out=buf2,
)
return (buf2,)

The long fused name exposes several derivative operations. Mean contributes a division by the number of elements. Squaring contributes multiplication by the forward value and the constant two. relu contributes a threshold mask. The scalar upstream gradient is expanded over the matrix. Those elementwise pieces become one generated function.

The two matrix products in the gradient follow from $x x^\mathsf{T}$. A small differential derivation makes the symmetry visible. Let:

$$M = X X^\mathsf{T}$$

and let $G$ be the gradient arriving at $M$. A small change in $X$ gives:

$$dM = dX X^\mathsf{T} + X dX^\mathsf{T}$$

The gradient contribution through the first occurrence of $X$ is $G X$. The contribution through the transposed occurrence is $G^\mathsf{T} X$. Therefore:

$$\frac{\partial L}{\partial X} = G X + G^\mathsf{T} X$$

One external mm computes one term. The addmm adds the second matrix product into the first output. The generated code and the derivative agree.

This is where AOTAutograd changes memory economics. Saving buf0 uses $64 \times 64 \times 4 = 16{,}384$ bytes for this example. For a large model, saved activations can dominate memory. Recomputing selected values during backward can reduce storage at the cost of more computation. That tradeoff belongs to graph partitioning and checkpointing, not to the small inference wrapper.

The backward also refutes a tempting fusion story. The entire derivative does not become one loop because matrix multiplications remain external operations with elementwise work around them. Compilation arranges and fuses what its lowerings permit. It does not abolish library boundaries or the data dependencies of calculus.

one shape can select several programs

Shape specialization reaches beyond static versus symbolic loop bounds. A matrix multiplication lowering can choose among algorithms based on dimensions, dtype, layout, device, and available workspace. A pointwise scheduler can choose vector width, tiling, or parallel structure. A GPU compiler can choose block size, warp count, and staging depth.

An autotuner makes some of those choices empirically. It generates candidates, runs them on representative inputs, checks or assumes correctness according to its contract, and records a winner. The search has a cost. If the shape runs once, search can cost more than the saved execution time. If the shape repeats millions of times, a small per-call improvement can repay substantial tuning.

The installed tuned_mm source includes candidate construction and checks configuration flags such as max_autotune and max_autotune_gemm. The CPU probe with max_autotune emitted no candidate table. I cannot tell from silence whether one choice survived, logging did not trigger, a cache supplied a decision, or this path skipped timed comparison. The article therefore records the source mechanism and the missing observation separately.

Shape specialization and autotuning also create a cache-cardinality problem. Suppose a service sees sequence lengths from 1 through 8,192, two dtypes, several batch sizes, and multiple layouts. A fully specialized candidate for every combination can consume compile time and storage while most entries receive little traffic. Bucketing shapes, padding, dynamic kernels, and selective tuning trade a little execution efficiency for fewer variants.

The relevant performance equation includes compilation:

$$T_{\text{total}} = T_{\text{capture}} + T_{\text{lower}} + T_{\text{generate}} + T_{\text{native compile}} + T_{\text{load}} + C \times T_{\text{execute}}$$

$C$ is the number of executions that reuse the artifact. If eager takes $6$ microseconds and compiled warm execution takes $24$ microseconds, no positive $C$ repays compilation because every execution is slower. At 1,024 rows, eager takes about $3{,}238$ microseconds and compiled takes $2{,}858$, saving $380$ microseconds per call. If a relevant cold path cost 8 seconds, the simple break-even count would be:

$$C = \frac{8\ \text{seconds}}{0.000380\ \text{seconds per call}} \approx 21{,}053\ \text{calls}$$

The units cancel to calls. This calculation mixes a one-run cold measurement with warm medians, so it is illustrative rather than a capacity estimate. It exposes the variable that matters: a speedup without reuse can still lose over the process lifetime.

If compilation occurs while a production request waits, the latency distribution gains a long tail. If compilation happens during controlled warmup, startup gets longer but request latency can remain stable. If a new shape appears after rollout, it can reintroduce the cold path at the worst moment. A serving system therefore needs visibility into graph counts, guard failures, compile latency, cache hits, and shape frequency, not only kernel duration.

the gpu path is still a prediction here

The PyTorch compiler overview states that Inductor uses Triton as a principal code-generation language for supported GPUs. Triton describes blocked programs over tensors rather than C++ loops over host pointers. Its compiler lowers those programs toward GPU-specific code. The same high-level FX region can therefore reach different lowerings and generated languages on CPU and GPU.

I did not run this expression on an NVIDIA or AMD GPU. No Triton source, PTX, GPU ISA, occupancy result, register count, memory-throughput measurement, or launch timeline in this investigation is local evidence. The CPU wrapper cannot answer how the GPU backend would fuse the matrix multiplication tail. It also cannot establish whether a vendor library, Triton template, or another backend wins for a given GPU shape.

Several costs would change. A GPU operation has host launch overhead and asynchronous execution. Correct timing needs synchronization or device events. Data transfer matters if the input begins on the CPU. Vector width becomes warp and block geometry. Registers are allocated per GPU thread, shared memory is per block, and occupancy limits how many warps can reside concurrently. None of those concepts can be read from Vectorized<float> or an OpenMP pragma.

The generated CPU code is still useful preparation because the compiler questions remain recognizable:

• Which frame region was captured?
• Which shapes and layouts do the guards admit?
• Which ATen operations were decomposed?
• Which operations went to libraries?
• Which pointwise or reduction operations fused?
• Which buffers were allocated, reused, or saved?
• Which choices were specialized and cached?
• Where did compilation time and execution time go?

Those are diagnostic questions, not a public reading path. On a GPU, their answers require GPU artifacts and GPU measurements.

MLIR is also outside the measured path. It is an infrastructure for defining multiple intermediate-representation dialects and conversions among them. Some compiler stacks use it to represent tensor, loop, vector, and device-specific levels. Finding the word MLIR in the broader compiler ecosystem would not prove that this PyTorch CPU run passed through MLIR. The artifacts and pinned source inspected here reach Dynamo, AOTAutograd, Inductor, generated C++, clang, and LLVM.

the wrapper predicts the original line

The wrapper at the beginning no longer looks like a mysterious translation:

assert_size_stride(arg0_1, (64, 64), (64, 1))
buf0 = empty_strided_cpu((64, 64), (64, 1), torch.float32)
extern_kernels.mm(
    arg0_1,
    reinterpret_tensor(arg0_1, (64, 64), (1, 64), 0),
    out=buf0,
)
buf1 = buf0
cpp_fused_add_relu_0(buf1)
return (buf1,)

The assertion belongs to one specialized entry. The reinterpretation expresses transpose through strides. The external matrix multiplication is a lowering inside the captured region. The generated C++ loop performs add and relu together. The buffer rename records reuse without proving SSA. Clang turns that C++ into a shared object that the wrapper loads and calls.

That model predicts the observations. A second $64 \times 64$ call can reuse the guarded entry. A changed shape first causes another capture, then automatic dynamic shapes can generalize it. A small tensor loses because fixed overhead exceeds avoided element traffic. A large elementwise tail approaches the explicit 2x traffic ratio. A graph break prevents cross-boundary fusion. A gradient request creates saved forward values and another compiled program. Four OpenMP threads enter a native failure that the Python graph cannot explain.

One detail remains deliberately unresolved. With OMP_NUM_THREADS=4, the generated call exits 139. With one thread, it returns the same numerical result twice. The line of Python has been followed far enough to place the failure after capture, lowering, C++ generation, compilation, and loading. The next evidence has to come from inside the native process, because the graph has already done exactly what it was asked to do.