IPDPS 2021 · arXiv:2010.10248

Temporal blocking of finite-difference stencil operators with sparse “off-the-grid” sources

This page is a companion to the paper (arXiv:2010.10248) — read that for the full derivations, proofs and benchmark methodology; this is a figure-by-figure walkthrough built for intuition. Moving data to and from memory costs time on every finite-difference wave solver, compute-bound or not, and temporal blocking cuts that cost by reusing cached data across timesteps. Real wave-propagation codes break the technique, because seismic sources and receivers sit between grid points. A short primer on time-tiling, then ten figures on why it breaks, how the paper puts it back together, and what the code actually says.

George Bisbas1 · Fabio Luporini2 · Mathias Louboutin3 · Rhodri Nelson1 · Gerard J. Gorman1 · Paul H. J. Kelly1

1 Imperial College London, UK 2 Devito Codes, London, UK 3 Georgia Institute of Technology, Atlanta, GA, USA

35th IEEE IPDPS, 17–21 May 2021, pp. 497–506  ·  arXiv:2010.10248  ·  IEEE Xplore

I wrote the paper five years ago. This page exists because I finally sat down and built the figures I wish I'd had back then — the tiling geometry, the sparse-source break, the code walkthrough — with Claude doing most of the drawing, instead of leaving it all as equations and prose.

How this was made

I had Claude Opus 5 build these figures, working from the published paper (arXiv:2010.10248), the IPDPS’21 slide deck, and the code listings and Devito snippets from the work. Every diagram here is redrawn for explanation rather than reproduced from the publication.

  • measurements Speedups, tile sizes, machine specifications and timestep counts trace to the paper and are cited in full at the foot of the page.
  • added analysis The tile-width model in Fig. 10, and the reading of Table I that goes with it, are not results from the paper. They are analysis added here and are flagged as such in place, including one point where the model and Table I disagree.
  • illustrative geometry Tile counts, block widths and skew values in the interactive figures are chosen for legibility, not to match any benchmarked configuration.

Check these claims against the paper before reusing any of this in a talk or draft.

The argument, in five moves
  1. 01

    Stencils tile well in time. Advance a cached block through several timesteps instead of one. Dense, affine, uniform dependences; decades of literature.

  2. 02

    Real propagators are not just stencils. Sources and receivers sit off the grid and are applied through a second loop with non-affine accesses.

  3. 03

    That second loop cannot be time-tiled. Not because the physics forbids it — because the loop has no valid position once tiles run ahead of one another.

  4. 04

    So remove the second loop. Observe which grid points a source touches, store the contribution on those points, and fold it into the stencil's own sweep.

  5. 05

    Now tile in time. Up to 1.6× over tuned spatially blocked code at low order, with a clear ceiling at high order.

Contents
  1. ··  Devito
  2. ··  The physics
  3. ··  What is time-tiling
  4. 01  The stencil
  5. 02  Off the grid
  6. 03  Skewing time
  7. 04  The violation
  8. 05  Both at once
  9. 06  Realignment
  10. 07  The code
  11. 08  The loop nest
  12. 09  What it bought
  13. 10  The squeeze
  14. ··  Open threads
SYMBOLS

The physics, written once

The ten figures that follow all turn on a single pattern: a residual written symbolically, rearranged into an explicit update, then handed to the compiler. The three propagators benchmarked in Fig. 09 are three instances of it, and the stencil radius that drives every legality argument on this page is set here, as space_order.

the model isotropic acoustic · continuous
  1. m · ∂²u/∂t² − ∇²u + damp · ∂u/∂t = q
  2. # m = 1/c² — squared slowness
  3. # q enters through src.inject (Fig. 07)
  4. # damp = absorbing boundary layer
the same, in Devito residual → solve → Eq
  1. eq = Eq(m * u.dt2 - u.laplace + damp * u.dt, 0)
  2. stencil = Eq(u.forward, solve(eq, u.forward))
  3. # solve() isolates u.forward — SymPy algebra,
  4. # not a hand-derived update (Fig. 07, b4)
What the two derivatives expand into
SymbolExpands toWhere it bites
u.dt2(u(t+1) − 2·u(t) + u(t−1)) / dt²time_order = 2 → three addressable timesteps
u.laplace∂²u/∂x² + ∂²u/∂y² + ∂²u/∂z², each a centred differencespace_order = so → stencil radius r = so/2
∂²u/∂x²Σ cₖ · u(x + k·h) for k ∈ [−r, r]the star of Fig. 01 · 6r + 1 points in 3D
one cₖone stencil weight on that arrowthe arrows whose reach the skew of Fig. 03 must cover
The other two propagators — same pattern, heavier symbols
KernelSymbolic formWhat changes
anisotropic acoustic (TTI) coupled u, v · second derivatives along tilted axes rotation by θ and anisotropy ε, δ inside the derivatives — the operation count climbs, which is why Fig. 09 gives it less back from cache
isotropic elastic velocity–stress system, first order in time 9 state fields instead of 1, and a shorter time dependence — the working set Fig. 10 counts multiplies by its 18 resident planes

Nothing on this page depends on the acoustic equation specifically. The tiling argument needs only that the update is explicit, local and affine — properties the compiler reads off the symbolic form rather than off any particular PDE. The same pipeline serves TTI and elastic in Fig. 09 without modification, and the precomputation of Fig. 06 observes the injection operator instead of modelling it, for the same reason.

PRIMER

Time-tiling, briefly

Temporal blocking reorders the loop nest so that a cached block is advanced through several timesteps before the next block is touched — one trip to memory serving many updates instead of one. It pays off because the field dwarfs the cache: 512³ in single precision is 512 MiB against roughly 48 MiB of last-level cache, with hundreds of timesteps to get through. The transformation is a strip-mine and interchange, sliding time inward:

spatial blocking time outermost
  1. for t in timesteps:
  2. for blk in blocks:
  3. for x in blk: update(t, x)
  4. # blk is reloaded on every t
temporal blocking time moved inside · naive
  1. for tt in time_tiles:
  2. for blk in blocks:
  3. for t in tt:
  4. for x in blk: update(t, x)
  5. # blk loaded once for all t in tt · its right edge reads the next blk, not yet advanced

Terminology follows the paper: a block is a spatial-only grouping; a tile is a block allowed to advance through several timesteps at once. Tiles therefore nest over blocks — the distinction the listings above, Table I and the loop nest of Fig. 08 all rely on.

Four ways to cut the space–time iteration space

Time-tiling is a family of methods, not a single one. Each panel below shows the same iteration space: every dot is one grid point at one timestep, and its colour says which tile owns it. The four drawn here are the ones this argument needs; the family also includes cache-oblivious space–time recursion — tiling recursively until a base case fits cache, with no cache size ever named in the code — (Frigo & Strumpen), hexagonal hybrid tiles (Grosser et al.) and 3.5D partial-dimension blocking (Nguyen et al.).

Sliding time inside the block loop — the naive interchange shown on the right — is not legal: a block's right-hand edge reads points in the next block, which runs after it and has not reached that timestep yet. The families above are the repairs, each restoring legality a different way — by leaning the tiles, by recomputing the shared margin, or by cutting cone-shaped tiles whose boundaries follow the dependence cone, so every tile can start at once — concurrent start, in the diamond panel below. This paper uses skewed wavefront tiling, which Fig. 03 covers. The problem it addresses is orthogonal to the repair you choose: none of these families can be applied unmodified once the simulation contains sources that are not on the grid: injection is additive, so tiles that overlap or lean across a seam would need ownership masks to avoid double-counting a contribution — the machinery Fig. 06 builds.

space — the grid, and dependences that hold time — timesteps, and space–time tiles sparse — sources and receivers, off the grid broken — a dependence that cannot be satisfied

Three figures step outside this: in the tiling families hue only tells tiles apart, the roofline follows the paper’s own red-and-yellow coding, and Fig. 09’s gains chart reuses these three colours for kernel identity — acoustic, TTI, elastic — not space/time/sparse.

FIG. 01

Every point depends on its neighbours

Discretise a PDE with explicit finite differences and every grid point at timestep t+1 becomes a weighted sum of its neighbours at t. The reach of that sum is the space order. It is dense, regular and predictable, and therefore tileable. The dependence pattern below is what every legality argument on this page refers back to.

1D · (2r+1)-point update u[t+1][x] = Σ cₖ·u[t][x+k] for k ∈ [−r, r] t+1 t 3D · star stencil radius r = space order / 2 · one arm per axis direction
both panels share one clock: the row's arrow-fan and the star's arms grow through the same space order together

Both panels share one clock: 2r + 1 neighbours at t feed one point at t+1 on the left, and the same radius grows a 3D star on the right — 6r + 1 points along its six arms. Every blocking scheme below has to respect those arrows. Watch the two grow together as the order cycles: 5 and 13 points at space order 4, 9 and 25 at order 8, 13 and 37 at order 12. The paper benchmarks all three, and Fig. 10 shows what that growth costs.

FIG. 02

Sources and receivers do not sit on the grid

A seismic or medical-imaging simulation is not a clean stencil benchmark. Energy is injected by sources and sampled by receivers whose physical coordinates are not aligned with the discretisation. Both sit on the same grid at the same time — a shot record has one source and thousands of receivers — and each spreads its contribution across the surrounding points through weights that vary continuously with position. Drag either marker and watch the weights move.

one grid, two kinds of operator, neither on it
drag either marker — the one nearest your pointer moves

Weights shown for a 2D slice; in 3D each operator touches eight points. Devito builds both from coordinates known only at runtime, so the accesses are non-affine.

Two iteration spaces share one field. The stencil walks the lattice in order; these two arrive through indirect accesses at coordinates that were never part of it. The arrows run opposite ways on purpose — injection scatters outward, interpolation gathers inward — because the receiver operator is the adjoint of the source operator, built from the same weights.

That symmetry matters later: in full waveform inversion (FWI) and reverse time migration (RTM) — the two seismic-imaging workloads this solver is built for — both appear in the same operator, so anything that fixes one has to fix the other. In time the two are aligned, acting on the same timestep index; in space neither is. They are not symmetric as schedules, though: injection writes, so it pins the order, while interpolation only reads and can be deferred — as far as the time buffer allows.

FIG. 03

Skewing the tiles

Spatial blocking cuts the domain into cache-sized tiles, then throws the data away at the end of every timestep. Temporal blocking keeps going — advance one tile through several timesteps before touching the next. That is only legal if everything a point reads has already been computed. Tiles run left to right, so the dangerous dependence is the one reaching right, into a tile that has not started. Skewing leans the tiles to pull it back.

Why the minimum skew is the stencil radius

Think of the stencil as a speed limit. A point at t+1 reads r cells to either side at t; those points read r cells further at t-1. Influence therefore spreads outward at exactly r cells per timestep — a discrete dependence cone, which the figure calls a light cone — and it widens whether you like it or not.

A tile is a promise to compute a region without consulting its right-hand neighbour. To keep that promise, the tile's right boundary has to retreat at least as fast as the cone widens. Retreat by less than r per timestep and the cone escapes; retreat by exactly r and the boundary and the cone travel together. Hence s ≥ r, where r = space_order / 2 — in tiling terms, the tile's shear must enclose the dependence vector (1, r). When a kernel has several dependences with different radii, the largest is the binding one.

The cost is the same geometry working against you. Skewing recomputes nothing — the iteration space is still partitioned, every point updated once. What it costs is footprint. Carried through T timesteps, a tile of width W sweeps across W + r·T columns while updating only W of them, so the volume it must stream grows with the lean even though the work does not. On top of that, tiles start in sequence rather than together, so there is pipeline fill and drain. That is the tension: a taller time block buys reuse, a wider stencil spends it, and by space order 12 the second term wins. Fig. 10 puts numbers on it.

The dependence cone against the tile boundary same point, same stencil (r = 1) — only the boundary differs tt+1t+2 boundary fixed · s = 0 2r r cells past the boundary 2r past, and growing The cone escapes after one timestep. Those values belong to the next tile, which has not run — the tile cannot keep its promise. boundary recedes r per step · s = r slope = r cells per timestep +r +2r The boundary gives ground at exactly the rate the cone gains it. Everything the tile reads, the tile already owns. Skew more and it still works — just wastefully.
0
cells per timestep
legal when s ≥ r
space x → time t →
not started
computed not yet computed point being updated now read that is available read of data that does not exist

Every dot is a grid point at one timestep; every parallelogram is one space–time tile, numbered in execution order. Press play and the sweep runs the way the generated code does: tile ① is carried through all four timesteps before tile ② begins, so its data stays in cache for the whole time block. The white marker is the point being updated, and the arrows are the 2r+1 values it reads from the row below. At skew 0, watch the right-hand arrows go red every time the sweep reaches a tile's trailing edge: it is asking for values from a tile that has not run. Raise the skew to r and the counter stays at zero for the whole run.

Raise it beyond r and the run is still clean, but the tile now sweeps across more columns than it updates for no extra reuse. The minimum legal skew is r, so a wider stencil forces a steeper lean whether or not you want one. The same geometry explains the paper's numbers: 1.6× at space order 4, about 1.13× at order 8, down to a token 5–10% or none at all by order 12, depending on the kernel (Fig. 09).

FIG. 04

The source lands on a tile seam

Skewing settles the stencil's own dependences. It does nothing for the source, because the source is a separate loop over a separate iteration space — one pass per timestep, over all sources, across the whole domain. That was harmless without time-tiling, when the whole domain advanced in lockstep. Once tiles run ahead of each other, a single source whose footprint straddles a tile seam has to be applied at two different moments in the schedule. One loop cannot be in two places. It is also why a general-purpose tiler cannot rescue this: the injection loop’s accesses are non-affine, so the polyhedral frameworks — PLUTO, Polly, Loo.py and CLooG — tile the stencil beside it and stop at this one.

Where the source writes one source · four affected grid points · a 2D slice space x → time t → t t+1 t+2 t+3 t+4 1 tile A · runs first 2 tile B · runs after source, off the grid 2 points in A 2 points in B all four must receive their contribution at t+2 — but A and B are never at t+2 at the same moment When the code actually runs wall-clock order, not simulated time wall clock → tile A pass t → t+4, all at once tile B pass t → t+4, all at once inject sources at t+2 one loop, one position in the schedule needs to be here… …and here No legal placement exists. The dependence is not wrong physics — it is unschedulable while the sparse operator lives in its own iteration space. Fig. 05 shows why skewing cannot.

Left, one source scatters into four grid points, and the tile seam runs straight through them. Right, the same situation in execution order: tile A is advanced through its whole time block before tile B starts, so the single per-timestep injection loop would have to fire once inside A's pass and once inside B's pass. This is the applicability barrier the paper identifies. Nothing about wave propagation forbids temporal blocking; the sparse operator simply cannot be scheduled until it is folded into the grid's own iteration space.

FIG. 05

Skewing fixes the stencil, not the source

Fig. 03 showed skewing repairing the stencil's own dependences. Fig. 04 showed the source failing for a different reason. Put them on the same axes and the asymmetry becomes visible: the grid and the source are fixed in physical space, but the tiles lean, so the seams sweep leftward as time advances. A source sitting at one position gets crossed. Turning up the skew does not remove the crossing — it just moves which timestep it happens at.

2.0
stencil radius fixed at r = 2 · space order 4
x = 15.70
physical space x → time t → the grid and the source do not move · only the tiles lean
stencil dependences
source dependences

Grid points sit in fixed columns; the source's two affected points are marked at every timestep. Tile seams lean left as t grows, since that lean is the skew. As the sweep climbs, a seam eventually slides across the source. On the timesteps where it does, the two points that must receive the same contribution at the same simulated time belong to tiles running at different moments.

Watch the two lamps as you drag: the left one turns green at s ≥ 2 and stays green; the right one stays red no matter what you do to the skew, because the injection loop's placement is not a geometry problem. Only switching the injection to precomputed turns it green — which is Fig. 06.

FIG. 06

Precompute the injection, then fuse it into the sweep

The fix removes the second iteration space entirely. Inject into an empty field for a single timestep (a few more only if the wavelet starts at zero) to find out exactly which grid points a source can ever touch; record those points as on-grid arrays; then rewrite the injection as a masked term evaluated inside the stencil loop. The arithmetic is algebraically unchanged — the weights are pre-applied, so results may differ in the last bits — but it is now aligned, vectorisable and legal to time-tile, and the compiler generates it from a symbolic description rather than the user writing it by hand.

A Discover the footprint Inject into a zero-valued field for one timestep — a few more only if the wavelet starts at zero. Whatever is non-zero afterwards is a point some source can reach. before · sparse, off the grid coords (n_src, 3) · data (nt, n_src) inject t = 0…3 after · non-zero cells found these are grid points · they have integer indices record affected points, aligned to the grid (12, 7, 31) → id 0 (13, 7, 31) → id 1 (12, 8, 31) → id 2 (13, 8, 31) → id 3 coords (n_aff_pts, 3) · data (nt, n_aff_pts) B Re-encode as three on-grid arrays A slice through x and z. Every cell is a grid point; shaded cells receive a source contribution. C Fuse the loops The mask decides whether a point takes a contribution; the id says which one. The scatter becomes a local read. before · two iteration spaces for t: for x, y, z: u[t+1][x][y][z] = stencil(u[t], …) for t: for s in sources: scatter s into 8 neighbours ← indirect after · one iteration space for t: for x, y: for k in 0 … nnz[x][y] - 1: z = sp_sid[x][y][k]; id = sid[x][y][z] u[t+1][x][y][z] += sm[x][y][z] · src[t][id] for z: u[t+1][x][y][z] = stencil(u[t], …)
StructureOff-the-gridAlignedWhat changed
countn_srcn_aff_ptsa few sources become the points they touch
coordinates(n_src, 3)(n_aff_pts, 3)fractional → integer grid indices
data(nt, n_src)(nt, n_aff_pts)weights pre-applied, per point per timestep

Band A pays a one-off cost — a single timestep of injection into an empty field — to learn the footprint; the paper reports this as negligible next to the gain, and it holds for non-linear injection too, because it observes the operator rather than modelling it. Band B stores that footprint three ways: a binary mask, an identifier into the pre-weighted source values, and a per-column count of non-zeros so the sweep can skip empty z slices instead of testing every point.

Band C is the result: one loop nest, no indirect scatter, the interpolation hoisted into the precomputation, and a single iteration space that Fig. 03's skew can now be applied to. Legality rests on three conditions the paper states: coordinates fixed across the time domain, the mask and ID structures read-only once built, and the injection ordered after the stencil at the same timestep — which is what Inc preserves. This is the classic inspector/executor split from the sparse-compilation literature: Band A is the inspector, a cheap pass that resolves the data-dependent structure once; the fused nest is the executor, which runs many times against what the inspector found. Fig. 07 shows the Devito that produces it, line by line.

FIG. 07

The same physics, written twice

This is what the transformation costs the person writing the solver. On the left, the textbook Devito acoustic propagator — five lines, and the source injection is a single method call. On the right, the same model rewritten so that the injection is expressible inside the grid's own loop nest. Click any highlighted line to see what it does and why it is there.

baseline untileable
  1. u = TimeFunction(name="u", grid=model.grid,
  2. space_order=so, time_order=2)
  3. src_term = src.inject(field=u.forward,
  4. expr=src * dt**2 / model.m)
  5. pde = model.m * u.dt2 - u.laplace + model.damp * u.dt
  6. stencil = Eq(u.forward, solve(pde, u.forward))
  7. op = Operator([stencil, src_term])
transformed tileable
  1. # 1 · observe the footprint on an empty field
  2. f = TimeFunction(name="f", grid=model.grid, space_order=so,
  3. time_order=2)
  4. src_f = src.inject(field=f.forward, expr=src * dt**2 / model.m)
  5. Operator([src_f]).apply(time=3)
  6. nzinds = np.nonzero(f.data[0])
  7.  
  8. # 2 · build mask, ids, per-point wavelets, column counts
  9. # source_mask · source_id · save_src · nnz_sp_source_mask
  10.  
  11. # 3 · fuse the injection into the grid iteration space
  12. eq0 = Eq(sp_zi.symbolic_max, nnz_sp_source_mask[x, y] - 1,
  13. implicit_dims=(time, x, y))
  14. eq1 = Eq(zind, sp_source_mask[x, y, sp_zi],
  15. implicit_dims=(time, x, y, sp_zi))
  16. mask_expr = source_mask[x, y, zind] \
  17. * save_src[time, source_id[x, y, zind]]
  18. eq2 = Inc(usol.forward[t+1, x, y, zind], mask_expr,
  19. implicit_dims=(time, x, y, sp_zi))
  20.  
  21. pde_2 = model.m * usol.dt2 - usol.laplace + model.damp * usol.dt
  22. stencil_2 = Eq(usol.forward, solve(pde_2, usol.forward))
Select a highlighted line to see what it does.
The Devito objects used above

Roughly in the order they appear. Everything here is ordinary Python — Devito builds a symbolic description first, then generates and compiles the C from it.

Fields and operators
Grid
The discretised domain: shape, spacing, and the dimensions x, y, z and time that index everything else. Reached here as model.grid.
TimeFunction
A time-varying field on the Grid. space_order fixes the stencil radius and the halo; time_order fixes how many timesteps stay addressable — u.backward, u, u.forward — held in a small rolling buffer, not one slot per timestep.
Function
The same, without a time dimension. The mask, identifier and count arrays of Fig. 06 are these: grid-shaped, built once, read every timestep.
SparseTimeFunction
Values living at coordinates that need not lie on the grid at all — coordinates of shape (n_src, 3) and data of shape (nt, n_src). Sources and receivers are these, and this class is where the whole difficulty starts.
.inject()
Expands symbolically into a scatter: write this value into the surrounding grid points, weighted by how near each one is. The addresses come from runtime coordinates, so the accesses are non-affine.
.interpolate()
The adjoint — a gather rather than a scatter, same weights, used by receivers. Both are drawn in Fig. 02. Both operators appear in the same FWI or RTM run, which is why a fix has to cover both.
Equations and compilation
Eq
A symbolic equation. Given a left-hand side that is a field access, it becomes an assignment in the generated code.
solve
Rearranges a residual to isolate the unknown, so solve(pde, u.forward) yields the explicit update. The algebra is SymPy's, not hand-derived.
Inc
Like Eq, but an accumulation. Marking the injection as a reduction stops the compiler treating it as a plain write and reordering it against the stencil beside it.
Operator
Takes a list of equations, derives an iteration space for each from the dimensions it touches, fuses what it can, applies its optimisations, then generates and just-in-time compiles the C. Two equations with mismatched dimensions become two loop nests — the fact this whole paper turns on.
.apply()
Runs the compiled operator, with any runtime arguments. apply(time=3) is what advances the throwaway field far enough to reveal the footprint.
Dimension
A loop index. Most come from the Grid, but one can be declared independently — sp_zi here — to iterate over something the Grid knows nothing about.
implicit_dims
Forces an expression into a loop nest whose dimensions it does not syntactically mention. This is the lever that fuses the injection into the stencil's sweep instead of letting it become a nest of its own.
symbolic_max
Sets a dimension's upper bound from a value read at runtime, giving a loop whose trip count is data — usually zero here, which is how sparsity is expressed without paying for it.

The physics is untouched — pde and pde_2 are the same equation. Everything added on the right exists to move one thing: the source's contribution, out of its own loop and into the grid's. Two mechanisms do the work. implicit_dims forces an expression into a loop nest it does not syntactically belong to; symbolic_max lets a loop bound be a value read from an array at runtime. Together they let a sparse, data-dependent operation be written as a dense one. In the published release the operator is then built as op2 = Operator([eqxb, eqyb, eqxb2, eqyb2, stencil_2, eq0, eq1, eq2], opt=('advanced')), where each eqx… copies a runtime scalar out of a block_sizes Function — the tile and block shapes become runtime arguments, so Fig. 09's tuning sweeps need no C edits and no recompiles. That line is the entire compiler-facing surface of the technique.

FIG. 08

The generated loop nest

The loop nest after both transformations. Reading outwards: two innermost loops share one x, y pair, which is the fusion from Fig. 06; the skew appears as an index shift, which is the lean from Fig. 03; then blocks, tiles, and the time-tile loop. Note where the parallelism sits, and what is not blocked.

for t_tile in time_tiles: for xtile, ytile in tiles: for t in t_tile: for xblk, yblk in blocks: for x in xblk: for y in yblk: for z = 1 … nz: A(t, x - time, y - time, z) stencil update for z2 = 1 … nnz_mask[x][y]: I(t, x - time, y - time, z2) source injection one x, y pair two z loops #pragma omp for collapse(2) #pragma omp simd · z only Four things to notice 1 · the skew is an index x - time, y - time. Fig. 03's lean is not a data movement — it is a shift applied when addressing the array. 2 · z is never blocked The innermost dimension keeps its full stride so the vector units see contiguous memory. 3 · the injection loop is short Its trip count is nnz_mask[x][y] — usually zero. Columns with no source cost one comparison, not a sweep. 4 · parallelism sits inside the tile, on the blocked x and y loops, collapsed into one OpenMP iteration space with dynamic scheduling. Structure after Listing 6 of the paper.
What the compiler generates, and what was done by hand

The released code (tag v0.9-alpha) shows where the line is. Everything inside the two fused inner loops is produced by the Devito compiler, from the equations of Fig. 07 alone:

generated — symbolic lowering (solve), iteration-space derivation and fusion of the injection into the stencil nest (implicit_dims + Inc + symbolic_max), spatial blocking with runtime block sizes (a block_sizes Function feeds the xb_size… scalars), SIMD vectorisation on the unblocked z, OpenMP with collapse(2), and autotuning — configuration['autotuning'] = 'aggressive' for the baseline, and a hand-rolled sweep in acoustic_tune_model.py that reassigns block_sizes.data between runs, one compiled kernel serving the entire sweep.

hand-applied — the wavefront skew itself: the outer tile loops and the x - time shift, read from print(op2.ccode) and edited into the generated C. Full automation of the skew inside the DSL is listed as future work, and the paper is explicit about it.

What has landed since: Devito now exposes time_m/time_M sub-timestep execution and save=Buffer(M) (in the v2.0-ipdps-accepted fork, with a walk-through in examples/seismic/tutorials/12_time_blocking.ipynb), which automates the blocking half — not yet the skewed wavefront schedule itself.

Two loops share the innermost x, y position: the dense stencil sweep over z, and a short injection loop whose length is read from nnz_mask. That adjacency turns the source's data dependence into a local, affine one that the outer tiling can reason about. It also explains why the sparsity reduction in Fig. 06 matters so much: without nnz_mask, the second loop would run the full nz and multiply by zero almost every time. The block sizes in the drawing are not constants: in the released code they arrive as scalars read from a block_sizes Function, so the autotuner re-shapes the nest at runtime on a single compiled binary. Note the order: t sits outside the block loops, so every block of a tile advances in lockstep at each step — the wave order is between tiles, which is both why the OpenMP parallelism on blocks is legal and where the pipeline fill and drain come from.

FIG. 09

Measured performance

Three propagators of industrial significance — isotropic acoustic, anisotropic acoustic (TTI) and isotropic elastic — on 512³ grids over 512 ms, benchmarked on Azure Broadwell and Skylake VMs against Devito's own aggressively auto-tuned, vectorised, spatially blocked kernels. Both sides were swept over the full tile and block parameter space, so this is a tuned-versus-tuned comparison.

The three propagators
KernelFormTime orderState fieldsTimestepsCharacter
isotropic acousticsingle scalar PDE21228 Jacobi-like (each point updates from last timestep's values only) Laplacian, low arithmetic intensity (few flops per byte moved), firmly bandwidth-bound — the best case for temporal blocking, and it gains the most.
anisotropic acoustic (TTI)coupled system, 2 scalar PDEs22587 Rotated Laplacian drives the operation count up sharply. More compute per byte to begin with, so less to win back from cache — but it is the industrial workhorse for RTM and FWI.
isotropic elasticcoupled vectorial + tensorial19436 Nine state parameters instead of one or two, so data movement dominates. First order in time, which shortens the temporal dependence — a different pattern along t on purpose, and evidence the scheme is not tuned to one shape.

512³ grid, 512 ms simulated, single precision. Grid spacing 10 m; 20 m for TTI.

Throughput speedup over spatially blocked code GPoints/s, after auto-tuning both schemes 1.0× 1.2× 1.4× 1.6× baseline — spatially blocked, vectorised, tuned space order 4 every model gains, on both machines 1.6× 1.44× 1.22–1.30× space order 8 the common production choice ≥1.13× ≥1.13× Skylake ≥1.13× Broadwell space order 12 Broadwell only · the technique runs out of room none ~1.05× ~1.05× isotropic acoustic anisotropic acoustic (TTI) isotropic elastic ↑ caps mark a reported floor, not a measured peak colour here is kernel identity, not the site-wide space/time/sparse key

A wider stencil needs a steeper wavefront angle, so each tile sweeps across more columns than it updates and its working set grows for no extra reuse — while the stencil itself demands more space updates per point advanced in time. Both terms move the wrong way at once, and neither is something tuning can recover. The paper is direct about this: gains of 13–60% at space orders 4 and 8, 5–10% at order 12, and it flags high-order kernels as open work, pointing at stencil retiming and data-layout transformations as possible routes. Fig. 10 puts numbers on the mechanism.

Cache-aware roofline · schematic after the paper’s own Fig. 11 — Broadwell, isotropic acoustic, measured with Intel Advisor. Redrawn for shape, not for coordinates. arithmetic intensity · FLOP/Byte · log scale → performance · GFLOP/s · log scale → vector FMA peak vector add peak scalar add peak L1 L2 L3 DRAM space order 4 — crosses the L3 roof order 12 — the pair barely separate; there is little left for the transformation to move space order 4 order 8 order 12 spatially blocked, vectorised temporally blocked

A cache-aware roofline plots one point per kernel against several bandwidth roofs — one for each level of the hierarchy — rather than a single DRAM roof, so a kernel’s position tells you which level is holding it back. Marker shape is space order and colour is scheme, following the paper’s own coding: red for the spatially blocked vectorised kernels, yellow for the temporally blocked ones.

Reusing a tile across timesteps cuts the bytes moved per flop, so each yellow marker sits up and to the right of its red partner. The vertical gap between the pair is the speedup; the horizontal gap is the traffic that reuse removes — wide at space order 4, where the kernel clears the L3 roof, and barely visible by order 12. The measured version was produced with Intel Advisor on Broadwell; the geometry here is redrawn to show that structure, with no measured coordinates.

Tuned tile and block shapes · Table I
ProblemBroadwellSkylake
Acoustic O(2,4)32, 32, 8, 864, 64, 8, 8
Acoustic O(2,8)64, 64, 8, 864, 64, 8, 8
Acoustic O(2,12)256, 256, 8, 8128, 128, 8, 8
Elastic O(1,4)32, 32, 8, 832, 32, 8, 8
Elastic O(1,8)32, 32, 8, 864, 56, 8, 12
Elastic O(1,12)256, 256, 8, 8256, 256, 8, 8
TTI O(2,4)40, 32, 4, 448, 48, 8, 8
TTI O(2,8)32, 32, 8, 864, 64, 8, 8
TTI O(2,12)256, 256, 8, 8256, 256, 8, 8

tile_x, tile_y, block_x, block_y. Note the jump to 256 at order 12 — the tiles grow until they stop being tiles.

The setup
machines
Azure E16s_v3 — 8-core Broadwell E5-2673 v4, AVX2, 50 MB L3. E32s_v3 — 16-core Skylake 8171M, AVX-512, 35.75 MB L3.
toolchain
Devito v4.2.3, GCC 7.5.0 and ICC 2021.1, OpenMP with dynamic scheduling, thread pinning on.
tuning
Tile and block shapes swept through a block_sizes Function with Devito's advanced mode (acoustic_tune_model.py, tag v0.9-alpha): the shapes are runtime arguments, so one compiled kernel serves the whole sweep. The spatially blocked baselines additionally use configuration['autotuning'] = 'aggressive'.
models
512³ points, single precision, 512 ms simulated. 228 timesteps acoustic, 436 elastic, 587 TTI. Grid spacing 10 m, 20 m for TTI. Zero initial conditions, damping absorbing layers.
more sources
Adding sparsely located sources barely moves the gain. Only when sources are packed densely through the whole volume does the scheme lose its structural sparsity — and even then it holds about 1.4× against 1.55×.
code
Devito fork, MIT licensed. Release v0.9-alpha (Feb 2021) holds the kernels that ran the paper's experiments — demo_temporal_sources.py, demo_tb_acoustic.py and the TTI/elastic wavesolver.py variants. Branch v2.0-ipdps-accepted carries the later time_m/time_M work, demonstrated in examples/seismic/tutorials/12_time_blocking.ipynb.

Up to 1.6× at space order 4 and around 1.13× or better at order 8, varying by propagator, on kernels that were already flop-optimised and tuned and that previously could not be time-tiled at all. The metric is GPoints/s, time to solution, rather than GFLOP/s: this technique adds no redundant computation (Fig. 08), so GFLOP/s moves with the same speedup shown here — but flop count isn't fixed across every tiling family (overlapped tiling trades redundant flops for concurrency), so points/s stays the metric of record, and the roofline is read as a diagnostic rather than a scoreboard. Temporal blocking is often reported at larger multiples, but against baselines tuned to a lesser degree and without the sparse-operator constraint that had to be lifted here; tuned against tuned is the conservative reading.

FIG. 10

Why the gains fall away at high space order

Fig. 09 shows the speedup collapsing between space order 4 and 12. The arithmetic behind it also accounts for the tile sizes in Table I. Two constraints pull the tile width in opposite directions. Skewing wants the tile wide, because a leaning tile touches W + r·T columns while only updating W of them. Cache wants the tile narrow, because z is never blocked, so a tile holds W² · nz elements for every field and every time buffer at once. The stencil radius sets how sharp the conflict is.

4
Tile width, squeezed from both sides left axis · skew efficiency η = W / (W + r·T) · right axis · working set per tile, log scale tile width W (tile_x = tile_y) skew efficiency working set
Assumptions

Working set counted as W² · nz · 4 bytes · planes with nz = 512, single precision, and planes = fields × time buffers: 3 for acoustic (one field, second order in time), 6 for TTI (two coupled fields), 18 for elastic (nine state parameters, first order in time). It also ignores the lean itself: a tile skewed by r over T timesteps spans (W + r·T)² columns rather than , so the true footprint is larger still at high order. It ignores the coefficient and parameter fields — m, damp, θ, φ, the Lamé fields — and ignores halo overlap between neighbouring tiles, so it is a floor, not an estimate. Efficiency η counts columns touched against columns updated and says nothing about pipeline fill and drain at the domain edges. This is a model for reading Table I, not a measurement from the paper. One disagreement remains: it calls elastic’s order-8 window shut, while Table I still lists a 32×​32 tile there with a measured gain — assuming all nine state parameters stay resident at once most likely over-counts elastic. Two corrections the literature suggests: the aspect-ratio argument — that a tile's footprint is minimised by a specific ratio of its space extent to its time extent, not by maximising either alone — is the classical one for skewed tiles (Wonnacott; Strzodka's cache-accurate time skewing bounds the same footprint with single-timestep passes), and partially blocking z — what 3.5D schemes do — would shrink the term at the cost of the contiguity SIMD needs.

The cyan curve is skew efficiency: it climbs with tile width because the r·T columns of lean amortise over a wider tile. The violet curve is the working set: it grows as and crosses the L3 ceiling fast, because the z dimension is left unblocked on purpose, for vectorisation. Between them sits a window of workable tile widths.

At space order 4 the window is comfortable. At order 12 the radius triples, efficiency needs a tile far wider than cache will hold, and the window shuts — which is the reading offered here of why Table I's order-12 rows jump to 256×256: the autotuner has effectively backed away from temporal blocking rather than found a better tile. Note also that elastic, with nine state parameters, hits the ceiling at a much narrower tile than acoustic does, which is consistent with its tiles staying at 32×32 while acoustic moves to 64×64.

NOTES

Open threads

What this page does not claim, and where the work goes next.

Care with the claims
the tiling is hand-applied
Devito generates the precomputation, the mask and identifier arrays, and the fused loop nest. The wavefront tiling on top was applied by transforming the generated C by hand. Full automation in the DSL is future work and should be described that way.
one source in the main results
The headline speedups use a single source. The corner-case study extends to many, and holds at roughly 1.4× even when sources are dense — but the main numbers should not be read as a multi-shot result.
both sides autotuned
The comparison sweeps the tile and block parameter space for the temporally blocked code and compares against Devito's own aggressively tuned spatially blocked kernels. That is the fair comparison, and it should be said up front: it is the first thing a reviewer will ask.
the metric is GPoints/s
Time to solution, not GFLOP/s. This technique adds no redundant computation, so GFLOP/s moves with the speedup here — but that isn't true of every tiling family, so points/s is the safer metric of record; the roofline is a diagnostic, not a scoreboard.
order 12 is a negative result
Reported as such. Fig. 10 argues it is structural rather than an artefact of tuning — an argument worth checking before leaning on it.
Where the work goes next
MPI-aware scheme
Tiles that cross rank boundaries need the same treatment the seam gets here, with halo exchange scheduled against the time-tile rather than the timestep.
GPUs and other architectures
ARM, and accelerators where the cache hierarchy and the vectorisation constraint on z both look different.
high space order
Stencil retiming and data-layout transformation are the two candidates named in the paper. Both act at register level — retiming reorders the intermediate computations so fewer temporaries stay live, vector folding re-packs the layout so each SIMD lane carries less halo. Fig. 10 suggests where to aim: the binding constraint is the working set, so anything that shrinks resident planes buys back window.
cache-oblivious schedules
Shapes here are autotuned per machine — Table I differs between Broadwell and Skylake. Recursive, cache-oblivious space–time schemes (Frigo & Strumpen) remove that tuning at some cost in peak; the tiling argument does not depend on which is used.
full DSL integration
The end state is a user writing the same five lines as the baseline and getting time-tiled code out. Since the paper, Devito gained time_m/time_M sub-timestep execution and save=Buffer(M) — present in the v2.0-ipdps-accepted fork and demonstrated in examples/seismic/tutorials/12_time_blocking.ipynb — so time blocks are now expressible as data, but the skewed wavefront schedule itself is still not generated.
moving sources
The precomputation assumes fixed source coordinates across the time domain. Devito's API supports moving sources and the algorithm does not depend on the assumption, but the footprint would need recomputing.
receivers
Interpolation is the adjoint of injection — a gather rather than a scatter — and the same alignment applies. In FWI and RTM both appear in the same operator, so a future draft should show the receiver path explicitly.