f8b21af4be
10_Wiki/Topics 대규모 정리: - 오류 캡처/미완성 stub 문서 227개 제거 - 교차폴더 중복 43클러스터 병합 (63파일 → redirect) - 링크명 정규화: 깨진 링크 수정·redirect 직결·개념 매핑 ~2,400건 - 카테고리 MOC 6개 신규 생성 - Graph 섹션 미해결 related-keyword 링크 10,058건 제거 Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
233 lines
7.4 KiB
Markdown
233 lines
7.4 KiB
Markdown
---
|
|
id: wiki-2026-0508-webgpu-compute-shaders
|
|
title: WebGPU Compute Shaders
|
|
category: 10_Wiki/Topics
|
|
status: verified
|
|
canonical_id: self
|
|
aliases: [WebGPU Compute, WGSL Compute, WebGPU Compute Shader]
|
|
duplicate_of: none
|
|
source_trust_level: A
|
|
confidence_score: 0.9
|
|
verification_status: applied
|
|
tags: [webgpu, compute, gpgpu, wgsl]
|
|
raw_sources: []
|
|
last_reinforced: 2026-05-10
|
|
github_commit: pending
|
|
tech_stack:
|
|
language: WGSL/JavaScript
|
|
framework: WebGPU
|
|
---
|
|
|
|
# WebGPU Compute Shaders
|
|
|
|
## 매 한 줄
|
|
> **"매 GPGPU 의 web — 매 finally"**. WebGPU compute shader 매 WGSL 의 modern Vulkan/Metal-style API 의 GPU 의 general compute. 2026 매 ML inference (ONNX Runtime Web, transformers.js GPU backend) / physics / image processing 의 standard.
|
|
|
|
## 매 핵심
|
|
|
|
### 매 vs WebGL
|
|
- **Compute shader**: WebGL X / WebGPU O — 매 game changer.
|
|
- **WGSL**: 매 statically typed, Rust-inspired — GLSL 보다 modern.
|
|
- **Bind groups**: explicit resource binding — 매 Vulkan-like.
|
|
- **Async pipeline**: command encoder + submit — 매 lower CPU overhead.
|
|
|
|
### 매 workgroup
|
|
- **`@workgroup_size(x, y, z)`**: thread block — typically 64 / 128 / 256 threads.
|
|
- **Dispatch**: `dispatchWorkgroups(gx, gy, gz)` — total threads = workgroup * dispatch.
|
|
- **Shared memory**: `var<workgroup>` — 매 fast on-chip.
|
|
- **Synchronization**: `workgroupBarrier()` — 매 within workgroup.
|
|
|
|
### 매 응용
|
|
1. **ML inference**: matrix multiply / attention — transformers.js / ONNX Runtime Web.
|
|
2. **Image processing**: blur / upscale (FSR-style) / segmentation.
|
|
3. **Physics**: cloth / particles / fluid sim.
|
|
4. **Sort / scan / reduce**: parallel primitives.
|
|
|
|
## 💻 패턴
|
|
|
|
### Init device + buffer
|
|
```js
|
|
const adapter = await navigator.gpu.requestAdapter();
|
|
const device = await adapter.requestDevice();
|
|
|
|
const input = new Float32Array(1024).fill(1);
|
|
const inBuf = device.createBuffer({
|
|
size: input.byteLength,
|
|
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
|
|
});
|
|
device.queue.writeBuffer(inBuf, 0, input);
|
|
|
|
const outBuf = device.createBuffer({
|
|
size: input.byteLength,
|
|
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
|
|
});
|
|
|
|
const readBuf = device.createBuffer({
|
|
size: input.byteLength,
|
|
usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
|
|
});
|
|
```
|
|
|
|
### Compute shader 매 element-wise
|
|
```js
|
|
const module = device.createShaderModule({
|
|
code: `
|
|
@group(0) @binding(0) var<storage, read> input: array<f32>;
|
|
@group(0) @binding(1) var<storage, read_write> output: array<f32>;
|
|
|
|
@compute @workgroup_size(64)
|
|
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
|
let i = gid.x;
|
|
if (i >= arrayLength(&input)) { return; }
|
|
output[i] = input[i] * 2.0 + 1.0;
|
|
}
|
|
`,
|
|
});
|
|
|
|
const pipeline = device.createComputePipeline({
|
|
layout: 'auto',
|
|
compute: { module, entryPoint: 'main' },
|
|
});
|
|
|
|
const bindGroup = device.createBindGroup({
|
|
layout: pipeline.getBindGroupLayout(0),
|
|
entries: [
|
|
{ binding: 0, resource: { buffer: inBuf } },
|
|
{ binding: 1, resource: { buffer: outBuf } },
|
|
],
|
|
});
|
|
|
|
const enc = device.createCommandEncoder();
|
|
const pass = enc.beginComputePass();
|
|
pass.setPipeline(pipeline);
|
|
pass.setBindGroup(0, bindGroup);
|
|
pass.dispatchWorkgroups(Math.ceil(1024 / 64));
|
|
pass.end();
|
|
enc.copyBufferToBuffer(outBuf, 0, readBuf, 0, input.byteLength);
|
|
device.queue.submit([enc.finish()]);
|
|
|
|
await readBuf.mapAsync(GPUMapMode.READ);
|
|
const result = new Float32Array(readBuf.getMappedRange().slice(0));
|
|
readBuf.unmap();
|
|
```
|
|
|
|
### Matmul 매 tiled (workgroup shared)
|
|
```wgsl
|
|
const TILE = 16u;
|
|
@group(0) @binding(0) var<storage, read> A: array<f32>;
|
|
@group(0) @binding(1) var<storage, read> B: array<f32>;
|
|
@group(0) @binding(2) var<storage, read_write> C: array<f32>;
|
|
@group(0) @binding(3) var<uniform> dims: vec3<u32>; // M, N, K
|
|
|
|
var<workgroup> Asub: array<array<f32, 16>, 16>;
|
|
var<workgroup> Bsub: array<array<f32, 16>, 16>;
|
|
|
|
@compute @workgroup_size(16, 16)
|
|
fn matmul(@builtin(global_invocation_id) gid: vec3<u32>,
|
|
@builtin(local_invocation_id) lid: vec3<u32>) {
|
|
let row = gid.y; let col = gid.x;
|
|
var sum = 0.0;
|
|
let tiles = (dims.z + TILE - 1u) / TILE;
|
|
for (var t = 0u; t < tiles; t++) {
|
|
Asub[lid.y][lid.x] = A[row * dims.z + t * TILE + lid.x];
|
|
Bsub[lid.y][lid.x] = B[(t * TILE + lid.y) * dims.x + col];
|
|
workgroupBarrier();
|
|
for (var k = 0u; k < TILE; k++) { sum += Asub[lid.y][k] * Bsub[k][lid.x]; }
|
|
workgroupBarrier();
|
|
}
|
|
C[row * dims.x + col] = sum;
|
|
}
|
|
```
|
|
|
|
### Reduction (parallel sum)
|
|
```wgsl
|
|
var<workgroup> shared_data: array<f32, 256>;
|
|
|
|
@compute @workgroup_size(256)
|
|
fn reduce(@builtin(local_invocation_id) lid: vec3<u32>,
|
|
@builtin(workgroup_id) wid: vec3<u32>) {
|
|
let tid = lid.x;
|
|
let i = wid.x * 256u + tid;
|
|
shared_data[tid] = select(0.0, input[i], i < arrayLength(&input));
|
|
workgroupBarrier();
|
|
|
|
var s = 128u;
|
|
loop {
|
|
if (s == 0u) { break; }
|
|
if (tid < s) { shared_data[tid] += shared_data[tid + s]; }
|
|
workgroupBarrier();
|
|
s = s / 2u;
|
|
}
|
|
if (tid == 0u) { partial[wid.x] = shared_data[0]; }
|
|
}
|
|
```
|
|
|
|
### Image processing — texture in, texture out
|
|
```wgsl
|
|
@group(0) @binding(0) var inputTex: texture_2d<f32>;
|
|
@group(0) @binding(1) var outputTex: texture_storage_2d<rgba8unorm, write>;
|
|
|
|
@compute @workgroup_size(8, 8)
|
|
fn blur(@builtin(global_invocation_id) gid: vec3<u32>) {
|
|
let dims = textureDimensions(inputTex);
|
|
if (gid.x >= dims.x || gid.y >= dims.y) { return; }
|
|
var sum = vec4<f32>(0.0);
|
|
for (var dy = -1; dy <= 1; dy++) {
|
|
for (var dx = -1; dx <= 1; dx++) {
|
|
sum += textureLoad(inputTex, vec2<i32>(gid.xy) + vec2(dx, dy), 0);
|
|
}
|
|
}
|
|
textureStore(outputTex, vec2<i32>(gid.xy), sum / 9.0);
|
|
}
|
|
```
|
|
|
|
### transformers.js GPU 매 inference
|
|
```js
|
|
import { pipeline, env } from '@xenova/transformers';
|
|
env.backends.onnx.wasm.simd = true;
|
|
const generator = await pipeline('text-generation', 'Xenova/Llama-3.2-1B', {
|
|
device: 'webgpu', // 매 WebGPU compute backend
|
|
dtype: 'q4',
|
|
});
|
|
const out = await generator('Hello', { max_new_tokens: 50 });
|
|
```
|
|
|
|
## 매 결정 기준
|
|
| 상황 | Approach |
|
|
|---|---|
|
|
| ML inference in browser | WebGPU compute (transformers.js / ORT-Web) |
|
|
| Particle / cloth sim | Compute shader |
|
|
| Image filter | Compute or fragment shader (compute 매 cleaner) |
|
|
| Sort / scan | Compute (no fragment hack) |
|
|
| Wide compatibility | WebGL fallback (no compute) |
|
|
| Mobile | WebGPU 매 iOS 18+ / Android Chrome 121+ |
|
|
|
|
**기본값**: WebGPU compute + WGSL + tiled algorithms + transformers.js for ML.
|
|
|
|
## 🔗 Graph
|
|
- 부모: [[WebGPU]] · [[GPGPU]]
|
|
- 변형: [[CUDA]]
|
|
- 응용: [[Threejs WebGPURenderer]]
|
|
- Adjacent: [[WGSL]] · [[Web Worker (웹 워커)]] · [[WebAssembly]]
|
|
|
|
## 🤖 LLM 활용
|
|
**언제**: GPU compute in browser — ML / physics / image / parallel reduce-scan.
|
|
**언제 X**: trivial work — JS / Worker 매 충분 / CPU / Safari iOS <18 — WebGPU 의 unavailable.
|
|
|
|
## ❌ 안티패턴
|
|
- **Tiny dispatch**: <1024 threads 매 launch overhead 의 net loss.
|
|
- **No barrier on shared memory**: race condition — `workgroupBarrier()` 의 between read/write.
|
|
- **`mapAsync` per frame**: GPU↔CPU sync 매 stall — pipeline + double-buffer.
|
|
- **Workgroup size 32**: 매 너무 small — 64-256 의 사용 (warp/wave occupancy).
|
|
- **No bounds check**: out-of-range invocation 매 garbage write.
|
|
|
|
## 🧪 검증 / 중복
|
|
- Verified (W3C WebGPU spec / WGSL spec / transformers.js docs).
|
|
- 신뢰도 A.
|
|
|
|
## 🕓 Changelog
|
|
| 날짜 | 변경 |
|
|
|---|---|
|
|
| 2026-05-08 | Phase 1 |
|
|
| 2026-05-10 | Manual cleanup — workgroup, tiled matmul, reduction, image, transformers.js |
|