Doctor Strange drifted through the river of time, searching 14 million possibilities for the one future that works. Did he see the monkey typing at a computer?
The first transaction: create one file in a fresh pool, one COW write. The crash states blew up to 262,165, and a single CPU thread chewed through them in 22.7 seconds. The second transaction is where the trouble started. At first the fixed script only ran up to point E, right after the rollback: 2.1 million states, and a single process took 3 hours 50 minutes for the two streams combined. After we went multithreaded, the script got filled in too: overwrite, free, rollback and reuse were all wired in, and the big segments now issue 24 or 28 writes in one go. Every extra write inside a segment doubles the number of states. We predicted the second transaction would blow up into the billions. When we actually counted, the stream up to E alone was 1.66 billion states, and all of Milestone 2 together came to 2.6 billion. At that order of magnitude, CPU alone stops being realistic.
So there were two roads in front of us. One: sampling. Two: find another way.
But under the singlefs philosophy, proving the code is stable comes first. Sampling, even with a fixed seed, can only reliably reproduce the part it happened to draw. Whether there's a bug in the states it didn't draw, it can't tell you. Exhaustive search judges every single crash state. If it comes out zero red, it really is zero red; nothing slips through. So sampling was off the table.
So how do you solve this?
Abstract the crash-blowup pipeline far enough and there are really just two steps. Step one: the CPU works out what each step wrote and the facts of every crash state. Step two: check those facts and judge whether they're right. Originally the two were mashed together: compute a state, judge a state.
We know a CPU is very efficient when it keeps doing the same kind of work. Why? A CPU doesn't just take one instruction and compute one instruction; it guesses ahead. The data it uses often sits in cache, within arm's reach, no trip to main memory. When it hits an if, the branch predictor guesses which way it'll go first. Guess right and it keeps running; guess wrong and everything it did ahead gets torn down and redone. Keep doing the same work, and the cache always holds the same set of data, the branches always take the same few paths. The longer it runs, the smoother it gets, and it guesses right almost every time.
But recording and judging are two completely different jobs. Recording simulates writing to disk: push forward one step at a time, work out which blocks this step wrote and what the hashes are. Judging is cleaning up after the crash: assemble the image, recover twice, then walk the checker from start to finish. Different data, different code, different branches. Record a bit, judge a bit, back and forth: you've just filled the cache with recording stuff and now you turn around to judge, so it all gets swapped out. The branch predictor has just figured out recording's habits, and now it has to start guessing judging's habits all over again. Every switch turns the CPU from an old hand back into a beginner.
It's like cooking. Chop one vegetable, stir-fry one dish, chop another, fry another. You keep running between the cutting board and the stove, picking up the knife, putting it down, picking up the wok. Chop everything first, then fry it all in one go, and the cutting board is the cutting board, the stove is the stove, and your hands stay in rhythm.
There's one more layer. Recording moves forward along the flow; each step has to wait for the one before it. Judging is different: every state is judged on its own, independent of the rest. When the two were mashed together, judging got dragged along by recording and could only go one at a time. Split them apart and judging is completely freed up: 31 threads each take a batch, and nobody waits for anybody.
So we split the two steps, making state computation and checking two separate stages, so each stage can run faster on its own.
States and results live in a database, and that database has to read and write fast, so we picked RocksDB. The pipeline became three stages: the CPU records and writes into the KV store; we read it back, check it, and write results back into the KV store; and finally the gate reads out the violation summary in one pass. The states themselves aren't stored; each state's mask is computed on the fly from its index. For the 51.77 million states up to D, the database is only 26 MB.
Splitting made the structure clean, but pure CPU still wasn't fast enough: two machines with 31 threads each took an hour and forty-five minutes to get through the 51.77 million states up to D, a combined eight thousand or so per second. At that rate, 2.6 billion would take three and a half days. And the run up to D is mostly small segments; the further you go, the bigger the segments and the heavier each state's work. Three and a half days is only the floor. So can we go faster?
Yes we can, bro. I happen to have a few graphics cards lying around.
The recording and segmenting in the first half can only happen on the CPU; the GPU can't help there. But the checking in the second half should be something we can hand to the GPU. This checking isn't just comparing two values: for each state you first assemble the post-crash image, recover it once consulting the journal and once ignoring it, then run the checker over it. But states don't depend on each other, the judging is exactly the same, only the input differs. That happens to be exactly the kind of work a GPU loves most.
A CPU and a GPU are completely different beasts. A CPU is a few dozen master craftsmen. Each one is smart: predicts branches, races ahead, can handle any complicated job. But there are only a few dozen of them. A GPU is thousands upon thousands of apprentices; a single 5090 has over twenty thousand small cores. Each one is dumb: can't guess, can't race ahead, can only work from the same instruction sheet. And they don't work one by one, either. They're lined up in rows of thirty-two, and at any given moment the whole row executes the same instruction.
That's an inkjet printer. The print head has a row of nozzles, and the printer sends just one command: spray. Every nozzle sprays a different spot, a different color, but there's only one command, and the whole row lands at once. The more nozzles, the wider each pass prints.
Our judging is exactly like that. There's only one copy of the judge's code; that's the instruction sheet. Each GPU thread takes one state index; that's the spot each nozzle is aimed at. One instruction goes out, and tens of thousands of states take a step forward together: assemble images together, recover together, run the checker together. States don't need to talk to each other or wait for each other. Each judges its own, and writes its verdict back into its own slot.
Of course, an inkjet has its quirks: the row of nozzles moves in lockstep. If one state in the row has an especially long road, the other 31 can finish and still have to sit there waiting for it.
So we rewrote the judge for the GPU in WGSL, running on Vulkan through wgpu. Three kernels, recovery, walk and scan, around nine thousand lines, covering the pool-level checker's 49 invariants, the record checker, and seven classes of oracle. Each GPU thread takes one crash state. Here's the entry point (crates/singlefs-checker-tier/src/crash_judge_recovery.wgsl):
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
let local = id.x;
if (local >= dispatch.state_count) {
return;
}
begin_state(local);
// 判
let cfg_choice = choose_system_configuration();
let consulted = recover(true, cfg_choice);
let ignored = recover(false, cfg_choice);
let newest = newest_persisted_root();
let consulted_oracle = classify_oracle(consulted, newest);
let ignored_oracle = classify_oracle(ignored, newest);
let records = check_records(consulted);
var bits = 0u;
if (consulted_oracle != ORACLE_NONE) {
bits = bits | (1u << (BIT_CONSULTED_ORACLE + consulted_oracle));
}
if (ignored_oracle != ORACLE_NONE) {
bits = bits | (1u << (BIT_IGNORED_ORACLE + ignored_oracle));
}
// …… the record checker's two criteria and the "cannot judge" flag, omitted ……
let out = local * 12u;
verdicts[out] = bits;
// …… write back the root it landed on and the outcomes of both recoveries, omitted ……
}
begin_state turns a state index back into "which writes hit the disk" (crash_judge_common.wgsl):
// index → persisted mask (same digit order as crash_facts::persisted_writes_of_state)
// …… first find which segment the index falls in; all earlier segments count as persisted, omitted ……
for (var k = 0u; k < count; k = k + 1u) {
let write = tab(T_SEGMENT_WRITES, first + k);
let choices = tab(T_RECORDED, write * 4u);
let digit = remaining % choices;
remaining = remaining / choices;
if (digit == choices - 1u) {
set_persisted(write);
} else if (digit == 1u) {
torn[write >> 5u] = torn[write >> 5u] | (1u << (write & 31u));
}
}
Rewrite the judge, and you have to prove it judges the same as the original CPU judge. We cross-checked state by state across 17 streams and over 520,000 states: zero differences.
Someone will ask: 520,000 states is two ten-thousandths of 2.6 billion. Isn't that sampling too? It's not the same thing. The target of the exhaustive search is the filesystem's crash states, and all 2.6 billion of them, without missing one, went onto the cards and got judged. The cross-check is validating the judge itself, the tool, so those 520,000 states were hand-picked: 11 of the 17 streams are deliberately broken and must come out red, and both sides have to go red on exactly the same state, for exactly the same rule. On top of that, every full run includes a separately written GPU verification kernel that re-judges the root self-proof, pointer targets, and unit checksums, and compares against the judge's verdicts state by state: all 1.66 billion states in the stream up to E were compared, zero mismatches.
And to be honest about it: that the GPU judge and the CPU judge agree state by state on every rule has so far only been proven on those 520,000 states. In the full run, every state is also backstopped by that verification kernel, but it only covers some of the rules. Midway through development we'd added a "CPU re-judges 1 of every 16 blocks" check ourselves, and later deleted the whole thing. Sampling moved somewhere else is still sampling.
We tested it, and the first version was bad: dispatch one job per state, build buffers, submit, wait for the result, one round trip per state. It ended up 18 times slower than the CPU. It's like a printer that, for every single dot it sprays, feeds in a sheet, warms up, and spits the sheet out again. However many nozzles you have, all the time goes into feeding paper. We switched to dispatching a big batch at once, thousands or tens of thousands of states on the card together, and it worked brilliantly. For the 51.77 million states up to D, the whole item took 91 seconds, with the GPU judging over 600,000 states per second. That genuinely shocked us. The final result: 2.6 billion, to be exact 2,563,727,926 states, two machines and five cards, 7 hours 35 minutes, zero red.
At over 600,000 per second, 2.6 billion should have taken a little over an hour. Why did it take seven and a half? The time went to three places:
First, the item up to E took 4 hours 31 minutes. Its 1.6 billion states sit almost entirely in the 28-write big segments, where the GPU only manages forty to seventy thousand per second (72,429 on this machine's two cards, 44,471 on the second machine's three). We haven't broken down the cause yet; our guess is that in big segments each state has many more units to walk. And the work was split unevenly between the two machines: the second one finished an hour and a half early and just sat there idle.
Second, after each stream is judged, a single thread has to re-record the stream, import it into the second machine's database, and reconcile, five to nine minutes per stream. For the other six items' sixteen streams, that added up to 2 hours 10 minutes.
Third, the remaining 53 minutes was the other six items actually being judged on the cards, plus each item's startup and shutdown.
So the bottleneck has moved from "judging" to "recording" and "big segments." Even so, it's more than eleven times faster than the pure-CPU floor of three and a half days.
So can we go even faster?
Yes we can, bro. Yes we can.
Suppose we store this data in a structured way and encode every crash point by a fixed rule. The name says "which entry point, what this segment writes, which segment number." The identity is the hash of two things: the bytes this step wrote (barriers included) and the image this step started from. The logic of COW is: same input, same output. Once encoded, these crash points link up in flow order into a prefix tree: whether each write in a segment hit the disk or not is a binary choice (in-place overwrites add a third, "torn halfway"), and when two streams share the same opening stretch, that stretch is only judged once. So as long as a node's input doesn't change, those values don't need to be run again. They can be reused directly.
And we don't cut corners when the judging changes: the judge version is recorded on every verdict, and when the judging changes, everything gets re-judged. GPUs are cheap anyway. That 2.6-billion full run above was exactly a full re-judge, after the judge widened one landing point's combination mask from 32 bits to 64 bits.
The second run, speed is going to go through the roof.
Here's our final code (crates/singlefs-checker-tier/src/crash_identity.rs):
/// The two attributes of a step (each a SHA-256): the bytes this step wrote (the write table with barriers)
/// and the image it started from. Only the flow goes into the reuse key:
/// the judging method is an attribute of the verdict row (`verdict_store::JudgeVersion`), not of this.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub struct StepAttributes {
pub write_table_with_barriers: [u8; DIGEST_BYTES],
pub parent_output: [u8; DIGEST_BYTES],
}
impl StepAttributes {
/// Join the two in a fixed order, each tagged with its name, and take the SHA-256.
#[must_use]
pub fn reuse_key(&self) -> ReuseKey {
let mut message = Vec::new();
for (name, digest) in [
("write_table_with_barriers", &self.write_table_with_barriers),
("parent_output", &self.parent_output),
] {
message.extend_from_slice(name.as_bytes());
message.push(b'=');
message.extend_from_slice(digest);
message.push(b'\n');
}
ReuseKey(sha256_digest(&message))
}
}
And sure enough, the second run went through the roof: running the same database up to D again, not one of the 51.77 million states was re-judged. The judging step took 2.2 seconds, and even with the reconciliation between the two machines, the whole item took about five minutes. The five position-addressed streams, 550 million states, went by in one to three seconds each.
In some earlier week post, I wrote about brute-force development.
This time, brute-force testing.
I can't help but think: at this point, verifying correctness has become a game of resources.
In a cage of compute and graphics cards, a troop of monkeys is writing out Hamlet after Hamlet.
Are you ready?
Top comments (0)