Phase 13 / Pallas kernels
Iterate with correctness and performance evidence
Revision v0.7.0 · Updated since your last visit. Review the changes; your progress is saved.
Lesson · about 120 minutes · CPU practice
Build a representative shape/dtype correctness matrix for an actual Pallas pipeline. Separate kernel implementation error from input/output precision differences.
Phase study guide: preparation, milestones and project evidence
Before you start
Before this: Tiling, memory, and pipelining
Hardware: CPU Pallas interpretation; actual TPU lowering and performance unverified
Setup help: Python, notebook, or Cloud TPU VM
Use the tested CPU packages (the recorded environment uses Python 3.14.3). Other Python versions and accelerator environments require separate checks. The optional Cloud TPU specialization explains the target-specific execution boundaries. Logical CPU devices practice partitioning; they do not simulate TPU performance.
Download the repository ZIP or use your existing checkout. Extract it and open a terminal in the folder containing requirements-cpu.txt. The beginner ZIP is separate; these lessons use the full repository folder.
python3 -m venv .venv
source .venv/bin/activate
python -m pip install -r requirements-cpu.txtpy -3 -m venv .venv
.\.venv\Scripts\python.exe -m pip install -r requirements-cpu.txtAfter activation, run the complete worked example from that same folder:
python phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.py.\.venv\Scripts\python.exe phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.pyFor an incremental build, create main.py in your own practice folder, append each build block in order, and run python main.py with this environment active. Keep that file separate from the downloaded reference solution.
Download this lesson’s notebook. Open it with a Python kernel using the course environment. Run cells from the top after restarting the kernel; experiments depend on earlier definitions.
For local Jupyter, first install the course requirements using the Local Python steps, then use your environment’s Python to start Jupyter:
python -m pip install jupyterlab
python -m jupyter lab.\.venv\Scripts\python.exe -m pip install jupyterlab
.\.venv\Scripts\python.exe -m jupyter labSelect the downloaded kernels-04.ipynb file in JupyterLab. The notebook contains reference solutions; make your prediction and attempt each practice task before opening them. Restart the kernel before changing logical device setup.
The worked reference for this lesson is CPU-validated. Some references explicitly configure logical CPU devices; setting JAX_PLATFORMS does not adapt them to a TPU. Follow the execution matrix and reviewed multi-host example before attempting a target run.
For accelerator preparation, study TPU setup, observed runtime receipts, checkpointed jobs and precision and retained traces. These target paths remain unqualified on physical TPU hardware until an actual receipt is recorded. Cloud provisioning is optional for CPU learning.
Check: All eight shape/dtype cases exactly match the represented-input oracle. bfloat16 differs from original float32 values by a separately reported precision gap. The benchmark refuses CPU execution; no target speedup is reported.
Revisit the guided CPU setup · Open 19: TPU Setup. Package instructions are CPU instructions; Windows and accelerator environments need their own validation.
What you will be able to do
- Build a representative shape/dtype correctness matrix for an actual Pallas pipeline.
- Separate kernel implementation error from input/output precision differences.
- Implement a guarded compile/warmup/synchronized target benchmark.
- Connect profiling, partitioning, tiling and numerical evidence to an optimization decision.
The problem
A custom kernel is useful only if it stays correct when shapes and dtypes change and improves an actual workload. What evidence is enough to keep an optimization, and what should we do when no target hardware is available? We will build a repeatable audit and an executable target-only measurement protocol.
The idea
A low-precision kernel can differ from a high-precision calculation because inputs changed representation, arithmetic changed, or the implementation is wrong. Separate these causes with references that answer different questions.
Use two references to locate a precision gap
First compare the kernel with an independent computation using the same represented inputs. That isolates implementation behavior under the chosen numeric contract. Then compare the represented-input reference with the original high-precision calculation to estimate representation effects.
If the first comparison passes while the second shows a gap, changing indexing is unlikely to fix that gap. Investigate the precision policy and whether downstream quality tolerates it.
The error bars should keep those comparisons separate and state scale, dtype and tolerance. A passed numerical check also says nothing about speed. Keep workload identity, synchronization and target-device measurements with any optimization claim.
Choose cases that challenge the contract
The fixtures include a singleton matrix, one exact tile, a small tail and a larger tail. The same frozen random data is tested in float32 and bfloat16. A benchmark-sized divisible matrix alone would miss many boundary bugs. Every implementation must return the original logical shape; padding is internal work, not extra observations in the output.
Use two references for two different questions
The represented-input oracle promotes the actual stored inputs to float32, computes , then applies the declared output cast. Kernel agreement with that reference checks implementation correctness. Comparing against the original float32 dataset instead reveals input and output rounding. These differences should not be merged into one unexplained error number or hidden by arbitrarily loose tolerances.
Make target timing a complete operation
The target protocol compiles separately, warms up five times and synchronizes every measured output. Inputs are already placed on the target before timing. The measured boundary includes logical-shape padding, the Pallas call and cropping, because those are part of this wrapper’s user-visible cost. A kernel-only benchmark would need a different explicit boundary and could not be compared directly with these results.
Compare a strong baseline with the same contract
The baseline is jitted ordinary JAX with float32 arithmetic and the same final dtype cast. XLA can fuse simple expressions, so the custom kernel may lose. Report every sample, median and ninetieth percentile, plus compilation time and actual device. The percentile is a distribution summary, not the maximum. Repeated samples expose variability but do not erase cache, thermal or system-load effects.
Relate local kernels to the distributed workload
A kernel optimizes an operation on local arrays; it does not automatically reduce all-reduce or resharding costs. Use a profile to determine whether the workload is limited by local compute, memory movement, input staging or communication. Keep the partitioning decision explicit and verify the same global numerical objective. A faster local operation may have little end-to-end benefit when another part dominates.
Accept, revise or decline the optimization
Keep the change only when representative correctness and the declared target boundary improve enough to justify maintenance and specialization costs. If performance is unchanged or worse, retain the simpler baseline and record the experiment. If target hardware is unavailable, retain the tested kernel and its runnable target protocol, mark performance unverified, and do not fill the gap with interpret-mode timings. The scale synthesis assessment asks for profiling and partitioning evidence as well as the local kernel audit.
Precision gaps remain after kernel correctness passes
Predict before looking: Can a kernel have zero implementation error and still differ from the original higher-precision data?
Read the figure
The horizontal axis names four frozen matrix fixtures. The vertical axis is maximum absolute difference between the actual simulated Pallas output and the original float32 dataset calculation. Float32 bars have zero height; bfloat16 bars are positive because stored inputs and final outputs are rounded.
The bfloat16 gaps are approximately , , and . The largest matrix does not have the largest gap in this run. These are different frozen inputs, so the bar heights are separate from a monotonic law relating shape and error.
Connect it to the computation
The separate represented-input oracle checks in the code report zero implementation error for every shape and dtype. This plot answers the additional precision question: what changed relative to the original float32 values despite correct execution of the chosen contract?
The figure contains no runtime or target-memory data. It is distinct from throughput, pipeline overlap or TPU correctness. Those claims require the guarded target runner and actual completed accelerator execution.
Open full-size figureDownload data and provenance
The notebook contains the data experiment, plotting code, and recorded figure output. Rerun it to test your changes.
Where to put & run this code: Work in phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.py (or downloaded exercises/kernels-04.py / exercises/kernels-04.ipynb) and execute from the repository root with python3 phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.py.
1 Keep reference semantics, target mode and benchmark boundary together
Create main.py for the first block, then append each block in order in your course CPU environment.
# Step 1 — Keep reference semantics, target mode and benchmark boundary together: The candidate uses the same checked Pallas pipeline.
import jax
import jax.numpy as jnp
import numpy as np
from jax.experimental import pallas as pl
from jax.experimental.pallas import tpu as pltpu
def validate_pair(x,y,block):
# Guard input contract (`x.ndim != 2 or x.shape != y.shape or min(x.shape) < 1`) and fail fast if violated.
if x.ndim!=2 or x.shape!=y.shape or min(x.shape)<1:
raise ValueError("inputs must be nonempty matrices with identical shapes")
# Guard input contract (`x.dtype != y.dtype or x.dtype not in (jnp.float32, jnp.bfloat16)`) and fail fast if violated.
if x.dtype!=y.dtype or x.dtype not in (jnp.float32,jnp.bfloat16):
raise ValueError("matching float32 or bfloat16 inputs required")
# Guard input contract (`len(block) != 2 or any((not isinstance(n, int) or n < 1 for n in block))`) and fail fast if violated.
if len(block)!=2 or any(not isinstance(n,int) or n<1 for n in block):
raise ValueError("two positive integer block dimensions required")
def pad_pair(x,y,block):
validate_pair(x,y,block)
m,n=x.shape
bm,bn=block
padded=((m+bm-1)//bm*bm,(n+bn-1)//bn*bn)
# Combine or mask array elements to form `pads`.
pads=((0,padded[0]-m),(0,padded[1]-n))
return jnp.pad(x,pads),jnp.pad(y,pads),padded
def axpy_body(x_ref,y_ref,out_ref):
result=2.0*x_ref[...].astype(jnp.float32)+y_ref[...].astype(jnp.float32)
out_ref[...]=result.astype(out_ref.dtype)
def blocked_axpy(x,y,block=(2,4)):
# Combine or mask array elements to form `(px, py, padded)`.
px,py,padded=pad_pair(x,y,block)
bm,bn=block
# Invoke custom Pallas kernel or tile specification (`spec`).
spec=pl.BlockSpec(block,lambda i,j:(i,j))
# Combine or mask array elements to form `out`.
out=pl.pallas_call(axpy_body,
out_shape=jax.ShapeDtypeStruct(padded,x.dtype),
grid=(padded[0]//bm,padded[1]//bn),
in_specs=(spec,spec),out_specs=spec,interpret=True)(px,py)
return out[:x.shape[0],:x.shape[1]]
def pipelined_axpy(x,y,block=(8,128),buffers=2,no_pipelining=False,mode="simulate"):
# Combine or mask array elements to form `(px, py, padded)`.
px,py,padded=pad_pair(x,y,block)
bm,bn=block
# Guard input contract (`bm % 8 or bn % 128`) and fail fast if violated.
if bm%8 or bn%128:
raise ValueError("pipeline blocks must be multiples of (8,128)")
# Guard input contract (`buffers not in (2, 3)`) and fail fast if violated.
if buffers not in (2,3):
raise ValueError("this lab supports two or three buffers")
# Guard input contract (`mode not in ('simulate', 'tpu')`) and fail fast if violated.
if mode not in ("simulate","tpu"):
raise ValueError("mode must explicitly be simulate or tpu")
# Guard input contract (`mode == 'tpu' and (jax.default_backend() != 'tpu' or any((not isinstance(a, jax.core.Tracer) and any((d.platform != 'tpu' for d in a.devices())) for a in (x, y))))`) and fail fast if violated.
if mode=="tpu" and (jax.default_backend()!="tpu" or any(not isinstance(a,jax.core.Tracer) and any(d.platform!="tpu" for d in a.devices()) for a in (x,y))):
raise RuntimeError("Real TPU inputs/backend required; simulation fallback is disabled")
# Invoke custom Pallas kernel or tile specification (`spec`).
spec=pl.BlockSpec(block,lambda i,j:(i,j),pipeline_mode=pl.Buffered(buffer_count=buffers))
# Invoke custom Pallas kernel or tile specification (`output_spec`).
output_spec=pl.BlockSpec(block,lambda i,j:(i,j),pipeline_mode=pl.Buffered(buffer_count=2))
def outer(x_hbm,y_hbm,out_hbm):
# Combine or mask array elements to form ``.
pltpu.emit_pipeline(axpy_body,grid=(padded[0]//bm,padded[1]//bn),
in_specs=(spec,spec),out_specs=output_spec,
no_pipelining=no_pipelining)(x_hbm,y_hbm,out_hbm)
# Invoke custom Pallas kernel or tile specification (`whole`).
whole=pl.BlockSpec(memory_space=pl.ANY)
interpretation=pltpu.InterpretParams(detect_races=True) if mode=="simulate" else False
# Combine or mask array elements to form `call`.
call=pl.pallas_call(outer,out_shape=jax.ShapeDtypeStruct(padded,x.dtype),
in_specs=(whole,whole),out_specs=whole,interpret=interpretation)
if mode=="simulate":
# This describes a SIMULATED TPU layout, not the machine running the code.
abstract=jax.sharding.AbstractMesh((1,),("simulated_device",),
abstract_device=jax.sharding.AbstractDevice(device_kind="TPU v5 lite",num_cores=1))
with jax.sharding.use_abstract_mesh(abstract):
output=call(px,py)
else:
output=call(px,py)
return output[:x.shape[0],:x.shape[1]]
def target_benchmark(candidate,baseline,args,repeats=20):
# Guard input contract (`jax.default_backend() != 'tpu' or any((d.platform != 'tpu' for a in args for d in a.devices()))`) and fail fast if violated.
if jax.default_backend()!="tpu" or any(d.platform!="tpu" for a in args for d in a.devices()):
raise RuntimeError("TPU benchmark requires real TPU inputs; interpretation timings are not accepted")
# Guard input contract (`repeats < 5`) and fail fast if violated.
if repeats<5:
raise ValueError("at least five repeated target measurements required")
import time
records={}
for label,function in (("candidate",candidate),("baseline",baseline)):
started=time.perf_counter()
# Wrap with `jax.jit` (`compiled`) so XLA traces and compiles the function.
compiled=jax.jit(function).lower(*args).compile()
compile_seconds=time.perf_counter()-started
# Synchronize host execution until asynchronous device computation completes.
for _ in range(5):compiled(*args).block_until_ready()
samples=[]
for _ in range(repeats):
started=time.perf_counter()
# Synchronize host execution until asynchronous device computation completes.
compiled(*args).block_until_ready()
samples.append((time.perf_counter()-started)*1000)
records[label]={"compile_seconds":compile_seconds,"samples_ms":samples,
"median_ms":float(np.median(samples)),"p90_ms":float(np.percentile(samples,90))}
# Check which hardware backend (`cpu`, `gpu`, or `tpu`) JAX selected for `records['actual_backend']`.
records["actual_backend"]=jax.default_backend()
# Query the active JAX devices into `records['device_kind']`.
records["device_kind"]=jax.devices()[0].device_kind
# Combine or mask array elements to form `records['boundary']`.
records["boundary"]="already-placed logical inputs; complete padded/cropped wrapper output; warm synchronized execution"
return recordsCode building blocks in this step
jax.jit(fn) / @jax.jit— Tracesfnwith abstract shapes and compiles a fused XLA executable cached by input shape and dtype.jax.block_until_ready(output)— Synchronizes with the accelerator/CPU device so asynchronous dispatch finishes before wall-clock timing.
The candidate uses the same checked Pallas pipeline. The benchmark rejects CPU inputs before compiling or timing. It measures the complete wrapper, including padding/cropping, and records compilation separately.
2 Audit a matrix of shapes and precisions
Create main.py for the first block, then append each block in order in your course CPU environment.
# Step 2 — Audit a matrix of shapes and precisions: A zero kernel error means the implementation honored the declared...
rng=np.random.default_rng(42)
correctness=[]
precision_gaps=[]
for shape in [(1,1),(8,128),(9,129),(17,257)]:
original_x=rng.normal(size=shape).astype(np.float32)
original_y=rng.normal(size=shape).astype(np.float32)
ideal=2*original_x+original_y
row=[]
gaps=[]
for dtype in (jnp.float32,jnp.bfloat16):
# Create device-backed JAX array `a`.
# Create device-backed JAX array `b`.
a=jnp.asarray(original_x,dtype=dtype)
b=jnp.asarray(original_y,dtype=dtype)
actual=pipelined_axpy(a,b)
represented=(2*a.astype(jnp.float32)+b.astype(jnp.float32)).astype(dtype)
# Convert `` to a host NumPy array for inspection or verification.
np.testing.assert_array_equal(np.asarray(actual),np.asarray(represented))
# Convert `error` to a host NumPy array for inspection or verification.
error=float(np.max(np.abs(np.asarray(actual,dtype=np.float32)-np.asarray(represented,dtype=np.float32))))
# Convert `gap` to a host NumPy array for inspection or verification.
gap=float(np.max(np.abs(np.asarray(actual,dtype=np.float32)-ideal)))
row.append(error)
gaps.append(gap)
correctness.append(row)
precision_gaps.append(gaps)
print("Kernel errors versus represented-input oracle:",correctness)
print("Gaps versus original float32 data:",precision_gaps)Code building blocks in this step
random.default_rng(...)— Callrandom.default_rngwith your updated parameters or inputs from this lesson's workspace.rng.normal(...)— Callrng.normalwith your updated parameters or inputs from this lesson's workspace.
A zero kernel error means the implementation honored the declared precision contract. A nonzero gap to the original float32 dataset measures representation and output rounding, which must be reported separately.
3 Check the evidence boundary before requesting a speed claim
Create main.py for the first block, then append each block in order in your course CPU environment.
# Step 3 — Check the evidence boundary before requesting a speed claim: The runnable target entry point is the project target_tpu.py script.
if jax.default_backend()=="cpu":
probe=jnp.ones((8,128),dtype=jnp.float32)
rejected=False
try:target_benchmark(lambda a,b:pipelined_axpy(a,b,mode="tpu"),lambda a,b:2*a+b,(probe,probe))
except RuntimeError as error:
rejected=True
print("Expected benchmark refusal:",error)
assert rejected
print("CPU correctness audit complete; no TPU latency or speedup measured")Code building blocks in this step
jnp.zeros / jnp.ones(shape, dtype=...)— Allocates a tensor of the givenshapeinitialized with constants.
The runnable target entry point is the project target_tpu.py script. A missing target receipt remains a missing target receipt; the CPU result is not relabeled as performance data.
View the cumulative build in main.py
These blocks append in order to the same file.
# Step 1 — Keep reference semantics, target mode and benchmark boundary together: The candidate uses the same checked Pallas pipeline.
import jax
import jax.numpy as jnp
import numpy as np
from jax.experimental import pallas as pl
from jax.experimental.pallas import tpu as pltpu
def validate_pair(x,y,block):
# Guard input contract (`x.ndim != 2 or x.shape != y.shape or min(x.shape) < 1`) and fail fast if violated.
if x.ndim!=2 or x.shape!=y.shape or min(x.shape)<1:
raise ValueError("inputs must be nonempty matrices with identical shapes")
# Guard input contract (`x.dtype != y.dtype or x.dtype not in (jnp.float32, jnp.bfloat16)`) and fail fast if violated.
if x.dtype!=y.dtype or x.dtype not in (jnp.float32,jnp.bfloat16):
raise ValueError("matching float32 or bfloat16 inputs required")
# Guard input contract (`len(block) != 2 or any((not isinstance(n, int) or n < 1 for n in block))`) and fail fast if violated.
if len(block)!=2 or any(not isinstance(n,int) or n<1 for n in block):
raise ValueError("two positive integer block dimensions required")
def pad_pair(x,y,block):
validate_pair(x,y,block)
m,n=x.shape
bm,bn=block
padded=((m+bm-1)//bm*bm,(n+bn-1)//bn*bn)
# Combine or mask array elements to form `pads`.
pads=((0,padded[0]-m),(0,padded[1]-n))
return jnp.pad(x,pads),jnp.pad(y,pads),padded
def axpy_body(x_ref,y_ref,out_ref):
result=2.0*x_ref[...].astype(jnp.float32)+y_ref[...].astype(jnp.float32)
out_ref[...]=result.astype(out_ref.dtype)
def blocked_axpy(x,y,block=(2,4)):
# Combine or mask array elements to form `(px, py, padded)`.
px,py,padded=pad_pair(x,y,block)
bm,bn=block
# Invoke custom Pallas kernel or tile specification (`spec`).
spec=pl.BlockSpec(block,lambda i,j:(i,j))
# Combine or mask array elements to form `out`.
out=pl.pallas_call(axpy_body,
out_shape=jax.ShapeDtypeStruct(padded,x.dtype),
grid=(padded[0]//bm,padded[1]//bn),
in_specs=(spec,spec),out_specs=spec,interpret=True)(px,py)
return out[:x.shape[0],:x.shape[1]]
def pipelined_axpy(x,y,block=(8,128),buffers=2,no_pipelining=False,mode="simulate"):
# Combine or mask array elements to form `(px, py, padded)`.
px,py,padded=pad_pair(x,y,block)
bm,bn=block
# Guard input contract (`bm % 8 or bn % 128`) and fail fast if violated.
if bm%8 or bn%128:
raise ValueError("pipeline blocks must be multiples of (8,128)")
# Guard input contract (`buffers not in (2, 3)`) and fail fast if violated.
if buffers not in (2,3):
raise ValueError("this lab supports two or three buffers")
# Guard input contract (`mode not in ('simulate', 'tpu')`) and fail fast if violated.
if mode not in ("simulate","tpu"):
raise ValueError("mode must explicitly be simulate or tpu")
# Guard input contract (`mode == 'tpu' and (jax.default_backend() != 'tpu' or any((not isinstance(a, jax.core.Tracer) and any((d.platform != 'tpu' for d in a.devices())) for a in (x, y))))`) and fail fast if violated.
if mode=="tpu" and (jax.default_backend()!="tpu" or any(not isinstance(a,jax.core.Tracer) and any(d.platform!="tpu" for d in a.devices()) for a in (x,y))):
raise RuntimeError("Real TPU inputs/backend required; simulation fallback is disabled")
# Invoke custom Pallas kernel or tile specification (`spec`).
spec=pl.BlockSpec(block,lambda i,j:(i,j),pipeline_mode=pl.Buffered(buffer_count=buffers))
# Invoke custom Pallas kernel or tile specification (`output_spec`).
output_spec=pl.BlockSpec(block,lambda i,j:(i,j),pipeline_mode=pl.Buffered(buffer_count=2))
def outer(x_hbm,y_hbm,out_hbm):
# Combine or mask array elements to form ``.
pltpu.emit_pipeline(axpy_body,grid=(padded[0]//bm,padded[1]//bn),
in_specs=(spec,spec),out_specs=output_spec,
no_pipelining=no_pipelining)(x_hbm,y_hbm,out_hbm)
# Invoke custom Pallas kernel or tile specification (`whole`).
whole=pl.BlockSpec(memory_space=pl.ANY)
interpretation=pltpu.InterpretParams(detect_races=True) if mode=="simulate" else False
# Combine or mask array elements to form `call`.
call=pl.pallas_call(outer,out_shape=jax.ShapeDtypeStruct(padded,x.dtype),
in_specs=(whole,whole),out_specs=whole,interpret=interpretation)
if mode=="simulate":
# This describes a SIMULATED TPU layout, not the machine running the code.
abstract=jax.sharding.AbstractMesh((1,),("simulated_device",),
abstract_device=jax.sharding.AbstractDevice(device_kind="TPU v5 lite",num_cores=1))
with jax.sharding.use_abstract_mesh(abstract):
output=call(px,py)
else:
output=call(px,py)
return output[:x.shape[0],:x.shape[1]]
def target_benchmark(candidate,baseline,args,repeats=20):
# Guard input contract (`jax.default_backend() != 'tpu' or any((d.platform != 'tpu' for a in args for d in a.devices()))`) and fail fast if violated.
if jax.default_backend()!="tpu" or any(d.platform!="tpu" for a in args for d in a.devices()):
raise RuntimeError("TPU benchmark requires real TPU inputs; interpretation timings are not accepted")
# Guard input contract (`repeats < 5`) and fail fast if violated.
if repeats<5:
raise ValueError("at least five repeated target measurements required")
import time
records={}
for label,function in (("candidate",candidate),("baseline",baseline)):
started=time.perf_counter()
# Wrap with `jax.jit` (`compiled`) so XLA traces and compiles the function.
compiled=jax.jit(function).lower(*args).compile()
compile_seconds=time.perf_counter()-started
# Synchronize host execution until asynchronous device computation completes.
for _ in range(5):compiled(*args).block_until_ready()
samples=[]
for _ in range(repeats):
started=time.perf_counter()
# Synchronize host execution until asynchronous device computation completes.
compiled(*args).block_until_ready()
samples.append((time.perf_counter()-started)*1000)
records[label]={"compile_seconds":compile_seconds,"samples_ms":samples,
"median_ms":float(np.median(samples)),"p90_ms":float(np.percentile(samples,90))}
# Check which hardware backend (`cpu`, `gpu`, or `tpu`) JAX selected for `records['actual_backend']`.
records["actual_backend"]=jax.default_backend()
# Query the active JAX devices into `records['device_kind']`.
records["device_kind"]=jax.devices()[0].device_kind
# Combine or mask array elements to form `records['boundary']`.
records["boundary"]="already-placed logical inputs; complete padded/cropped wrapper output; warm synchronized execution"
return records
# Step 2 — Audit a matrix of shapes and precisions: A zero kernel error means the implementation honored the declared...
rng=np.random.default_rng(42)
correctness=[]
precision_gaps=[]
for shape in [(1,1),(8,128),(9,129),(17,257)]:
original_x=rng.normal(size=shape).astype(np.float32)
original_y=rng.normal(size=shape).astype(np.float32)
ideal=2*original_x+original_y
row=[]
gaps=[]
for dtype in (jnp.float32,jnp.bfloat16):
# Create device-backed JAX array `a`.
# Create device-backed JAX array `b`.
a=jnp.asarray(original_x,dtype=dtype)
b=jnp.asarray(original_y,dtype=dtype)
actual=pipelined_axpy(a,b)
represented=(2*a.astype(jnp.float32)+b.astype(jnp.float32)).astype(dtype)
# Convert `` to a host NumPy array for inspection or verification.
np.testing.assert_array_equal(np.asarray(actual),np.asarray(represented))
# Convert `error` to a host NumPy array for inspection or verification.
error=float(np.max(np.abs(np.asarray(actual,dtype=np.float32)-np.asarray(represented,dtype=np.float32))))
# Convert `gap` to a host NumPy array for inspection or verification.
gap=float(np.max(np.abs(np.asarray(actual,dtype=np.float32)-ideal)))
row.append(error)
gaps.append(gap)
correctness.append(row)
precision_gaps.append(gaps)
print("Kernel errors versus represented-input oracle:",correctness)
print("Gaps versus original float32 data:",precision_gaps)
# Step 3 — Check the evidence boundary before requesting a speed claim: The runnable target entry point is the project target_tpu.py script.
if jax.default_backend()=="cpu":
probe=jnp.ones((8,128),dtype=jnp.float32)
rejected=False
try:target_benchmark(lambda a,b:pipelined_axpy(a,b,mode="tpu"),lambda a,b:2*a+b,(probe,probe))
except RuntimeError as error:
rejected=True
print("Expected benchmark refusal:",error)
assert rejected
print("CPU correctness audit complete; no TPU latency or speedup measured")Run the example
python3 phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.pyAbout this command
This shortcut uses an activated macOS/Linux course environment in the repository folder. For Windows, use the explicit interpreter command in “Setup help” above. The script includes worked examples and reference solutions.
Inspect the complete assembled example
# Step 1 — Keep reference semantics, target mode and benchmark boundary together: The candidate uses the same checked Pallas pipeline.
import jax
import jax.numpy as jnp
import numpy as np
from jax.experimental import pallas as pl
from jax.experimental.pallas import tpu as pltpu
def validate_pair(x,y,block):
# Guard input contract (`x.ndim != 2 or x.shape != y.shape or min(x.shape) < 1`) and fail fast if violated.
if x.ndim!=2 or x.shape!=y.shape or min(x.shape)<1:
raise ValueError("inputs must be nonempty matrices with identical shapes")
# Guard input contract (`x.dtype != y.dtype or x.dtype not in (jnp.float32, jnp.bfloat16)`) and fail fast if violated.
if x.dtype!=y.dtype or x.dtype not in (jnp.float32,jnp.bfloat16):
raise ValueError("matching float32 or bfloat16 inputs required")
# Guard input contract (`len(block) != 2 or any((not isinstance(n, int) or n < 1 for n in block))`) and fail fast if violated.
if len(block)!=2 or any(not isinstance(n,int) or n<1 for n in block):
raise ValueError("two positive integer block dimensions required")
def pad_pair(x,y,block):
validate_pair(x,y,block)
m,n=x.shape
bm,bn=block
padded=((m+bm-1)//bm*bm,(n+bn-1)//bn*bn)
# Combine or mask array elements to form `pads`.
pads=((0,padded[0]-m),(0,padded[1]-n))
return jnp.pad(x,pads),jnp.pad(y,pads),padded
def axpy_body(x_ref,y_ref,out_ref):
result=2.0*x_ref[...].astype(jnp.float32)+y_ref[...].astype(jnp.float32)
out_ref[...]=result.astype(out_ref.dtype)
def blocked_axpy(x,y,block=(2,4)):
# Combine or mask array elements to form `(px, py, padded)`.
px,py,padded=pad_pair(x,y,block)
bm,bn=block
# Invoke custom Pallas kernel or tile specification (`spec`).
spec=pl.BlockSpec(block,lambda i,j:(i,j))
# Combine or mask array elements to form `out`.
out=pl.pallas_call(axpy_body,
out_shape=jax.ShapeDtypeStruct(padded,x.dtype),
grid=(padded[0]//bm,padded[1]//bn),
in_specs=(spec,spec),out_specs=spec,interpret=True)(px,py)
return out[:x.shape[0],:x.shape[1]]
def pipelined_axpy(x,y,block=(8,128),buffers=2,no_pipelining=False,mode="simulate"):
# Combine or mask array elements to form `(px, py, padded)`.
px,py,padded=pad_pair(x,y,block)
bm,bn=block
# Guard input contract (`bm % 8 or bn % 128`) and fail fast if violated.
if bm%8 or bn%128:
raise ValueError("pipeline blocks must be multiples of (8,128)")
# Guard input contract (`buffers not in (2, 3)`) and fail fast if violated.
if buffers not in (2,3):
raise ValueError("this lab supports two or three buffers")
# Guard input contract (`mode not in ('simulate', 'tpu')`) and fail fast if violated.
if mode not in ("simulate","tpu"):
raise ValueError("mode must explicitly be simulate or tpu")
# Guard input contract (`mode == 'tpu' and (jax.default_backend() != 'tpu' or any((not isinstance(a, jax.core.Tracer) and any((d.platform != 'tpu' for d in a.devices())) for a in (x, y))))`) and fail fast if violated.
if mode=="tpu" and (jax.default_backend()!="tpu" or any(not isinstance(a,jax.core.Tracer) and any(d.platform!="tpu" for d in a.devices()) for a in (x,y))):
raise RuntimeError("Real TPU inputs/backend required; simulation fallback is disabled")
# Invoke custom Pallas kernel or tile specification (`spec`).
spec=pl.BlockSpec(block,lambda i,j:(i,j),pipeline_mode=pl.Buffered(buffer_count=buffers))
# Invoke custom Pallas kernel or tile specification (`output_spec`).
output_spec=pl.BlockSpec(block,lambda i,j:(i,j),pipeline_mode=pl.Buffered(buffer_count=2))
def outer(x_hbm,y_hbm,out_hbm):
# Combine or mask array elements to form ``.
pltpu.emit_pipeline(axpy_body,grid=(padded[0]//bm,padded[1]//bn),
in_specs=(spec,spec),out_specs=output_spec,
no_pipelining=no_pipelining)(x_hbm,y_hbm,out_hbm)
# Invoke custom Pallas kernel or tile specification (`whole`).
whole=pl.BlockSpec(memory_space=pl.ANY)
interpretation=pltpu.InterpretParams(detect_races=True) if mode=="simulate" else False
# Combine or mask array elements to form `call`.
call=pl.pallas_call(outer,out_shape=jax.ShapeDtypeStruct(padded,x.dtype),
in_specs=(whole,whole),out_specs=whole,interpret=interpretation)
if mode=="simulate":
# This describes a SIMULATED TPU layout, not the machine running the code.
abstract=jax.sharding.AbstractMesh((1,),("simulated_device",),
abstract_device=jax.sharding.AbstractDevice(device_kind="TPU v5 lite",num_cores=1))
with jax.sharding.use_abstract_mesh(abstract):
output=call(px,py)
else:
output=call(px,py)
return output[:x.shape[0],:x.shape[1]]
def target_benchmark(candidate,baseline,args,repeats=20):
# Guard input contract (`jax.default_backend() != 'tpu' or any((d.platform != 'tpu' for a in args for d in a.devices()))`) and fail fast if violated.
if jax.default_backend()!="tpu" or any(d.platform!="tpu" for a in args for d in a.devices()):
raise RuntimeError("TPU benchmark requires real TPU inputs; interpretation timings are not accepted")
# Guard input contract (`repeats < 5`) and fail fast if violated.
if repeats<5:
raise ValueError("at least five repeated target measurements required")
import time
records={}
for label,function in (("candidate",candidate),("baseline",baseline)):
started=time.perf_counter()
# Wrap with `jax.jit` (`compiled`) so XLA traces and compiles the function.
compiled=jax.jit(function).lower(*args).compile()
compile_seconds=time.perf_counter()-started
# Synchronize host execution until asynchronous device computation completes.
for _ in range(5):compiled(*args).block_until_ready()
samples=[]
for _ in range(repeats):
started=time.perf_counter()
# Synchronize host execution until asynchronous device computation completes.
compiled(*args).block_until_ready()
samples.append((time.perf_counter()-started)*1000)
records[label]={"compile_seconds":compile_seconds,"samples_ms":samples,
"median_ms":float(np.median(samples)),"p90_ms":float(np.percentile(samples,90))}
# Check which hardware backend (`cpu`, `gpu`, or `tpu`) JAX selected for `records['actual_backend']`.
records["actual_backend"]=jax.default_backend()
# Query the active JAX devices into `records['device_kind']`.
records["device_kind"]=jax.devices()[0].device_kind
# Combine or mask array elements to form `records['boundary']`.
records["boundary"]="already-placed logical inputs; complete padded/cropped wrapper output; warm synchronized execution"
return records
# Step 2 — Audit a matrix of shapes and precisions: A zero kernel error means the implementation honored the declared...
rng=np.random.default_rng(42)
correctness=[]
precision_gaps=[]
for shape in [(1,1),(8,128),(9,129),(17,257)]:
original_x=rng.normal(size=shape).astype(np.float32)
original_y=rng.normal(size=shape).astype(np.float32)
ideal=2*original_x+original_y
row=[]
gaps=[]
for dtype in (jnp.float32,jnp.bfloat16):
# Create device-backed JAX array `a`.
# Create device-backed JAX array `b`.
a=jnp.asarray(original_x,dtype=dtype)
b=jnp.asarray(original_y,dtype=dtype)
actual=pipelined_axpy(a,b)
represented=(2*a.astype(jnp.float32)+b.astype(jnp.float32)).astype(dtype)
# Convert `` to a host NumPy array for inspection or verification.
np.testing.assert_array_equal(np.asarray(actual),np.asarray(represented))
# Convert `error` to a host NumPy array for inspection or verification.
error=float(np.max(np.abs(np.asarray(actual,dtype=np.float32)-np.asarray(represented,dtype=np.float32))))
# Convert `gap` to a host NumPy array for inspection or verification.
gap=float(np.max(np.abs(np.asarray(actual,dtype=np.float32)-ideal)))
row.append(error)
gaps.append(gap)
correctness.append(row)
precision_gaps.append(gaps)
print("Kernel errors versus represented-input oracle:",correctness)
print("Gaps versus original float32 data:",precision_gaps)
# Step 3 — Check the evidence boundary before requesting a speed claim: The runnable target entry point is the project target_tpu.py script.
if jax.default_backend()=="cpu":
probe=jnp.ones((8,128),dtype=jnp.float32)
rejected=False
try:target_benchmark(lambda a,b:pipelined_axpy(a,b,mode="tpu"),lambda a,b:2*a+b,(probe,probe))
except RuntimeError as error:
rejected=True
print("Expected benchmark refusal:",error)
assert rejected
print("CPU correctness audit complete; no TPU latency or speedup measured")Expected result: All eight shape/dtype cases exactly match the represented-input oracle. bfloat16 differs from original float32 values by a separately reported precision gap. The benchmark refuses CPU execution; no target speedup is reported.
Recorded reference execution
Reference output captured from executing the complete lesson script (code/main.py) and its assertions in the course CPU environment.
Kernel errors versus represented-input oracle: [[0.0, 0.0], [0.0, 0.0], [0.0, 0.0], [0.0, 0.0]]
Gaps versus original float32 data: [[0.0, 0.000862419605255127], [0.0, 0.025217056274414062], [0.0, 0.030788421630859375], [0.0, 0.0292205810546875]]
Expected benchmark refusal: TPU benchmark requires real TPU inputs; interpretation timings are not accepted
CPU correctness audit complete; no TPU latency or speedup measured
Kernel errors versus represented-input oracle: [[0.0, 0.0], [0.0, 0.0], [0.0, 0.0], [0.0, 0.0]]
Gaps versus original float32 data: [[0.0, 0.000862419605255127], [0.0, 0.025217056274414062], [0.0, 0.030788421630859375], [0.0, 0.0292205810546875]]
Expected benchmark refusal: TPU benchmark requires real TPU inputs; interpretation timings are not accepted
CPU correctness audit complete; no TPU latency or speedup measured
Joint row permutation preserves output identity
Early/final-only rounding: 1.0 1.0078125
Configuration, padded elements: (8, 128) 2 4096
Configuration, padded elements: (8, 128) 3 4096
Configuration, padded elements: (16, 256) 2 4096
Configuration, padded elements: (16, 256) 3 4096
Mixed and integer dtypes rejected
Modeled twofold-local speedup and zero-cost limit: 1.1764705882352942 1.4285714285714286
PASS: kernels-04
Download execution record · The notebook includes executed cells and figure outputs. Compare your run with these values; timings depend on the machine.
Preserve correctness under a joint row permutation
# Experiment — Preserve correctness under a joint row permutation: This operation has no coupling between rows.
a=jnp.asarray(rng.normal(size=(9,129)),dtype=jnp.float32)
# Create device-backed JAX array `b`.
b=jnp.asarray(rng.normal(size=(9,129)),dtype=jnp.float32)
order=jnp.array([8,0,7,1,6,2,5,3,4])
inverse=jnp.argsort(order)
original=pipelined_axpy(a,b)
permuted=pipelined_axpy(a[order],b[order])[inverse]
# Convert `` to a host NumPy array for inspection or verification.
np.testing.assert_array_equal(np.asarray(permuted),np.asarray(original))
print("Joint row permutation preserves output identity")Expected: The restored output is identical.
This operation has no coupling between rows. A failure suggests ownership or data-pairing errors rather than ordinary floating-point variability.
Make a precision-policy counterexample
# Experiment — Make a precision-policy counterexample: The order of representation changes is part of the numerical...
first=jnp.array([1.00390625],dtype=jnp.float32)
second=jnp.array([0.00390625],dtype=jnp.float32)
early=(first.astype(jnp.bfloat16).astype(jnp.float32)+second).astype(jnp.bfloat16)
late=(first+second).astype(jnp.bfloat16)
assert float(early[0])!=float(late[0])
print("Early/final-only rounding:",float(early[0]),float(late[0]))Expected: Early rounding yields 1.0 while one final cast yields 1.0078125.
The order of representation changes is part of the numerical algorithm. A reference must reflect the declared contract rather than a superficially similar expression.
Make it yours
Where to put your code: Append your solution at the bottom of phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.py (or inside the Make it yours cell of exercises/kernels-04.ipynb) and run python3 phases/13-kernels/04-iterate-with-correctness-and-performance-evidence/code/main.py to verify your prediction before opening the reference solution.
Add a (15,255) fixture, compare two tile choices and two buffer counts, and retain the same frozen inputs for every candidate. Report errors and padded work separately.
Show a writing guide and starter scaffold
How to write this exercise code: step-by-step recipe & starter scaffold
1. Key functions & syntax to use:
jnp.allclose(actual, expected, rtol=..., atol=...)— Checks that two arrays match elementwise within floating-point tolerance.
2. Step-by-step implementation plan:
- Create device-backed JAX array
b. - Convert
expectedto a host NumPy array for inspection or verification.
# Exercise solution: Add a (15,255) fixture, compare two tile choices and two buffer...
a = jnp.asarray(...) # TODO: compute a
# Create device-backed JAX array `b`.
b = jnp.asarray(...) # TODO: compute b
# Convert `expected` to a host NumPy array for inspection or verification.
expected = ... # TODO: compute expected
for block in [(8,128),(16,256)]:
for buffers in (2,3):
result = pipelined_axpy(...) # TODO: compute result
np.testing.assert_allclose(result,expected,rtol = ... # TODO: compute np.testing.assert_allclose(result,expected,rtol
gm = ... # TODO: compute gm
gn = ... # TODO: compute gn
print("Configuration, padded elements:",block,buffers,gm*gn*block[0]*block[1])
Reveal the reference solution
# Exercise solution: Add a (15,255) fixture, compare two tile choices and two buffer...
a=jnp.asarray(np.random.default_rng(8).normal(size=(15,255)),dtype=jnp.float32)
# Create device-backed JAX array `b`.
b=jnp.asarray(np.random.default_rng(9).normal(size=(15,255)),dtype=jnp.float32)
# Convert `expected` to a host NumPy array for inspection or verification.
expected=2*np.asarray(a)+np.asarray(b)
for block in [(8,128),(16,256)]:
for buffers in (2,3):
result=pipelined_axpy(a,b,block,buffers=buffers)
np.testing.assert_allclose(result,expected,rtol=1e-6,atol=1e-6)
gm=(15+block[0]-1)//block[0]
gn=(255+block[1]-1)//block[1]
print("Configuration, padded elements:",block,buffers,gm*gn*block[0]*block[1])Reject unsupported dtype mixing
Practice
Pass a float32 matrix and a bfloat16 matrix together, then pass integer arrays. Verify neither is silently coerced into an undocumented kernel contract.
Show a writing guide and starter scaffold
How to write this practice solution: step-by-step recipe & starter scaffold
1. Key functions & syntax to use:
jnp.zeros / jnp.ones(shape, dtype=...)— Allocates a tensor of the givenshapeinitialized with constants.
2. Step-by-step implementation plan:
- Run the boundary check and catch the expected exception:
# Reject unsupported dtype mixing (Practice): Supported dtype boundaries are a correctness decision.
for first,second in [(jnp.ones((8,128),jnp.float32),jnp.ones((8,128),jnp.bfloat16)),(jnp.ones((8,128),jnp.int32),jnp.ones((8,128),jnp.int32))]:
rejected = ... # TODO: compute rejected
# Run the boundary check and catch the expected exception:
try:pipelined_axpy(first,second)
except ValueError:rejected=True
assert rejected # TODO: complete assertion check
print("Mixed and integer dtypes rejected")
Show a hint
The wrapper promises matching float32 or bfloat16 inputs.
Reveal reference solution and reasoning
# Reject unsupported dtype mixing (Practice): Supported dtype boundaries are a correctness decision.
for first,second in [(jnp.ones((8,128),jnp.float32),jnp.ones((8,128),jnp.bfloat16)),(jnp.ones((8,128),jnp.int32),jnp.ones((8,128),jnp.int32))]:
rejected=False
# Run the boundary check and catch the expected exception:
try:pipelined_axpy(first,second)
except ValueError:rejected=True
assert rejected
print("Mixed and integer dtypes rejected")Supported dtype boundaries are a correctness decision. Adding a new dtype requires a reference, target support and justified tolerances.
Bound the end-to-end opportunity
Challenge
A profile attributes of time to a local operation. Derive the best end-to-end speedup if that operation becomes twice as fast, then the impossible-to-exceed limit if its cost vanished.
Show a writing guide and starter scaffold
How to write this practice solution: step-by-step recipe & starter scaffold
1. Key functions & syntax to use:
jnp.allclose(actual, expected, rtol=..., atol=...)— Checks that two arrays match elementwise within floating-point tolerance.
2. Step-by-step implementation plan:
- Execute
np.testing.assert_allclose(end_to_end,1.1764705882352942) - Execute
np.testing.assert_allclose(limit,1.4285714285714286)
# Bound the end-to-end opportunity (Challenge): These are conditional analytic estimates, not measured speedups.
fraction = ... # TODO: compute fraction
local_speedup = ... # TODO: compute local_speedup
end_to_end = ... # TODO: compute end_to_end
limit = ... # TODO: compute limit
# Execute `np.testing.assert_allclose(end_to_end,1.1764705882352942)`
np.testing.assert_allclose(end_to_end,1.1764705882352942)
# Execute `np.testing.assert_allclose(limit,1.4285714285714286)`
np.testing.assert_allclose(limit,1.4285714285714286)
print("Modeled twofold-local speedup and zero-cost limit:",end_to_end,limit)
Show a hint
Keep the remaining workload time unchanged in the calculation.
Reveal reference solution and reasoning
# Bound the end-to-end opportunity (Challenge): These are conditional analytic estimates, not measured speedups.
fraction=.3
local_speedup=2.0
end_to_end=1/((1-fraction)+fraction/local_speedup)
limit=1/(1-fraction)
# Execute `np.testing.assert_allclose(end_to_end,1.1764705882352942)`
np.testing.assert_allclose(end_to_end,1.1764705882352942)
# Execute `np.testing.assert_allclose(limit,1.4285714285714286)`
np.testing.assert_allclose(limit,1.4285714285714286)
print("Modeled twofold-local speedup and zero-cost limit:",end_to_end,limit)These are conditional analytic estimates, not measured speedups. Communication or scheduling changes can alter the fraction, so reprofile the complete workload.
Diagnose the result
If bfloat16 disagrees only with the original float32 data, check whether the represented-input oracle still passes before calling it a kernel bug. If timings are implausibly tiny, inspect synchronization and device placement. If a local speedup fails to improve the training step, reprofile communication and other unchanged work. Never replace missing target evidence with simulated-device labels.
Carry forward
- Choose cases that challenge the contract
- Use two references for two different questions
- Make target timing a complete operation
- Compare a strong baseline with the same contract
- Relate local kernels to the distributed workload
- Accept, revise or decline the optimization
Keep your evidence
Shape/dtype reference checks, separate precision gaps, guarded target benchmark and a bounded end-to-end speedup argument. Keep the environment, observed outputs and your explanation of the figure.
Save your modified code, its output, and an explanation of your prediction. Checkpoint answers and exercise evidence are tracked separately.
Saved in this browser. Export or back up your portfolio anytime in My learning.
Verified on local CPU. Stage to a Cloud TPU VM via Cloud TPU setup to record accelerator execution receipts (TODO: add automated Cloud TPU CI execution receipts).
Primary references
- Pallas grids and block specifications ↗
- Writing TPU kernels: layout and memory restrictions ↗
- TPU pipelining and emit_pipeline ↗
- TPU interpretation is simulation, not target execution ↗
Phase 13 Integration Project
Bring this phase together: A checked Pallas kernel with an explicit target qualification plan
Checkpoint to prove: Compare against a reference across shapes and justify each optimization with evidence.
Copy projects/kernel-audit/starter/model.py to projects/kernel-audit/my_model.py and write your code in projects/kernel-audit/my_model.py. Run python3 projects/kernel-audit/tests/check.py --implementation projects/kernel-audit/my_model.py --stage 1 from the top-level folder to verify all 3 stages.