ドクター・ストレンジは時の流れを渡り歩き、1400万通りの可能性の中からたった一つの未来を探した。彼には、パソコンの前でキーを叩く猿が見えていただろうか。
最初のトランザクションは、新しいプールにファイルを一つ作り、COWで一回書き込むだけ。クラッシュ状態は262,165個に展開され、CPUのシングルスレッドで22.7秒。問題は二つ目のトランザクションから始まった。最初の固定スクリプトはロールバック後のE地点までしか走らず、状態数は210万個。それでもシングルプロセスで2本のストリーム合計3時間50分かかった。マルチスレッドに切り替えたあと、スクリプトもフルになった。上書き、解放、ロールバック、再利用がすべて入り、大きなセグメントでは一度に24回、28回と書き込む。セグメント内で書き込みが1回増えるごとに、状態数は倍になる。二つ目のトランザクションのクラッシュ展開は10億オーダーと見積もっていたが、実際に数えるとE地点までのストリーム1本だけで16.6億個、マイルストーン2全体では26億個になった。この桁になると、CPUだけではもう現実的ではない。
だから目の前の道は二つ。一つ、サンプリング。二つ、別の方法を考える。
しかし、singlefsの理念ではコードの安定性の証明が最優先だ。サンプリングはシードを固定しても、安定して再現できるのは引いた部分だけ。引かなかった状態に誤りがあるかどうかには答えられない。全数探索ならクラッシュ状態を一つ残らず判定する。0レッドと出れば、それは本当に0レッドで、見逃しはない。だからサンプリングは考えない。
では、どう解決するか。
クラッシュ展開の流れをぐっと抽象化すると、実は二つのステップしかない。一つ目、CPUが各ステップで何を書いたか、各クラッシュ状態の事実を計算する。二つ目、その事実を突き合わせて、正しいかどうかを判定する。もともとこの二つは混ざっていた。一つ計算して、一つ判定する。
CPUはずっと同じ仕事をしているときに効率が高い。なぜか。CPUは命令が来たら一つずつ処理しているわけではなく、先読みしている。よく使うデータはキャッシュに置いてあって、手を伸ばせば届く。メモリまで取りに行かなくていい。ifに当たれば、分岐予測がどっちに行くかを先に当てにいく。当たればそのまま進み、外れればそれまでの仕事は全部やり直し。同じ仕事をずっとしていれば、キャッシュには同じデータが入り続け、分岐はいつも同じ道を通る。走るほど滑らかになり、ほぼ毎回当たる。
ところが、記録と判定はまったく別の仕事だ。記録はディスクへの書き込みをシミュレートすることで、一歩ずつ進めながら、このステップでどのブロックを書いたか、ハッシュは何かを計算する。判定はクラッシュ後の後始末で、イメージを組み立て、二回リカバリーし、checkerを頭から一通り走らせる。使うデータも、通るコードも、分岐も違う。少し記録して、少し判定して、と行ったり来たりすると、キャッシュに記録の中身を詰め終えたところで判定に移って全部入れ替え。分岐予測が記録のクセをつかんだところで、また判定のクセを当てにいく。切り替えるたびに、CPUは熟練工から新人に戻ってしまう。
料理と同じだ。野菜を一本切って一品炒め、また一本切ってまた一品炒める。人はまな板とコンロの間を行ったり来たりし、包丁と鍋を交互に持ち替える。先に野菜を全部切ってから一気に全部炒めれば、まな板はまな板、コンロはコンロ。手の仕事が途切れない。
もう一つある。記録は流れに沿って前に進むもので、次のステップは前のステップを待たなければならない。判定は違う。状態ごとに独立して判定でき、互いに関係しない。混ざっていたときは判定も記録に引きずられ、一つずつしか進めなかった。分ければ判定は完全に解き放たれる。31スレッドがそれぞれひと塊ずつ受け持ち、誰も誰かを待たなくていい。
だから二つのステップを分け、状態の計算と突き合わせを別々の工程にする。そうすれば、それぞれの工程の中で速度を上げられる。
状態と結果はデータベースに置く。このデータベースには高速な読み書きが必要なので、RocksDBを選んだ。流れは三段になる。CPUが記録してKVに書き込み、読み出して突き合わせ、結果をKVに書き戻し、最後にゲートが違反の集計を一度で読み出す。状態そのものは保存しない。各状態のマスクは番号からその場で計算する。D地点までの5177万状態で、DBはたった26MBだ。
分けたことで構造ははっきりした。だが純CPUではまだ遅い。2台のマシンで各31スレッド、D地点までの5177万状態に1時間45分。合計で毎秒8000ちょっと。この速度だと26億には3日半かかる。しかもD地点までは小さなセグメントがほとんどで、後ろに行くほどセグメントは大きくなり、1状態あたりの仕事も重くなる。3日半は下限にすぎない。もっと速くできないか。
できるよ、兄弟。ちょうど手元にグラボが何枚かある。
前半の記録とセグメント分割はCPUでしかできない。GPUには手が出せない。でも後半の突き合わせなら、GPUに任せられるはずだ。ここでいう突き合わせは、二つの値を比べるだけの話ではない。各状態についてまずクラッシュ後のイメージを組み立て、journalを見る場合と見ない場合でそれぞれリカバリーし、さらにcheckerを一通り走らせる。だが状態同士は互いに依存せず、判定方法はまったく同じで、入力が違うだけ。これこそGPUがいちばん得意な仕事だ。
CPUとGPUはまったく別物だ。CPUは数十人のベテラン職人。一人ひとりが賢くて、分岐を読み、先回りして計算し、どんなに複雑な仕事もこなす。でも人数は数十人しかいない。GPUは何千、何万人の見習い。5090一枚で2万個以上の小さなコアがある。一人ひとりは不器用で、先読みも先回りもできず、同じ一枚の指示書どおりに動くだけ。しかも一人ずつバラバラに働くのではなく、32人で一列を組み、同じ瞬間に列全体が同じ命令を実行する。
これがインクジェットプリンターだ。ヘッドにはノズルが一列に並んでいて、プリンターが出す命令は「吹け」の一つだけ。ノズルごとに位置も色も違うのに、命令は一つで、一列が同時に吹きつける。ノズルが多いほど、一回のスキャンで刷れる幅が広くなる。
僕らの判定はまさにこれだ。判定器のコードは一つだけ。それが指示書。GPUの各スレッドが状態番号を一つ受け取る。それが各ノズルの狙う位置。一つの命令が下りると、何万もの状態が同時に一歩進む。同時にイメージを組み、同時にリカバリーし、同時にcheckerを走らせる。状態同士は相談も待ち合わせもいらない。それぞれ判定して、結論を自分のマスに書き戻すだけ。
もちろん、インクジェットにはインクジェットのクセがある。一列のノズルは一緒に進み、一緒に止まる。列の中に一つだけ道のりが特別に長い状態があれば、残りの31個は判定を終えていても、付き合って空回りするしかない。
そこで判定器をWGSLでGPU向けに書き直し、wgpu経由でVulkan上で動かした。リカバリー、ウォーク、スキャンの三つのカーネルで九千行ほど。プールレベルのcheckerの49個の不変条件、レコード照合器、7種類のoracleをカバーする。GPUの各スレッドがクラッシュ状態を一つ受け持つ。エントリーのコードは次のとおり（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));
}
/ …… レコード照合器の二つの判定基準と、判定できない状態のフラグ立て、省略 ……
let out = local * 12u;
verdicts[out] = bits;
// …… 着地したルートと二回のリカバリーの結末を書き戻す、省略 ……
}
begin_state は状態番号を「どの書き込みがディスクに落ちたか」に戻す（crash_judge_common.wgsl）。
// 番号 → 永続マスク（crash_facts::persisted_writes_of_state と同じ桁の順序）
// …… まず番号がどのセグメントに入るかを探し、それより前のセグメントはすべてディスクに落ちたものとする、省略 ……
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));
}
}
判定器を書き直したなら、元のCPU判定器と同じ判定をすることを証明しなければならない。17本のストリーム、52万個あまりの状態で一状態ずつ突き合わせ、相違は0。
こう聞かれるかもしれない。52万状態は26億の1万分の2。それもサンプリングじゃないのか。違う話だ。全数探索の対象はファイルシステムのクラッシュ状態で、26億個を一つ残らずGPUで判定した。突き合わせで検証しているのは判定器という道具そのもの。だからこの52万状態は選んである。17本のうち11本はわざと壊してあり、必ずレッドにならなければならない。どの状態でレッドになったか、どの条件でレッドになったか、両者で完全に一致しなければならない。さらに、毎回のフル実行には別に書いたGPU照合カーネルが一つ入っていて、ルートの自己証明、ポインタの着地点、ユニットのチェックサムを判定し直し、判定器の結論と一状態ずつ比べる。E地点までの16.6億状態はすべて比較済みで、不一致は0。
ここは正直に言っておく。GPU判定器とCPU判定器がすべての判定基準で一状態ずつ一致することは、今のところこの52万状態でしか証明していない。フル実行では各状態に照合カーネルが付いているが、カバーしているのは判定基準の一部だけだ。開発の途中で「CPUが16ブロックに1ブロックを抜き取って再判定する」仕組みも入れたが、あとでまるごと消した。場所を変えても、サンプリングはサンプリングだ。
試してみると、初版はよくなかった。一状態ごとに一回ジョブを投げ、バッファを作り、submitし、結果を待つ。一状態ごとに往復していたので、CPUより18倍遅くなった。プリンターが一点吹くたびに、紙を送り、予熱し、紙を排出するようなもので、ノズルがいくら多くても時間は全部紙送りに消える。まとめて一度に投げるように変え、何千何万の状態を一緒にGPUに載せると、効果は抜群だった。D地点までの5177万状態で、全体91秒。GPUは毎秒60万以上の状態を判定した。これには本当に驚いた。最終結果は26億、正確には2,563,727,926状態。2台のマシンと5枚のカードで7時間35分、0レッド。
毎秒60万以上なら、26億は1時間ちょっとで済むはず。なぜ7時間半もかかったのか。時間は三か所に消えている。
一つ目、E地点までの項目が4時間31分を占めた。その16億状態はほぼすべて28回書き込みの大きなセグメントの中にあり、GPUは毎秒4万から7万しか判定できない（このマシンの2枚で72,429、2台目の3枚で44,471）。原因はまだ分解して測っていない。大きなセグメントでは1状態あたりにたどるユニットがずっと多いからだと推測している。しかも2台の分担が偏っていて、2台目は1時間半早く判定を終え、そのまま待ちぼうけだった。
二つ目、ストリームを1本判定し終えるたびに、シングルスレッドでストリームを記録し直し、2台目のDBに取り込み、突き合わせる。1本あたり5〜9分。残り6項目16本のストリームで、合計2時間10分。
三つ目、残りの53分が、ほかの6項目を実際にGPUで判定していた時間と、各項目の起動と停止だ。
つまりボトルネックは「判定」から「記録」と「大きなセグメント」に移った。それでも純CPUの3日半という下限より11倍以上速い。
では、もっと速くできないか。
できるよ、兄弟。できる。
このデータを構造化して保存し、各クラッシュ点を決まった規則で符号化する。名前は「どのエントリーか、このセグメントで何を書いたか、何番目のセグメントか」。アイデンティティは二つのもののハッシュ。このステップで書いたバイト列（バリア込み）と、このステップ開始時のイメージだ。COWの理屈では、入力が変わらなければ出力も変わらない。符号化すると、クラッシュ点は流れの前後でつながって一本のプレフィックス木になる。セグメント内の各書き込みが落ちたか落ちなかったかが二択（インプレース上書きにはさらに「半分ちぎれた」がある）。2本のストリームの前半が同じなら、そこは一度しか判定しない。ノードの入力が変わらない限り、その値はもう走らせなくていい。そのまま再利用できる。
判定方法が変わってもごまかさない。判定方法のバージョンは各結論に記録してあり、判定方法が変われば全部判定し直す。どうせGPUは安い。上の26億のフル実行も、判定器が一つの着地点の組み合わせマスクを32ビットから64ビットに広げたあと、全部判定し直した結果だ。
2回目に走らせるときは、速度がぶっ飛ぶはずだ。
最終的なコードは次のとおり（crates/singlefs-checker-tier/src/crash_identity.rs）。
/// 1ステップの二つの属性（それぞれ SHA-256）：このステップが書いたバイト列（バリア付きの書き込み表）と、開始時のイメージ。
/// 再利用キーに入るのは流れだけ。判定方法は判定行の属性（`verdict_store::JudgeVersion`）で、ここには入らない。
#[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 {
/// 二つを固定の順序で、それぞれ名前を付けて連結し、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))
}
}
2回目は案の定ぶっ飛んだ。同じDBでD地点までをもう一度走らせると、5177万状態は一つも判定し直さず、判定ステップは2.2秒。2台の間の突き合わせを含めても項目全体で5分ほど。位置アドレス指定の5本のストリーム、5.5億状態は、1本あたり1〜3秒で通り過ぎた。
前に書いた記事には、暴力開発と書いた。
今回は、暴力テストだ。
つくづく思う。ここまで来ると、正しさの検証はもうリソースのゲームになった。
計算力とグラボの檻の中で、猿の群れがハムレットを一篇また一篇と書いている。
準備はいいか。