IPDPS 2021 · arXiv:2010.10248

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

These solvers are only fast when cached data is reused 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

How this was made

These figures were built by Claude Opus 5 (Anthropic) at the authors’ request, 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.

Worth checking before any of it is reused in a talk or a 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.

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 is worth doing 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 loop 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
  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 is loaded once for all t in tt
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.

That interchange is not legal as written: a block's edge needs neighbours from the block beside it, which has not reached the timestep yet. Each family above restores legality differently — by leaning, by recomputing, or by choosing a tile shape that never asks. This paper uses skewed wavefront tiling, which Fig. 03 covers. The problem this paper addresses is orthogonal to that choice: none of these families can be applied at all once the simulation contains sources that are not on the grid.

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

Two figures step outside this: in the tiling families hue only tells tiles apart, and the roofline follows the paper’s own red-and-yellow coding.

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 · 3-point update u[t+1][x] = a·u[t][x-1] + b·u[t][x] + c·u[t][x+1] t+1 t 3D · star stencil radius r = space order / 2 · one arm per axis direction
the sweep runs left to right; the star turns and grows with the space order

Left, the canonical case: three neighbours at t feed one point at t+1, and the update marches along the row. Every blocking scheme below has to respect those three arrows. Right, the same idea in 3D — a star reaching r = so/2 points along each axis, so the point count is 6r + 1. Watch the arms as the order cycles: 13 points at space order 4, 25 at order 8, 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 FWI and RTM 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.

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 light cone, and the cone 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.

The cost is the flip side of the same geometry, and it is worth being exact about what it is. 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 data it must hold 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. This is why the paper's 1.6× at space order 4 falls to about 1.13× at order 8 and vanishes by order 12.

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 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 · two tiles 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 unchanged, 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 (n_src, nt) 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 (n_aff_pts, nt) 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 = sid[x][y][k] 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(n_src, nt)(n_aff_pts, nt)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 same operation count, and a single iteration space that Fig. 03's skew can now be applied to. 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.

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 is automated, and what is not

Only the first half of this is automatic. The precomputation, the mask and identifier arrays, and the fused loop nest are all generated by Devito from the code in Fig. 07. The wavefront tiling on top — the outer tile loops and the x - time skew — was applied by manually transforming the generated C. Full automation inside the DSL is listed as future work, and the paper is explicit about it.

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.

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 Laplacian, low arithmetic intensity, 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 deliberately different pattern along t, 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 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

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 15–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 distance between the pair is the speedup: 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.
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 georgebisbas/devito v0.9-alpha, MIT licensed, with reproduction instructions.

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: a temporally blocked kernel can post a lower flop rate and still finish sooner, which is why the roofline is read as a diagnostic rather than a scoreboard.

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 our model for reading Table I, not a measurement from the paper. One disagreement worth chasing: it calls elastic’s order-8 window shut, while Table I still lists a 32×​32 tile there with a measured gain — we are almost certainly over-counting elastic by assuming all nine state parameters stay resident at once.

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 deliberately left unblocked 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 our reading 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 artifact deliberately 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. This is the fair framing and is worth stating loudly, because it is the first thing a reviewer will ask.
the metric is GPoints/s
Time to solution, not GFLOP/s. A temporally blocked kernel can post a lower flop rate and still finish sooner; 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 — worth checking that argument before we lean 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. Fig. 10 suggests where to aim: the binding constraint is the working set, so anything that shrinks resident planes buys back window.
full DSL integration
The end state is a user writing the same five lines as the baseline and getting time-tiled code out.
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 it is worth showing the receiver path explicitly in a future draft.