# Which Cell Am I?

*Task 3 of 6 · [Wavefronts: Aligning DNA on the Diagonal](https://gpu.rocks/learn/sequence-alignment-a85ca6d9.md) · GPU.js Learn*

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:

```js
// 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.

## Figures

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

## Goal

**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.

## Starter code

```js
// One thread per cell of the diagonal — no passengers.
const gpu = new GPU({ mode });

const diagonal = gpu.createKernel(function (H, a, b, d, iStart) {
  // TODO: which cell does this thread own?
  //   i walks DOWN the diagonal from iStart, one row per thread
  //   j is whatever makes i + j come to d
  const i = iStart;
  const j = d - iStart;
  let s = this.constants.mismatch;
  if (a[i - 1] === b[j - 1]) s = this.constants.match;
  const diag = H[i - 1][j - 1] + s;
  const up = H[i - 1][j] - this.constants.gap;
  const left = H[i][j - 1] - this.constants.gap;
  return Math.max(0, Math.max(diag, Math.max(up, left)));
}, { dynamicOutput: true, constants: { match: 3, mismatch: -3, gap: 2 } });

// Anti-diagonal d covers rows i = iStart … iStart + len - 1, and j = d - i.
// Clipped at both ends: i never passes 8 (= |A|), and j never passes 9 (= |B|).
function diagonalRange(d) {
  const iStart = Math.max(1, d - 9);
  const iEnd = Math.min(8, d - 1);
  return { iStart, len: iEnd - iStart + 1 };
}

// `partial` is the matrix mid-sweep: diagonals 2…9 are final, the rest is 0.
const range = diagonalRange(10);
diagonal.setOutput([range.len]);
const values = await diagonal(partial, codesA, codesB, 10, range.iStart);

console.log('diagonal 10 starts at row', range.iStart, 'and has', range.len, 'cells');
console.log(Array.from(values));
```

---

Interactive version: https://gpu.rocks/learn/sequence-alignment-a85ca6d9/3

[Previous task](https://gpu.rocks/learn/sequence-alignment-a85ca6d9/2.md) · [Next task](https://gpu.rocks/learn/sequence-alignment-a85ca6d9/4.md)