WebAssemblyでGPU計算は可能？

---

## はじめに

[前回の記事](https://qiita.com/sukimaengineer/items/e44dd3998f4a4d49422b)では、**ブラウザからWebAssembly（Wasm）モジュールを呼び出す**仕組みや、WebAssemblyが同一プロセスで動作するメリット・セキュリティを解説しました。その続編として、今回は**WebAssemblyでGPUを使って高速計算を行う**方法、そして**CUDA/ROCmなどのネイティブGPU技術やOpenMP/ACCのGPUオフロード**が使えるのかをまとめます。

- **WebGLやWebGPU**を通じたブラウザGPGPU  
- **WASIサーバーサイドWasm**でGPUアクセスが可能になる未来  
- **CUDAやROCm、OpenMP GPU offload、OpenACC**をWasmで活用できるか？  
- **HPC（High-Performance Computing）分野でのWasm活用**はどうなる？

現時点での現状・課題・展望を、なるべく詳しく紹介していきます。
なお、WebAssemblyでのSIMD対応についてはこちらの記事を書きました。[WebAssemblyのSIMD対応の現状とWasm2.0での進化](https://qiita.com/sukimaengineer/items/10d408d8a6ffcff0773a)

---

## 1. まず結論：ブラウザWasmでネイティブなGPU利用はまだ無い

### 1.1 WebAssemblyの仕様に「GPU命令」は含まれていない

- **現状のWasm**（2025年）には、ネイティブGPUドライバを直接呼び出す仕組みはありません。  
- ブラウザでは**WebGL**や**WebGPU**といった高レベルAPIを**JavaScript**が操作し、それをWasm側が**ラッパ（import）** として呼ぶだけです。

### 1.2 CUDAやROCmを直接呼べない

- **NVIDIA CUDA**や**AMD ROCm**などのGPUドライバレベルAPIは、ブラウザWasm環境では**直接使えません**。  
- セキュリティ的に、ブラウザがOSネイティブドライバを操るのは想定外で、**WebGPU**等の抽象化APIを介す設計だからです。

---

## 2. WebGLとWebGPUを使うのが実用的

### 2.1 ブラウザでのGPU計算は**WebGL/WebGPU**経由

- **WebGL**: OpenGL ES相当の3D APIをブラウザで使う仕組み。GPGPU的に応用する場合は、シェーダーを駆使して行列演算や画像処理を行うテクニックがあります。  
- **WebGPU**: 次世代のブラウザGPU API。Vulkan/Metal/D3D12に近い低レベル制御を提供し、Computeパイプラインも用意。2023〜2024年で主要ブラウザに実装が進行中。

```mermaid
flowchart TD
    subgraph Browser
    A[Wasm Module<br>GPGPU code] --Call JS Interop--> B[JavaScript Glue<br>WebGL/WebGPU]
    B --GPU commands--> C[GPU driver in browser]
    end
    C --> D[Hardware GPU]
```

- **Wasmモジュール**から**JavaScript**のWebGL/WebGPUラッパを呼ぶ構造  
- メモリコピーやシェーダー設定も基本的に**JSが司る**形になります。
- WebGPUはComputeパイプラインがあるため、**描画を伴わない科学計算**にも適しています。

---

### 2.2 **WebAssemblyからWebGLでGPU計算する**簡単なサンプルコード

以下は、**Emscripten**を使ってC/C++からWebGLを呼び出し、フラグメントシェーダーで簡単なGPGPU的処理を行う**最小限の例**です。  
シェーダーの内容は非常に単純ですが、**ピクセルごとに (r+0.1, g+0.1, b+0.1) のように色を計算**して返すイメージだと考えてください。

#### A. C++サイド（mygpgpu.cpp）
```cpp
#include <emscripten.h>
#include <emscripten/html5.h>
#include <string>
#include <iostream>

static const char* vertexShaderSource =
R"(
attribute vec2 aPos;
void main() {
  gl_Position = vec4(aPos, 0.0, 1.0);
}
)";

static const char* fragmentShaderSource =
R"(
precision mediump float;
uniform vec3 uColor;
void main() {
  // GPGPU的にピクセル単位で計算する例
  // 単純に uColor を加算するだけ
  gl_FragColor = vec4(uColor.r + 0.1, uColor.g + 0.1, uColor.b + 0.1, 1.0);
}
)";

// Compile a shader
GLuint compileShader(GLenum type, const char* src) {
    GLuint shader = glCreateShader(type);
    glShaderSource(shader, 1, &src, nullptr);
    glCompileShader(shader);
    // check compile errors, omitted for brevity
    return shader;
}

// Create a program from vertex & fragment shader
GLuint createProgram(const char* vsSrc, const char* fsSrc) {
    GLuint vs = compileShader(GL_VERTEX_SHADER, vsSrc);
    GLuint fs = compileShader(GL_FRAGMENT_SHADER, fsSrc);
    GLuint prog = glCreateProgram();
    glAttachShader(prog, vs);
    glAttachShader(prog, fs);
    glLinkProgram(prog);
    // check link errors, omitted
    glDeleteShader(vs);
    glDeleteShader(fs);
    return prog;
}

// A simple function to run "GPU calculation" and read back pixels
extern "C" {
EMSCRIPTEN_KEEPALIVE
void runWebGLCompute(float r, float g, float b) {
    // Create WebGL context
    EmscriptenWebGLContextAttributes attr;
    emscripten_webgl_init_context_attributes(&attr);
    attr.alpha = false;
    attr.depth = false;
    attr.stencil = false;
    EMSCRIPTEN_WEBGL_CONTEXT_HANDLE ctx = emscripten_webgl_create_context("#canvas", &attr);
    emscripten_webgl_make_context_current(ctx);

    // compile & link shaders
    GLuint program = createProgram(vertexShaderSource, fragmentShaderSource);
    glUseProgram(program);

    // Setup geometry: a simple fullscreen quad
    static const float verts[6] = {
      -1.f, -1.f,  // bottom-left
       1.f, -1.f,  // bottom-right
      -1.f,  1.f   // top-left
      // we could add top-right too, but a second triangle is enough for a rectangle
    };
    GLuint vao;
    glGenBuffers(1, &vao);
    glBindBuffer(GL_ARRAY_BUFFER, vao);
    glBufferData(GL_ARRAY_BUFFER, sizeof(verts), verts, GL_STATIC_DRAW);

    GLint aPosLoc = glGetAttribLocation(program, "aPos");
    glEnableVertexAttribArray(aPosLoc);
    glVertexAttribPointer(aPosLoc, 2, GL_FLOAT, GL_FALSE, 0, 0);

    // pass uniform color
    GLint colorLoc = glGetUniformLocation(program, "uColor");
    glUniform3f(colorLoc, r, g, b);

    // draw
    glClearColor(0,0,0,1);
    glClear(GL_COLOR_BUFFER_BIT);
    glDrawArrays(GL_TRIANGLES, 0, 3);

    // read back the GPU results
    const int width = 256, height = 256;
    unsigned char* pixels = new unsigned char[width * height * 4];
    glReadPixels(0, 0, width, height, GL_RGBA, GL_UNSIGNED_BYTE, pixels);

    // Just print out a few pixel values
    std::cout << "Pixel(0,0): ("
              << (int)pixels[0] << ","
              << (int)pixels[1] << ","
              << (int)pixels[2] << ")" << std::endl;

    delete[] pixels;

    // cleanup
    glDeleteProgram(program);
    emscripten_webgl_make_context_current(0);
    emscripten_webgl_destroy_context(ctx);
}
}
```

**ポイント**:

- `emscripten_webgl_create_context` で `<canvas id="canvas">` をWebGLコンテキストとして初期化  
- シェーダーをコンパイルし、フラグメントシェーダーで `(r+0.1, g+0.1, b+0.1)` な色を計算  
- `glReadPixels` でGPU側の結果を取り出し（ある意味GPGPUの出力）。実際は行列演算やテクスチャ計算に応用可能  
- `_cdecl`やC++呼び出しの形式は省略。`EMSCRIPTEN_KEEPALIVE`でJSからアクセス可能に

#### B. コンパイル＆HTML

```bash
emcc mygpgpu.cpp -o mygpgpu.js -s MODULARIZE=1 -s EXPORT_ALL=1 \
  -s USE_WEBGL2=1 -s FULL_ES3=1
```

生成された `mygpgpu.js` と `mygpgpu.wasm` をHTMLと同じフォルダに配置し、以下のHTMLで呼び出します:

```html
<!DOCTYPE html>
<html>
<head><meta charset="utf-8"></head>
<body>
<canvas id="canvas" width="256" height="256"></canvas>
<script src="mygpgpu.js"></script>
<script>
Mygpgpu().then(module => {
  // run GPU "calculation"
  module.ccall('runWebGLCompute', null, ['number','number','number'], [0.2, 0.5, 0.7]);
});
</script>
</body>
</html>
```

- `module.ccall('runWebGLCompute', ...)` でC++の関数 `runWebGLCompute(float r, float g, float b)` を呼び出し、WebGLを使ってピクセル計算→`glReadPixels`で結果を取得して一部を `std::cout` に出力。
- 実際はJavaScriptコンソールには出ず、**Emscriptenのstdout**としてDevToolsに表示されるか、console.logにリダイレクトされる場合があります。

この例は非常に単純化されていますが、**Wasm経由でWebGLシェーダーを呼び出し、GPU計算（擬似GPGPU）をする**流れを示しています。

---

### 2.3 **WebAssemblyからWebGPUでGPU計算する**簡単なサンプルコード

以下は、**Emscripten**を使ってC/C++コードをWasmにコンパイルし、**WebGPU**のComputeパイプラインをJavaScript経由で呼び出す**最小限の例**です。WebGPUがまだ実験的であるため、**ブラウザのフラグ設定**や**Emscriptenの実験オプション**が必要となります。

#### A. C++サイド（mywebgpu.cpp）
```cpp
#include <emscripten.h>
#include <string>
#include <iostream>

// このサンプルはC++からWebGPU APIを呼ぶというより、
// JavaScript側でWebGPUを扱い、Wasmコードは補助的な関数のみを持つ形。
// 実際には Emscripten が webgpu.h などを提供する実験APIがありますが、
// ここでは簡素化して JS 経由で "gpuCompute" を呼ぶデモを想定。

extern "C" {
EMSCRIPTEN_KEEPALIVE
float computeValue(float x, float y) {
    // ここでは単に JS(=WebGPU)に渡すための値を加工
    // 実用的にはWasm側でデータ生成→JavaScript→WebGPUへ転送などが考えられる
    return x * y + 1.0f;
}

// JavaScriptの"gpuComputeResult"を受け取り、stdoutに出力
EMSCRIPTEN_KEEPALIVE
void printComputeResult(float result) {
    std::cout << "WebGPU Compute returned: " << result << std::endl;
}
}
```

- この例では、**Wasmに組み込むC/C++コードはシンプル**で、`computeValue(x,y)`という関数で計算するだけ。  
- 実際の**WebGPU操作**はJavaScript側で行い、**Wasmが出力する`printComputeResult`**で受け取って出力するイメージ。

#### B. コンパイル

```bash
emcc mywebgpu.cpp -o mywebgpu.js -s MODULARIZE=1 -s EXPORT_ALL=1 \
  -s USE_WEBGPU=1 -s WASM=1 -O3
```

- `-s USE_WEBGPU=1` は現状**実験オプション**。ブラウザでWebGPUを有効にする必要あり（Chromeで `chrome://flags/#enable-unsafe-webgpu` など）。
- 生成された `mywebgpu.js` と `mywebgpu.wasm` をHTMLと同フォルダに置く。

#### C. HTML & JavaScript (mywebgpu.html)

```html
<!DOCTYPE html>
<html>
<head><meta charset="utf-8"></head>
<body>
<script src="mywebgpu.js"></script>
<script>
(async function() {
  // Initialize the Emscripten module
  const Module = await Mywebgpu();

  // 1) Call a Wasm function to compute some data
  const val = Module.ccall('computeValue', 'number', ['number','number'], [2.0, 3.5]);
  console.log("Wasm gave value: ", val);

  // 2) Use WebGPU to process 'val' with a simple compute shader
  if (!('gpu' in navigator)) {
    console.warn("WebGPU not supported or disabled");
    return;
  }
  const adapter = await navigator.gpu.requestAdapter();
  const device = await adapter.requestDevice();

  // Minimal compute shader to add 2.0 to the input
  const shaderCode = `
    @group(0) @binding(0) var<storage, read> inBuffer: array<f32>;
    @group(0) @binding(1) var<storage, write> outBuffer: array<f32>;

    @compute @workgroup_size(1)
    fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
      let idx = global_id.x;
      let inputVal = inBuffer[idx];
      outBuffer[idx] = inputVal + 2.0;
    }
  `;

  // Create shader module
  const shaderModule = device.createShaderModule({ code: shaderCode });

  // Create input/output buffers
  const arraySize = 1;
  const inputData = new Float32Array([val]);
  const inputBuffer = device.createBuffer({
    size: inputData.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
  });
  device.queue.writeBuffer(inputBuffer, 0, inputData);

  const outputBuffer = device.createBuffer({
    size: inputData.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
  });

  // Create bind group layout & pipeline
  const bindGroupLayout = device.createBindGroupLayout({
    entries: [
      { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
      { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
    ],
  });
  const pipeline = device.createComputePipeline({
    layout: device.createPipelineLayout({ bindGroupLayouts: [bindGroupLayout] }),
    compute: { module: shaderModule, entryPoint: 'main' }
  });

  // Create bind group
  const bindGroup = device.createBindGroup({
    layout: bindGroupLayout,
    entries: [
      { binding: 0, resource: { buffer: inputBuffer } },
      { binding: 1, resource: { buffer: outputBuffer } },
    ],
  });

  // Encode commands
  const commandEncoder = device.createCommandEncoder();
  {
    const pass = commandEncoder.beginComputePass();
    pass.setPipeline(pipeline);
    pass.setBindGroup(0, bindGroup);
    pass.dispatchWorkgroups(arraySize);
    pass.end();
  }

  // Create buffer for readback
  const readbackBuffer = device.createBuffer({
    size: inputData.byteLength,
    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
  });

  commandEncoder.copyBufferToBuffer(
    outputBuffer, 0, readbackBuffer, 0, inputData.byteLength
  );

  // Submit
  const commandBuffer = commandEncoder.finish();
  device.queue.submit([commandBuffer]);

  // Readback
  await readbackBuffer.mapAsync(GPUMapMode.READ);
  const arrayOut = new Float32Array(readbackBuffer.getMappedRange().slice(0));
  const gpuResult = arrayOut[0];
  readbackBuffer.unmap();

  console.log("WebGPU compute result: ", gpuResult);

  // 3) Pass the final GPU result back to Wasm
  Module.ccall('printComputeResult', null, ['number'], [gpuResult]);
})();
</script>
</body>
</html>
```

**説明**:

1. **Wasm (mywebgpu.wasm)** を初期化し、`computeValue(x,y)` を呼び出して `val` を取得（C/C++関数内部は簡単な計算のみ）。  
2. **WebGPU** でComputeシェーダーを定義し、`val`を入力バッファとして加算処理（`input + 2.0f`）するだけの最小プログラムを実行。  
3. GPU結果 (`gpuResult`) を**Wasm**側の関数 `printComputeResult` に渡して表示。結果的に **(x*y +1.0) + 2.0** になる。  

これにより、**Wasm→JS(WebGPU)→GPU** という一連の流れで、簡単な数値計算が行われます。現時点では**実験的API**かつ**特定ブラウザフラグ**が必要なので、動作確認にご注意ください。

---

## 3. サーバーサイドWasm（WASI）でのGPU利用は？

### 3.1 WASI標準でGPUアクセスはまだ議論中

- サーバーサイドのWasmランタイム（Wasmtime, WasmEdgeなど）で**GPUアクセス**を正式に定義するWASI拡張は未策定。  
- 一部独自拡張で**Vulkan/CUDA**に繋げる試みはあるが、広く使われているわけではありません。

### 3.2 将来の可能性

- **WASI GPU Extension**が標準化し、OpenCL/Metal/Vulkan等の抽象APIを提供できれば、**ネイティブGPUに近い利用**がWasm上で可能となる可能性があります。  
- ただし、まだ初期段階で、実用化には時間がかかりそうです。

---

## 4. CUDAやROCmをWebAssemblyで使いたい？

### 4.1 ブラウザでは無理

- **NVIDIA CUDA**や**AMD ROCm**を**ブラウザでWasm**として動かすのは、**OSドライバへの直接アクセス**が必要であり不可能。  
- EmscriptenでCUDA/ROCmコードをコンパイルしても、呼び出し先が無いのでGPU処理は失敗します。

### 4.2 サーバーサイドでの実験的取り組み

- 研究レベルで、**Wasmランタイム→CUDA**などを独自拡張するプロジェクトがある一方、**メインストリームのWASI**には入っていません。  
- したがって、**CUDA/ROCmを使うなら現状はネイティブ環境が無難**。

---

## 5. OpenMP GPUオフロードやOpenACCは使えるのか？

### 5.1 背景

- **OpenMP GPU offload**: `#pragma omp target`などを使い、C/C++/Fortranコードの一部をGPUカーネルとしてコンパイラが自動生成し、CUDA/OpenCL APIを呼ぶ仕組み。  
- **OpenACC**: 同様に、`#pragma acc`でアクセラレータ（GPU）向けコードを生成。

### 5.2 Wasm側の制限

- ブラウザWasmではネイティブGPUドライバがなく、OpenMP/ACCの`target`先が存在しない。  
- サーバーサイドWasmも**WASI標準でGPU無**いため、**OpenMP/ACCのGPUオフロード機能は機能しない**。

### 5.3 結論: 不可

- つまり、**OpenMP GPU offload**や**OpenACC**のGPU部分をWasmで動かすのは、現状不可能。CPU並列(`-fopenmp`のスレッド活用)は一部対応していますが、GPUには行けません。

---

## 6. HPC（High-Performance Computing）環境との比較

HPC分野では、CUDAやOpenMP、MPIなどを使い、**ネイティブOS上**でGPU・並列計算を行うのが主流です。Wasmはサンドボックスされ移植性が高いものの、**ネイティブドライバ呼び出し**や**GPUオフロード**が苦手です。

```mermaid
flowchart LR
    subgraph HPC Native
      A[C/C++ HPC code<br>OpenMP, CUDA, MPI...]
      A --> B[Native GPU driver<br>CUDA/ROCm...]
      B --> C[Hardware GPU]
    end

    subgraph HPC with Wasm
      D[C/C++ HPC code<br>Compiled to Wasm]
      D --> E[GPU Access]
      E --> F[Limited or <br>no direct driver call]
      F --> G[Possible fallback<br>WebGPU/WebGL or no GPU]
    end
```

- **ネイティブHPC**: OS上で自由にGPUドライバを操作。高性能かつ柔軟  
- **Wasm HPC**: 移植性やセキュリティは高いが、**GPU制御API**がないため、**実用性は限定的**（特にOpenMP GPU offloadやCUDA/ROCmは不可）

---

## 7. 全体まとめ表

| 項目                         | ブラウザWasm                             | サーバーサイドWasm(WASI)                             | 現状の可能性                                        |
|-----------------------------|-----------------------------------------|------------------------------------------------------|-----------------------------------------------------|
| **WebGL / WebGPU**          | JSラッパ経由で呼び出し                  | 通常は提供されない<br>(ブラウザAPI特有)               | ブラウザGPU計算の実用的手段                         |
| **CUDA / ROCm**             | 直接利用不可（ネイティブドライバ呼び出し不可） | 公式サポート無し<br>一部独自拡張は実験レベル           | 将来のWASI拡張に期待だが現状無理                    |
| **OpenMP GPU offload**      | CPU並列のみ可。GPUターゲットは不可       | 同様に不可<br>GPUドライバAPI無                       | 実質使えない                                        |
| **OpenACC**                 | 同上                                     | 同上                                                 | 事実上利用不可                                      |
| **GPGPU via WebGL/WebGPU**  | Wasm→JS→(WebGL/WebGPU)でComputeシェーダー | サーバーサイドでは一般に無                            | ブラウザでの高速並列処理                            |
| **WASI + GPU**（将来案）    | ブラウザ外の話                           | 仕様議論中で標準化されていない                        | 将来的発展に期待。現状は実用的実装少                |
| **HPCネイティブ環境**       | そもそもWasmとは別                       | CUDA/ROCmやOpenMP GPU offload可                      | 高性能だが移植性低い                                 |

---

## よくある質問（FAQ）

1. **Q. WebAssemblyブラウザ環境でCUDAのようなネイティブGPUを直接叩けませんか？**  
   **A.** 不可です。ブラウザはセキュリティ的にOSドライバを隠蔽しており、**WebGL/WebGPU**を通じた高レベル制御が唯一の道。

2. **Q. サーバーサイドWasmでCUDA/ROCmが動く日は来るのでしょうか？**  
   **A.** 公式WASIには未定義ですが、一部独自拡張の研究はあります。**標準的に使える形**になるにはまだ時間が必要です。

3. **Q. OpenMPの`target`機能でGPUオフロードしたC/C++をWasmへコンパイルすると？**  
   **A.** GPU向けコードが無いのでエラーになるか、CPU側だけがビルドされる形となり、GPUオフロードは動作しません。

4. **Q. HPCクラスタでWasmを導入するメリットは？**  
   **A.** セキュアなサンドボックスや移植性が得られますが、**ネイティブGPU対応やMPI機能**は不十分。現在のHPC向けワークロードとは相性が悪い面が多いです。

5. **Q. WebGPUのComputeシェーダーをC/C++で書きたい**  
   **A.** Emscriptenなどで`--use-webgpu`の実験オプションを使うと、C/C++からWebGPUを呼ぶラッパがある程度使える。**WGSL**シェーダー記述をどう管理するかなど要検討。

6. **Q. メモリ転送のオーバーヘッドはどの程度？**  
   **A.** Wasm→JS→GPUでのバッファ転送は多少のオーバーヘッドがあるが、**大規模行列演算など**でGPU計算のメリットが大きければペイする可能性が高い。

---

## まとめ

- **ブラウザ上でGPUを使う**なら、現状は**WebGLまたはWebGPU**を**JavaScriptラッパ**を通じて呼び出す方法がメイン。  
- **CUDA/ROCm/OpenMP GPU offload/OpenACC**といったネイティブ技術は、Wasmブラウザ環境では使えない。サーバーサイドWasm（WASI）でも標準サポートなし。  
- 今後のWASI拡張や独自ランタイムでネイティブGPU対応が進めば、**将来的には可能性**が広がるかもしれないが、現状は研究・実験段階。  
- HPC分野（大規模並列計算）でもWasmが脚光を浴びつつあるが、**GPU活用**の面ではネイティブ環境に追いついていないのが実情。  

:::note info
- もしブラウザで高速並列計算をしたいなら、**WebGPU Computeシェーダー**をJS越しに呼ぶ形が最も有望です。また、SIMDを使う方法もあります。なお、WebAssemblyでのSIMD対応についてはこちらの記事を書きました。[WebAssemblyのSIMD対応の現状とWasm2.0での進化](https://qiita.com/sukimaengineer/items/10d408d8a6ffcff0773a)
:::
---

### 参考リンク

- [WebGPU Explainer (GitHub)](https://github.com/gpuweb/gpuweb)  
- [EmscriptenとWebGPU](https://emscripten.org/docs/porting/graphics.html#webgpu)  
- [WASI proposals for GPU access (wgpu-native)](https://github.com/gfx-rs/wgpu)  
- [OpenACC公式サイト](https://www.openacc.org/)  
- [OpenMP.org - GPU Offload Docs](https://www.openmp.org/)  
- [CUDA / ROCm 公式ドキュメント](https://developer.nvidia.com/cuda-zone), [https://rocmdocs.amd.com/en/latest/](https://rocmdocs.amd.com/en/latest/)

---

以上