WGSLで協調計算を実装する — アトミック集計とステップ反復

CONTENTS

協調計算は何が違うのか
全体構造 — 2カーネル × ステップ反復
需要計算カーネル — 各エージェントの注文
集約の核心 — アトミック加算でfloatを集める
価格更新カーネル — 1スレッドで集計結果を反映
ステップ反復 — JS側の同期ループ
なぜ同期コストが効くのか

01協調計算は何が違うのか

#1（独立並列）と#2（行列演算）では、各スレッドが自分の出力を計算して書き込めば終わりだった。スレッド同士は互いを気にしない。だが協調計算は違う。全スレッドの結果を1つに集め、それを全員に配り直し、次のステップに進む。

独立並列：各スレッドが結果を書く → 終了
協調計算：各スレッドが注文 → 全部を集約 → 配布 → 次ステップ（反復）

今回の市場モデルでは、各エージェント（スレッド）が「買いたい/売りたい量」を出し、その合計（超過需要）から価格が決まる。価格が動けば全エージェントが新価格を見て次の判断をする。この「集約して配る」を時間ステップ分くり返すのが、協調計算の実装の本質だ。難所は2つ——どうやって全スレッドの値を集約するかと、どうやってステップ間で同期するかである。

02全体構造 — 2カーネル × ステップ反復

1ステップを2つのコンピュートシェーダー（カーネル）に分けた。

1ステップ＝需要計算→集約→価格更新。これを500ステップ反復する。緑の破線が「新価格を全員が参照」というフィードバック。

カーネル①はN体ぶんのスレッドで並列に走り、各エージェントの注文を計算してアトミックに加算する。カーネル②はたった1スレッドで、集まった超過需要から新価格を計算する。この非対称さ（大量並列→1スレッド）も協調計算の特徴だ。

03需要計算カーネル — 各エージェントの注文

カーネル①では、各スレッド（=1エージェント）が自分のタイプに応じて注文量 d を計算する。

WGSL — 需要計算カーネル（要点）

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
  let i = gid.x;
  if (i >= D.N) { return; }
  let cur = LOGP[D.step];              // 現在の価格（対数）
  var prevIdx : u32 = 0u;
  if (D.step > D.k) { prevIdx = D.step - D.k; }
  let trend = cur - LOGP[prevIdx];     // 直近kステップのトレンド
  let t = TYPE[i];
  var d : f32 = 0.0;
  if (t < 0.5) {
    d = COEF[i] * (D.logF - cur);        // ファンダ：適正値へ逆張り
  } else if (t < 1.5) {
    d = COEF[i] * trend;                  // トレンド：流れに順張り
  } else {
    let g = gauss(i);                    // ノイズ：正規乱数（GPU内生成）
    d = COEF[i] * g * 0.1;
  }
  atomicAdd(&DEMAND, i32(d * D.scale));   // 超過需要に集約
}

価格履歴 LOGP は全スレッドが読む共有データだ。現在価格とkステップ前の価格からトレンドを計算し、タイプ別に注文を出す。ファンダは適正値 logF との差で逆張り、トレンドは trend で順張り、ノイズは乱数。最後の atomicAdd が集約の入口だ。

GPU内乱数ノイズトレーダーの乱数を毎ステップ生成するには、N体×500ステップ分の乱数が要る。10万体なら200MBにもなり、事前生成は非現実的。そこで各エージェントが乱数状態を持ち、GPU内でPCG法により逐次生成する（#1で確立した手法）。STATE[i] を更新しながら一様乱数を作り、Box-Muller法で正規乱数に変換している。

04集約の核心 — アトミック加算でfloatを集める

協調計算の最大の難所がこれだ。N体すべての注文 d を、1つの超過需要に合計する。何千ものスレッドが同時に同じ変数へ足し込むので、単純な DEMAND = DEMAND + d では競合（レースコンディション）が起きて値が壊れる。

解決策がアトミック操作だ。atomicAdd は「読んで足して書く」を不可分（atomic）に行い、複数スレッドが同時に呼んでも正しく合計される。だが問題がある——WGSLのアトミックは整数（i32/u32）専用で、floatには使えない。注文量 d は小数なのに。

WGSL — floatを整数化してアトミック加算

// DEMAND は atomic<i32> として宣言
@group(0) @binding(4) var<storage, read_write> DEMAND : atomic<i32>;

// 各スレッド：floatを SCALE 倍して整数化 → アトミック加算
atomicAdd(&DEMAND, i32(d * D.scale));   // scale = 1e6

注文量を ×1,000,000 して整数に丸め、その整数をアトミック加算する。集計後に ÷1,000,000 して元のスケールに戻す。固定小数点のようなテクニックだ。d が 0.000003 でも、×1e6 すれば 3 という整数になり、アトミック加算できる。これは#1で扱った「アトミックu32でカウントする」手法の応用で、カウントから加重和へ拡張したものだ。

精度とオーバーフローの注意 SCALEが大きいほど小数の精度は上がるが、合計がi32の上限（約21億）を超えるとオーバーフローする。N体×注文量×SCALEが上限内に収まるSCALEを選ぶ必要がある。今回は注文量が小さく、1e6で精度とオーバーフロー耐性のバランスが取れた。実際、CPU版と数値レベルで一致した（結果編参照）。

05価格更新カーネル — 1スレッドで集計結果を反映

集約が終われば、価格更新は単純だ。超過需要から新価格を計算する。これは1スレッドだけで行う（更新する値は1つだから並列化の必要がない）。

WGSL — 価格更新カーネル（全文に近い）

@compute @workgroup_size(1)
fn main() {
  let totalScaled = atomicLoad(&DEMAND);     // 集計結果を読む
  let total = f32(totalScaled) / D.scale;   // 元のスケールに戻す
  var dlog = D.lambda * (total / f32(D.N));  // 1人あたりの超過需要
  dlog = clamp(dlog, -0.2, 0.2);             // 1ステップの変化を制限
  LOGP[D.step + 1u] = LOGP[D.step] + dlog;     // 次ステップの価格を書く
}

atomicLoad で集計値を読み、SCALEで割って元に戻す。超過需要をN（エージェント数）で正規化し、価格感応度 lambda を掛けて価格変化 dlog を得る。clamp は1ステップの変化を±0.2に制限するもので、フィードバックループの暴走（価格が無限大やゼロに飛ぶ）を防ぐ安全弁だ。最後に次ステップの対数価格を書き込む。対数で扱うのは、価格が必ず正に保たれ、変化率が対称になるからだ。

06ステップ反復 — JS側の同期ループ

ここが#1#2と決定的に違う。これまでは「1回ディスパッチして終わり」だったが、協調計算ではJS側でステップ数だけループを回し、毎回ディスパッチする。

JavaScript — ステップ反復ループ

for (let t = 0; t < steps; t++) {
  // ① 現ステップ番号をuniformに書く
  dev.queue.writeBuffer(uParams, 8, new Uint32Array([t]));
  // ② 集約用アトミック変数を0にクリア
  dev.queue.writeBuffer(demandBuf, 0, new Int32Array([0]));
  const enc = dev.createCommandEncoder();
  // ③ カーネル①：需要計算（N体を並列）
  const p1 = enc.beginComputePass();
  p1.setPipeline(pipeD); p1.setBindGroup(0, bgD);
  p1.dispatchWorkgroups(Math.ceil(N / WG)); p1.end();
  // ④ カーネル②：価格更新（1スレッド）
  const p2 = enc.beginComputePass();
  p2.setPipeline(pipeP); p2.setBindGroup(0, bgP);
  p2.dispatchWorkgroups(1); p2.end();
  dev.queue.submit([enc.finish()]);   // このステップを実行
}

毎ステップ、現在のステップ番号をuniformに書き、アトミック変数をクリアし、2つのカーネルを順に実行する。同じコマンドエンコーダ内でカーネル①→②を並べることで、①の全スレッドが終わってから②が走る（GPU内で順序が保証される）。500ステップなら、このループが500回回る。

ここが性能を左右するこの反復ループこそが、協調計算の性能特性を決める。各ステップで writeBuffer（CPU→GPU転送）と submit（GPUへの命令発行）が入り、CPUとGPUが500回やり取りする。1回あたりのオーバーヘッドは小さくても、500回積み重なると無視できない。これが結果編で見た「GPU時間が同期回数で決まる」「デバイスによって速度が大きく違う」の正体だ。

07なぜ同期コストが効くのか

結果編のベンチマークで、協調計算は独立並列・行列演算と違う振る舞いを見せた。エージェント数を増やしてもGPU時間がほぼ一定で、しかもデバイスによって速度が桁違いに違った。実装を見れば、その理由がわかる。

計算タイプ	ディスパッチ	律速
#1 独立並列	1回	試行数（並列度）
#2 行列演算	1回	行列サイズ（並列度）
#3 協調計算	500回（ステップ数）	同期回数

#1#2は1回のディスパッチで計算が完結する。だから「どれだけ並列に処理できるか」が性能を決め、最適化（タイル化など）はGPUを問わず通用した。一方、協調計算は500回のディスパッチと同期を伴う。1ステップの計算量は小さく（エージェントが増えてもGPUは並列に吸収する）、支配的なのはCPU-GPU間を500回往復するオーバーヘッドだ。

このオーバーヘッドは、GPUアーキテクチャに強く依存する。Apple GPU（iPad）はCPUとGPUがメモリを共有する設計で、往復が軽い。一方、別アーキテクチャのGPUでは、同期のたびのコストが積み上がる。だから同じコードが、iPadで8倍速、別のデバイスでほぼ等速という結果になった。計算の中身ではなく、同期の実装とハードウェアが性能を決めるのだ。

改善の方向性同期回数を減らせれば速くなる。たとえば「数ステップ分をGPU内で一気に進める」「ワークグループ内に全エージェントを収めてworkgroupBarrierでGPU内同期する」といった手法がある。ただし後者はエージェント数がワークグループ上限（1024）に制限される。汎用性（エージェント数を増やせる）と速度（同期を減らす）のトレードオフ——協調計算の設計は、この綱引きになる。

以上が、エージェント市場のバブルを生んだ協調計算の実装だ。要点をまとめると——スレッド間の集約はアトミック加算（floatは整数化して対応）、ステップ間の同期はJS側の反復ディスパッチ。そして、この反復同期こそが性能とデバイス依存性の源泉だった。#1の独立並列が「並列度がすべて」だったのに対し、協調計算は「いかに同期を減らすか」が腕の見せ所になる。計算の質が変われば、最適化の発想もまるごと変わる——それがこのシリーズで一貫して見えてきたことだ。

▶ このページで実際に試す — 協調計算ベンチ

あなたのデバイスで CPU と WebGPU の市場シミュレーション（500ステップ）を実行します。協調計算はGPUアーキテクチャ依存が強い——あなたのGPUではどうか、確かめてください。

WebGPU CPU

エージェント数:

※ 初回はシェーダーのコンパイル時間を含み遅めに出ることがあります。2回目以降が実力値に近くなります。小さいエージェント数ではCPUが速く、大きくなるとGPUが有利に——そのクロスオーバーがどこに来るかは、デバイス次第です。

WGSLで協調計算を実装する— アトミック集計とステップ反復