LoginSignup
1
0

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

@sukimaengineer

WebAssemblyでGPU計算は可能？

Last updated at Posted at 2025-01-29

はじめに

前回の記事では、ブラウザから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での進化

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年で主要ブラウザに実装が進行中。
  • 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）

#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

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で呼び出します:

<!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）

#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. コンパイル

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)

<!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オフロードが苦手です。

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

7. 全体まとめ表

項目 ブラウザWasm サーバーサイドWasm(WASI) 現状の可能性
WebGL / WebGPU JSラッパ経由で呼び出し 通常は提供されない
(ブラウザAPI特有)		 ブラウザGPU計算の実用的手段
CUDA / ROCm 直接利用不可（ネイティブドライバ呼び出し不可） 公式サポート無し
一部独自拡張は実験レベル		 将来のWASI拡張に期待だが現状無理
OpenMP GPU offload CPU並列のみ可。GPUターゲットは不可 同様に不可
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活用の面ではネイティブ環境に追いついていないのが実情。
  • もしブラウザで高速並列計算をしたいなら、WebGPU ComputeシェーダーをJS越しに呼ぶ形が最も有望です。また、SIMDを使う方法もあります。なお、WebAssemblyでのSIMD対応についてはこちらの記事を書きました。WebAssemblyのSIMD対応の現状とWasm2.0での進化

参考リンク

以上

1
0
1

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
1
0

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?