CPU vs WebGPU
並列化の威力を測る

シリーズ第2回（速度編）です。基本編では「乱数が理論値に収束する」ことを確認しました。 今回は同じ計算をCPUとWebGPUで実行し、速度差を測定します。 さらに、全パスの平均をGPU上で高速に集計する並列リダクションを実装します。

オプション価格評価の速度は、実務で直接的な価値を持ちます。 「正確だが遅い」計算は、変化の速い市場では使い物になりません。

基本編では各パスのペイオフをバッファに書き出し、CPU側で合計していました。 しかしパスが数百万になると、この「GPUからCPUへの全データ転送 + CPU集計」がボトルネックになります。 そこで合計をGPU上で求めるのが並列リダクションです。

並列リダクションの実装を中心に解説します。基本編の乱数生成・株価計算はそのまま使います。

// ワークグループ内で共有するメモリ（64スレッド分）
var<workgroup> partialSum: array<f32, 64>;

// 各ワークグループの部分和を書き出す配列（要素数=ワークグループ数）
@group(0) @binding(0) var<storage, read_write> partials: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3u,
        @builtin(local_invocation_id)  lid: vec3u,
        @builtin(workgroup_id) wid: vec3u,
        @builtin(num_workgroups) nwg: vec3u) {
  let totalThreads = nwg.x * 64u;

  // ① このスレッドが担当する全パスのペイオフを合計（grid-stride）
  var localSum = 0.0;
  var i = gid.x;
  loop {
    if (i >= p.numPaths) { break; }
    localSum = localSum + computePayoff(i);  // 乱数→株価→ペイオフ
    i = i + totalThreads;
  }
  partialSum[lid.x] = localSum;

各スレッドはまず自分の担当パスのペイオフを localSum に集めます。 grid-stride で複数パスを担当するので、ここで既にかなりのデータが圧縮されます。

  // ② ワークグループ内でツリー状に合計（64→32→16→...→1）
  workgroupBarrier();
  var stride = 32u;
  loop {
    if (stride == 0u) { break; }
    if (lid.x < stride) {
      partialSum[lid.x] = partialSum[lid.x] + partialSum[lid.x + stride];
    }
    workgroupBarrier();  // 各段で全スレッドの完了を待つ
    stride = stride / 2u;
  }

  // ③ 代表スレッド（lid.x==0）が全体合計に加算
  if (lid.x == 0u) {
    // f32 を u32 にスケール変換して atomicAdd（後述のハマりどころ）
    atomicAdd(&globalSum, u32(partialSum[0] * 1000.0));
  }
}

ハイライト部分がツリー状リダクションの核心です。 stride を半分ずつ減らしながら、各段で隣接要素を加算します。 workgroupBarrier() で全スレッドの足並みを揃えないと、 まだ書かれていない値を読んでしまい結果が壊れます。

  // ③ 代表スレッドが「このワークグループの部分和」を配列に書き出す
  if (lid.x == 0u) {
    partials[wid.x] = partialSum[0];  // atomic不要・f32精度を維持
  }
}

// CPU側: 部分和の配列を合計（要素数 = ワークグループ数 ≤ 65535）
let sum = 0;
for (let i = 0; i < numWG; i++) sum += partials[i];
const price = (sum / numPaths) * Math.exp(-r * t);

当初は atomicAdd で1つの合計に足し込む設計でしたが、 WGSLのatomicは整数型のみで、f32を固定小数点化すると大パス数で u32オーバーフローが起きます（平均50・2000万パスで合計1e12 ≫ 43億）。
そこで各ワークグループの部分和を配列の別スロットに書き出し、 CPU側で合計する方式にしました。書き出すのは最大65,535個だけなので、 CPU合計も一瞬です。GPUで大部分を圧縮し、最後の僅かな集計だけCPUが担う実用的な折衷案です。