the path · 0/15
start the path

the kernel path · XLA · lesson 06 of 7

The one-level world

Below every tensor compiler sits an IR with no tensors in it. Reading one small dump tells you exactly which facts get dropped on the way down, and why the wall this unit teaches has to exist where it does.

the goal Read a small LLVM IR dump without a reference open, then name, from the dump itself, the two things the level above knew for free that this level has to be told or has to prove at runtime.

mastery work · this chapter0/3
  1. go →
manual items are your word; auto items complete from your streaks, labs, and can-you ticks · stored in your browser only
§ 01

Four lines of C, and what clang makes of them

Start with a loop small enough to hold in your head. Three float pointers, a length, one add per element, nothing clever anywhere in it.

Compile that with clang -S -emit-llvm -O0 and roughly sixty lines of text come back. LLVM IR is what that text is called, and it is the representation every LLVM-based backend optimizes before it emits machine code. Clang emits it. Rust emits it. XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla →'s CPU backend emits it, and its GPU backend emits it too, which is why this floor is worth an hour of your time even if you never write C again.

One property to hold onto while reading: everything in the dump came from those four lines. Nothing was added by a library, and nothing is hidden in a runtime call.

vadd.c, the whole input for the clang captures in this lesson
void vadd(const float *a, const float *b, float *out, int n) {
  for (int i = 0; i < n; i++) {
    out[i] = a[i] + b[i];
  }
}
§ 02

Values, blocks, and address arithmetic

Three habits get you through any dump at this level. Read %N as a value that is defined exactly once, on the left of one instruction, and referenced by name after that. The jaxprThe traced program: one equation per primitive in single-assignment form, every shape and dtype stated.taught in /l/jaxpr → chapter at /l/jaxpr teaches that same one-name-one-definition property where you first met it, and it means here what it means there.

Read a label with a ; preds = comment after it as a basic block, a straight run of instructions with one entry and one exit. There is no for at this level. The loop is three blocks, one that tests, one that does the work, one that increments, wired together by explicit branches back to the test. Five blocks in the whole function at -O0, counting the entry and the exit.

Read getelementptr as arithmetic that touches no memory. It computes an address from a base pointer, an element type, and an index, and hands it to a separate load or store. That leaves exactly one instruction in the body doing what the C source asked for, %25 = fadd float %19, %24. Everything around it is bookkeeping, most of it because -O0 keeps every variable in an alloca slot and reloads it each time round.

clang -S -emit-llvm -O0 vadd.c: the three loop blocks, verbatim (Apple clang 21.0.0, target x86_64-apple-macosx26.0.0)
10:                                               ; preds = %30, %4
  %11 = load i32, ptr %9, align 4
  %12 = load i32, ptr %8, align 4
  %13 = icmp slt i32 %11, %12
  br i1 %13, label %14, label %33

14:                                               ; preds = %10
  %15 = load ptr, ptr %5, align 8
  %16 = load i32, ptr %9, align 4
  %17 = sext i32 %16 to i64
  %18 = getelementptr inbounds float, ptr %15, i64 %17
  %19 = load float, ptr %18, align 4
  %20 = load ptr, ptr %6, align 8
  %21 = load i32, ptr %9, align 4
  %22 = sext i32 %21 to i64
  %23 = getelementptr inbounds float, ptr %20, i64 %22
  %24 = load float, ptr %23, align 4
  %25 = fadd float %19, %24
  %26 = load ptr, ptr %7, align 8
  %27 = load i32, ptr %9, align 4
  %28 = sext i32 %27 to i64
  %29 = getelementptr inbounds float, ptr %26, i64 %28
  store float %25, ptr %29, align 4
  br label %30

30:                                               ; preds = %14
  %31 = load i32, ptr %9, align 4
  %32 = add nsw i32 %31, 1
  store i32 %32, ptr %9, align 4
  br label %10, !llvm.loop !6
the whole 67-line module, allocas and metadata included · 67 lines
; ModuleID = 'vadd.c'
source_filename = "vadd.c"
target datalayout = "e-m:o-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-apple-macosx26.0.0"

; Function Attrs: noinline nounwind optnone ssp uwtable
define void @vadd(ptr noundef %0, ptr noundef %1, ptr noundef %2, i32 noundef %3) #0 {
  %5 = alloca ptr, align 8
  %6 = alloca ptr, align 8
  %7 = alloca ptr, align 8
  %8 = alloca i32, align 4
  %9 = alloca i32, align 4
  store ptr %0, ptr %5, align 8
  store ptr %1, ptr %6, align 8
  store ptr %2, ptr %7, align 8
  store i32 %3, ptr %8, align 4
  store i32 0, ptr %9, align 4
  br label %10

10:                                               ; preds = %30, %4
  %11 = load i32, ptr %9, align 4
  %12 = load i32, ptr %8, align 4
  %13 = icmp slt i32 %11, %12
  br i1 %13, label %14, label %33

14:                                               ; preds = %10
  %15 = load ptr, ptr %5, align 8
  %16 = load i32, ptr %9, align 4
  %17 = sext i32 %16 to i64
  %18 = getelementptr inbounds float, ptr %15, i64 %17
  %19 = load float, ptr %18, align 4
  %20 = load ptr, ptr %6, align 8
  %21 = load i32, ptr %9, align 4
  %22 = sext i32 %21 to i64
  %23 = getelementptr inbounds float, ptr %20, i64 %22
  %24 = load float, ptr %23, align 4
  %25 = fadd float %19, %24
  %26 = load ptr, ptr %7, align 8
  %27 = load i32, ptr %9, align 4
  %28 = sext i32 %27 to i64
  %29 = getelementptr inbounds float, ptr %26, i64 %28
  store float %25, ptr %29, align 4
  br label %30

30:                                               ; preds = %14
  %31 = load i32, ptr %9, align 4
  %32 = add nsw i32 %31, 1
  store i32 %32, ptr %9, align 4
  br label %10, !llvm.loop !6

33:                                               ; preds = %10
  ret void
}

attributes #0 = { noinline nounwind optnone ssp uwtable "darwin-stkchk-strong-link" "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "probe-stack"="___chkstk_darwin" "stack-protector-buffer-size"="8" "target-cpu"="penryn" "target-features"="+cmov,+cx16,+cx8,+fxsr,+mmx,+sahf,+sse,+sse2,+sse3,+sse4.1,+ssse3,+x87" "tune-cpu"="generic" }

!llvm.module.flags = !{!0, !1, !2, !3, !4}
!llvm.ident = !{!5}

!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 26, i32 5]}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 8, !"PIC Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{i32 7, !"frame-pointer", i32 2}
!5 = !{!"Apple clang version 21.0.0 (clang-2100.1.1.101)"}
!6 = distinct !{!6, !7}
!7 = !{!"llvm.loop.mustprogress"}
§ 03

The type system, and the word missing from it

LLVM's type system is short enough to list from memory. Integers of any width, a handful of floating-point types, ptr, vectors, labels, tokens, metadata, and two aggregates: arrays and structs. Fetch the language reference, search it for the word tensor, and you get zero hits. That absence is the cleanest statement of this lesson available.

The closest LLVM comes is an intrinsic. llvm.matrix.multiply.* takes two flat vectors and three constant integers, and the reference describes it as treating %A as a <OuterRows> x <Inner> matrix. So the shape sits in the argument list of one call rather than in the type of any value, and nothing in the verifier connects the two.

A shape at this level is either an array type used for address arithmetic or three integers handed to an intrinsic. It is never a property of a value.

Notice what ptr does not say, either. It carries no element type at all, so load float, ptr %18 is where the float appears, and the same address could be read as an i32 by the next instruction without the type system objecting.

written aswhat it fixeswhat it leaves open
tensor<8x1024xf32>element type, rank, both extents, and that this is one whole valueeverything about memory: no address, no stride, and no aliasing question to ask
[8 x [1024 x float]]element type and both extents, as arithmetic for one getelementptrwhether two such pointers overlap, and whether anything downstream still treats this as one value
ptrthat this is an addresselement type, extent, alignment, provenance, all of which move to the load, the store, or metadata
<8 x float>a register-width value, one machine operation wideany relation to the array it was read out of
one 8 by 1024 array of f32, and what each notation still knows about it
§ 04

The same add, emitted by XLA

None of this is theoretical for a JAX program. Set XLA_FLAGS=--xla_dump_to=DIR around a CPU compile and the dump directory gains two .ll files per module, one as XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla → emitted it and one after LLVM has optimized it. Which flag does what belongs to another lesson in this unit; what matters here is what the emission looks like.

Lower a + b on two arrays of shape (8, 1024) and XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla → writes the loop nest itself, two levels deep, with block names that carry the HLO instruction that produced them. The shape survives into getelementptr inbounds [8 x [1024 x float]], and that array type is the last trace of it. No value in the module has that type.

The metadata is the part worth staring at. Both loads carry !noalias !3, the store carries !alias.scope !3, and the domain string is XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla →'s own: !5 = !{!"XLA host kernel add.3 AA domain"}. Buffer assignment already knew these three slices do not overlap, and metadata is the only channel that fact had left. The !dereferenceable note says i64 32768, which is 8 by 1024 by 4 bytes, spelled out because the type no longer says it.

module_0005.jit__lambda_.ir-no-opt.ll, the inner loop body XLA:CPU emitted for one stablehlo.add (jax 0.4.38, CPU backend, XLA_FLAGS=--xla_dump_to)
add.3.loop_body.dim.1:                            ; preds = %add.3.loop_header.dim.1
  %4 = getelementptr inbounds [8 x [1024 x float]], ptr %arg0, i64 0, i64 %add.3.indvar.dim.0, i64 %add.3.indvar.dim.1
  %5 = load float, ptr %4, align 4, !invariant.load !0, !noalias !3
  %6 = getelementptr inbounds [8 x [1024 x float]], ptr %arg1, i64 0, i64 %add.3.indvar.dim.0, i64 %add.3.indvar.dim.1
  %7 = load float, ptr %6, align 4, !invariant.load !0, !noalias !3
  %add.3 = fadd float %5, %7
  %8 = getelementptr inbounds [8 x [1024 x float]], ptr %arg2, i64 0, i64 %add.3.indvar.dim.0, i64 %add.3.indvar.dim.1
  store float %add.3, ptr %8, align 4, !alias.scope !3
  %invar.inc5 = add nuw nsw i64 %add.3.indvar.dim.1, 1
  store i64 %invar.inc5, ptr %add.3.invar_address.dim.1, align 4
  br label %add.3.loop_header.dim.1
the whole emitted kernel, 91 lines, call frame unpacking included · 91 lines
; ModuleID = '__compute_module'
source_filename = "__compute_module"
target datalayout = "e-m:o-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-darwin25.5.0"

%XLA_CPU_KernelCallFrame = type { ptr, ptr, i64, ptr }
%XLA_CPU_KernelThreadDim = type { i64, i64, i64 }
%XLA_CPU_KernelThread = type { i64, i64, i64 }
%XLA_CPU_KernelArg = type { ptr, i64 }

; Function Attrs: uwtable
define ptr @add.3(ptr %0) #0 {
  %add.3.invar_address.dim.1 = alloca i64, align 8
  %add.3.invar_address.dim.0 = alloca i64, align 8
  %tdims_gep = getelementptr inbounds nuw %XLA_CPU_KernelCallFrame, ptr %0, i32 0, i32 0
  %tdims = load ptr, ptr %tdims_gep, align 8
  %tdim_x_gep = getelementptr inbounds nuw %XLA_CPU_KernelThreadDim, ptr %tdims, i32 0, i32 0
  %tdim_y_gep = getelementptr inbounds nuw %XLA_CPU_KernelThreadDim, ptr %tdims, i32 0, i32 1
  %tdim_z_gep = getelementptr inbounds nuw %XLA_CPU_KernelThreadDim, ptr %tdims, i32 0, i32 2
  %tdim_x = load i64, ptr %tdim_x_gep, align 4
  %tdim_y = load i64, ptr %tdim_y_gep, align 4
  %tdim_z = load i64, ptr %tdim_z_gep, align 4
  %tid_gep = getelementptr inbounds nuw %XLA_CPU_KernelCallFrame, ptr %0, i32 0, i32 1
  %tids = load ptr, ptr %tid_gep, align 8
  %tid_x_gep = getelementptr inbounds nuw %XLA_CPU_KernelThread, ptr %tids, i32 0, i32 0
  %tid_y_gep = getelementptr inbounds nuw %XLA_CPU_KernelThread, ptr %tids, i32 0, i32 1
  %tid_z_gep = getelementptr inbounds nuw %XLA_CPU_KernelThread, ptr %tids, i32 0, i32 2
  %tid_x = load i64, ptr %tid_x_gep, align 4
  %tid_y = load i64, ptr %tid_y_gep, align 4
  %tid_z = load i64, ptr %tid_z_gep, align 4
  %args_gep = getelementptr inbounds nuw %XLA_CPU_KernelCallFrame, ptr %0, i32 0, i32 3
  %args = load ptr, ptr %args_gep, align 8
  %arg0_gep = getelementptr %XLA_CPU_KernelArg, ptr %args, i32 0, i32 0
  %arg0 = load ptr, ptr %arg0_gep, align 8, !invariant.load !0, !dereferenceable !1, !align !2
  %args_gep1 = getelementptr inbounds nuw %XLA_CPU_KernelCallFrame, ptr %0, i32 0, i32 3
  %args2 = load ptr, ptr %args_gep1, align 8
  %arg1_gep = getelementptr %XLA_CPU_KernelArg, ptr %args2, i32 1, i32 0
  %arg1 = load ptr, ptr %arg1_gep, align 8, !invariant.load !0, !dereferenceable !1, !align !2
  %args_gep3 = getelementptr inbounds nuw %XLA_CPU_KernelCallFrame, ptr %0, i32 0, i32 3
  %args4 = load ptr, ptr %args_gep3, align 8
  %arg2_gep = getelementptr %XLA_CPU_KernelArg, ptr %args4, i32 2, i32 0
  %arg2 = load ptr, ptr %arg2_gep, align 8, !invariant.load !0, !dereferenceable !1, !align !2
  store i64 0, ptr %add.3.invar_address.dim.0, align 4
  br label %add.3.loop_header.dim.0

add.3.loop_header.dim.0:                          ; preds = %add.3.loop_exit.dim.1, %1
  %add.3.indvar.dim.0 = load i64, ptr %add.3.invar_address.dim.0, align 4
  %2 = icmp uge i64 %add.3.indvar.dim.0, 8
  br i1 %2, label %add.3.loop_exit.dim.0, label %add.3.loop_body.dim.0

add.3.loop_body.dim.0:                            ; preds = %add.3.loop_header.dim.0
  store i64 0, ptr %add.3.invar_address.dim.1, align 4
  br label %add.3.loop_header.dim.1

add.3.loop_header.dim.1:                          ; preds = %add.3.loop_body.dim.1, %add.3.loop_body.dim.0
  %add.3.indvar.dim.1 = load i64, ptr %add.3.invar_address.dim.1, align 4
  %3 = icmp uge i64 %add.3.indvar.dim.1, 1024
  br i1 %3, label %add.3.loop_exit.dim.1, label %add.3.loop_body.dim.1

add.3.loop_body.dim.1:                            ; preds = %add.3.loop_header.dim.1
  %4 = getelementptr inbounds [8 x [1024 x float]], ptr %arg0, i64 0, i64 %add.3.indvar.dim.0, i64 %add.3.indvar.dim.1
  %5 = load float, ptr %4, align 4, !invariant.load !0, !noalias !3
  %6 = getelementptr inbounds [8 x [1024 x float]], ptr %arg1, i64 0, i64 %add.3.indvar.dim.0, i64 %add.3.indvar.dim.1
  %7 = load float, ptr %6, align 4, !invariant.load !0, !noalias !3
  %add.3 = fadd float %5, %7
  %8 = getelementptr inbounds [8 x [1024 x float]], ptr %arg2, i64 0, i64 %add.3.indvar.dim.0, i64 %add.3.indvar.dim.1
  store float %add.3, ptr %8, align 4, !alias.scope !3
  %invar.inc5 = add nuw nsw i64 %add.3.indvar.dim.1, 1
  store i64 %invar.inc5, ptr %add.3.invar_address.dim.1, align 4
  br label %add.3.loop_header.dim.1

add.3.loop_exit.dim.1:                            ; preds = %add.3.loop_header.dim.1
  %invar.inc = add nuw nsw i64 %add.3.indvar.dim.0, 1
  store i64 %invar.inc, ptr %add.3.invar_address.dim.0, align 4
  br label %add.3.loop_header.dim.0

add.3.loop_exit.dim.0:                            ; preds = %add.3.loop_header.dim.0
  br label %return

return:                                           ; preds = %add.3.loop_exit.dim.0
  ret ptr null
}

attributes #0 = { uwtable "frame-pointer"="all" "prefer-vector-width"="256" }

!0 = !{}
!1 = !{i64 32768}
!2 = !{i64 64}
!3 = !{!4}
!4 = !{!"result slice: {index:0, offset:0, size:32768}", !5}
!5 = !{!"XLA host kernel add.3 AA domain"}
§ 05

Told, or proved at runtime

Compile the same four-line C loop at -O2 and the vectorizer does not start with vectors. It starts with a guard: two pointer subtractions, two unsigned compares against 32, and a branch that skips the vector path if either distance is under 32 bytes. That is the compiler proving at runtime what it could not prove at compile time, that the output does not overlap either input. The function goes from five basic blocks to eleven, and most of the new ones are the guard and the scalar tails it needs.

XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla →'s version of the same add never asks. Its optimized dump walks from the entry block straight to br label %vector.ph and loads <8 x float> in the body, because the !noalias metadata was already sitting there. Same optimizer, same machine, same arithmetic.

Same vectorizer, same machine. One of them was told.

That gap is the shape of every question this unit asks. The level above knows things by construction, that a value is a whole array, that two buffers are distinct, that one op covers every element. The level below can only be told those things or re-derive them. So the decisions the unit reads out of a compiled dump have to be made while that knowledge still exists in the IR, and by the time you are reading LLVM IR they have been made for you. The next lesson is about the design that refuses to accept one level at all.

clang -S -emit-llvm -O2 vadd.c: the runtime aliasing guard, ahead of any vector instruction
9:                                                ; preds = %4
  %10 = zext nneg i32 %3 to i64
  %11 = icmp ult i32 %3, 8
  br i1 %11, label %38, label %12

12:                                               ; preds = %9
  %13 = sub i64 %7, %6
  %14 = icmp ult i64 %13, 32
  %15 = sub i64 %7, %5
  %16 = icmp ult i64 %15, 32
  %17 = or i1 %14, %16
  br i1 %17, label %38, label %18
the whole -O2 function: guard, unrolled vector body, and two scalar tails · 121 lines
; ModuleID = 'vadd.c'
source_filename = "vadd.c"
target datalayout = "e-m:o-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-apple-macosx26.0.0"

; Function Attrs: nofree norecurse nosync nounwind ssp memory(argmem: readwrite) uwtable
define void @vadd(ptr noundef readonly captures(none) %0, ptr noundef readonly captures(none) %1, ptr noundef writeonly captures(none) %2, i32 noundef %3) local_unnamed_addr #0 {
  %5 = ptrtoint ptr %1 to i64
  %6 = ptrtoint ptr %0 to i64
  %7 = ptrtoint ptr %2 to i64
  %8 = icmp sgt i32 %3, 0
  br i1 %8, label %9, label %58

9:                                                ; preds = %4
  %10 = zext nneg i32 %3 to i64
  %11 = icmp ult i32 %3, 8
  br i1 %11, label %38, label %12

12:                                               ; preds = %9
  %13 = sub i64 %7, %6
  %14 = icmp ult i64 %13, 32
  %15 = sub i64 %7, %5
  %16 = icmp ult i64 %15, 32
  %17 = or i1 %14, %16
  br i1 %17, label %38, label %18

18:                                               ; preds = %12
  %19 = and i64 %10, 2147483640
  br label %20

20:                                               ; preds = %20, %18
  %21 = phi i64 [ 0, %18 ], [ %34, %20 ]
  %22 = getelementptr inbounds nuw float, ptr %0, i64 %21
  %23 = getelementptr inbounds nuw i8, ptr %22, i64 16
  %24 = load <4 x float>, ptr %22, align 4, !tbaa !6
  %25 = load <4 x float>, ptr %23, align 4, !tbaa !6
  %26 = getelementptr inbounds nuw float, ptr %1, i64 %21
  %27 = getelementptr inbounds nuw i8, ptr %26, i64 16
  %28 = load <4 x float>, ptr %26, align 4, !tbaa !6
  %29 = load <4 x float>, ptr %27, align 4, !tbaa !6
  %30 = fadd <4 x float> %24, %28
  %31 = fadd <4 x float> %25, %29
  %32 = getelementptr inbounds nuw float, ptr %2, i64 %21
  %33 = getelementptr inbounds nuw i8, ptr %32, i64 16
  store <4 x float> %30, ptr %32, align 4, !tbaa !6
  store <4 x float> %31, ptr %33, align 4, !tbaa !6
  %34 = add nuw i64 %21, 8
  %35 = icmp eq i64 %34, %19
  br i1 %35, label %36, label %20, !llvm.loop !10

36:                                               ; preds = %20
  %37 = icmp eq i64 %19, %10
  br i1 %37, label %58, label %38

38:                                               ; preds = %12, %9, %36
  %39 = phi i64 [ 0, %12 ], [ 0, %9 ], [ %19, %36 ]
  %40 = and i64 %10, 3
  %41 = icmp eq i64 %40, 0
  br i1 %41, label %54, label %42

42:                                               ; preds = %38, %42
  %43 = phi i64 [ %51, %42 ], [ %39, %38 ]
  %44 = phi i64 [ %52, %42 ], [ 0, %38 ]
  %45 = getelementptr inbounds nuw float, ptr %0, i64 %43
  %46 = load float, ptr %45, align 4, !tbaa !6
  %47 = getelementptr inbounds nuw float, ptr %1, i64 %43
  %48 = load float, ptr %47, align 4, !tbaa !6
  %49 = fadd float %46, %48
  %50 = getelementptr inbounds nuw float, ptr %2, i64 %43
  store float %49, ptr %50, align 4, !tbaa !6
  %51 = add nuw nsw i64 %43, 1
  %52 = add i64 %44, 1
  %53 = icmp eq i64 %52, %40
  br i1 %53, label %54, label %42, !llvm.loop !14

54:                                               ; preds = %42, %38
  %55 = phi i64 [ %39, %38 ], [ %51, %42 ]
  %56 = sub nsw i64 %39, %10
  %57 = icmp ugt i64 %56, -4
  br i1 %57, label %58, label %59

58:                                               ; preds = %54, %59, %36, %4
  ret void

59:                                               ; preds = %54, %59
  %60 = phi i64 [ %88, %59 ], [ %55, %54 ]
  %61 = getelementptr inbounds nuw float, ptr %0, i64 %60
  %62 = load float, ptr %61, align 4, !tbaa !6
  %63 = getelementptr inbounds nuw float, ptr %1, i64 %60
  %64 = load float, ptr %63, align 4, !tbaa !6
  %65 = fadd float %62, %64
  %66 = getelementptr inbounds nuw float, ptr %2, i64 %60
  store float %65, ptr %66, align 4, !tbaa !6
  %67 = add nuw nsw i64 %60, 1
  %68 = getelementptr inbounds nuw float, ptr %0, i64 %67
  %69 = load float, ptr %68, align 4, !tbaa !6
  %70 = getelementptr inbounds nuw float, ptr %1, i64 %67
  %71 = load float, ptr %70, align 4, !tbaa !6
  %72 = fadd float %69, %71
  %73 = getelementptr inbounds nuw float, ptr %2, i64 %67
  store float %72, ptr %73, align 4, !tbaa !6
  %74 = add nuw nsw i64 %60, 2
  %75 = getelementptr inbounds nuw float, ptr %0, i64 %74
  %76 = load float, ptr %75, align 4, !tbaa !6
  %77 = getelementptr inbounds nuw float, ptr %1, i64 %74
  %78 = load float, ptr %77, align 4, !tbaa !6
  %79 = fadd float %76, %78
  %80 = getelementptr inbounds nuw float, ptr %2, i64 %74
  store float %79, ptr %80, align 4, !tbaa !6
  %81 = add nuw nsw i64 %60, 3
  %82 = getelementptr inbounds nuw float, ptr %0, i64 %81
  %83 = load float, ptr %82, align 4, !tbaa !6
  %84 = getelementptr inbounds nuw float, ptr %1, i64 %81
  %85 = load float, ptr %84, align 4, !tbaa !6
  %86 = fadd float %83, %85
  %87 = getelementptr inbounds nuw float, ptr %2, i64 %81
  store float %86, ptr %87, align 4, !tbaa !6
  %88 = add nuw nsw i64 %60, 4
  %89 = icmp eq i64 %88, %10
  br i1 %89, label %58, label %59, !llvm.loop !16
}
before you move on

Check yourself

01 In `%18 = getelementptr inbounds float, ptr %15, i64 %17`, which part of the original array reached this instruction and which part did not?

The element type reached it, as the `float` the index gets scaled by. Nothing else did: the extent is absent, the rank is absent, and `ptr` does not say what it points at. In the XLA emission the extents survive one step further, inside the array type `[8 x [1024 x float]]`, but only as arithmetic; no value in that module has that type.

02 The -O2 loop opens with two pointer subtractions and two unsigned compares against 32. What is that code for, and why has the XLA emission no equivalent?

It is a runtime overlap check. The vectorizer may only use wide loads and stores if the output does not alias the inputs, and at this level nothing proves that, so it tests the byte distances and falls back to a scalar loop when they are too small. XLA attaches `!noalias` and `!alias.scope` from buffer assignment instead, so the same vectorizer enters the vector body unconditionally.

03 Someone proposes skipping the tensor compiler and letting LLVM fuse the loops instead. Which facts would LLVM have to re-derive first?

That the buffers are distinct, that a loop nest is one operation over one whole array, and that a producer has exactly one consumer. All three are free above and cost a dependence analysis over scalar address math here. LLVM does have loop fusion and vectorization; what it lacks is the information. And the algebraic rewrite the chapter at /l/xla is about, the one that turns two-pass softmax into a streaming kernel, is out of reach at every level anyway.

assigned

Readings