Task 3 of 6

Which Cell Am I?

The last task launched 90 threads to compute at most 8 cells. Every thread not on the diagonal woke up, read a value and wrote it straight back out — 1,440 thread-launches across the whole sweep to fill 72 cells. Correct, and embarrassing.

The fix is to launch the diagonal itself: output: [len], one thread per cell that actually exists. The price is that a thread no longer arrives knowing its coordinates. It gets this.thread.x — a number from 0 to len − 1 — and has to work out which matrix cell that is. Two lines, and they are the fiddly heart of every wavefront implementation ever written:

// walk DOWN the diagonal, one row per thread
const i = iStart + this.thread.x;
// and j is whatever makes i + j come to d
const j = d - i;

iStart is the row where the diagonal enters the matrix, and it moves. Early on it is row 1 — the top edge — and the diagonal runs out at the left edge, where j would fall below 1. Later the top edge stops mattering and the right edge takes over: j can never exceed |B|, so the diagonal cannot start any higher than row d − |B|. Hence iStart = max(1, d − |B|) and iEnd = min(|A|, d − 1) — handed to you below, because getting them wrong is a JavaScript bug, while getting the two lines above wrong is a GPU bug that reads whatever happens to be next to the matrix.

a thread arrives knowing only its number — two lines make it a cell

Goal: write the compact diagonal kernel — len threads, each computing exactly one cell of anti-diagonal d — and pull diagonal 10 out of a half-filled matrix with it.

Requirements

Map the thread index to a row: i = iStart + this.thread.x
Map the row to a column: j = d - i
Run the same recurrence as before, and return one number per thread
Create the kernel with dynamicOutput: true so setOutput([len]) can resize it per diagonal

Hint 1 — no boundary guard is needed here

iStart and len already keep i inside 1 … |A| and j inside 1 … |B|. That is what they are for: the launch shape does the clipping, so the kernel body is nothing but the recurrence.

Hint 2 — which end is thread 0?

Thread 0 owns the cell nearest the top of the matrix — the smallest row index on the diagonal, which is iStart. Thread len − 1 owns the one furthest down. Get that backwards and every value is right but in the wrong order.

Hint 3 — resizing between launches

dynamicOutput: true is what makes a kernel resizable; diagonal.setOutput([len]) then sets the thread count for the next call, the same move the halving ladder in Reductions makes on its way down.

Same idea elsewhere

Turning a thread id into a coordinate is most of what a GPU programmer does all day: CUDA's blockIdx * blockDim + threadIdx, WGSL's global_invocation_id and Metal's thread_position_in_grid all hand you a flat number and leave the geometry to you. Launching exactly the work that exists — rather than a rectangle with a guard in it — is the same instinct behind compacted launches, persistent-thread kernels and indirect dispatch.

All tasks in Wavefronts: Aligning DNA on the Diagonal

This page is an interactive exercise — the editor, the GPU runner and your saved progress need JavaScript. The text above is the full brief.