diff --git a/benchmarks/webgpu-scaling.html b/benchmarks/webgpu-scaling.html new file mode 100644 index 0000000..98cb928 --- /dev/null +++ b/benchmarks/webgpu-scaling.html @@ -0,0 +1,30 @@ + + + + + + WebGPU scaling benchmark · micarrsim + + + + +
+

192-channel reproducible benchmark

+

WebGPU DAS scaling

+

fixture를 준비하고 있습니다…

+
+ + + 전체 테스트 열기 +
+ 

+    

+    
+  
+
diff --git a/benchmarks/webgpu-scaling.js b/benchmarks/webgpu-scaling.js
new file mode 100644
index 0000000..99ef8ac
--- /dev/null
+++ b/benchmarks/webgpu-scaling.js
@@ -0,0 +1,183 @@
+import { distance3 } from "/packages/acoustics-core/math.js";
+import { synthesizeMicrophoneSignals } from "/packages/acoustics-core/propagation.js";
+import { computeDasWebGpu } from "/packages/webgpu-kernels/das.js";
+
+const status = document.querySelector("#benchmark-status");
+const output = document.querySelector("#benchmark-output");
+const runButton = document.querySelector("#run-benchmark");
+const downloadButton = document.querySelector("#download-result");
+const query = new URLSearchParams(location.search);
+const requestedRuns = Number.parseInt(query.get("runs") ?? "3", 10);
+const runCount = Number.isInteger(requestedRuns) ? Math.min(5, Math.max(1, requestedRuns)) : 3;
+let latestResult = null;
+
+async function loadJson(path) {
+  const response = await fetch(path);
+  if (!response.ok) throw new Error(`${path}를 불러오지 못했습니다 (${response.status}).`);
+  return response.json();
+}
+
+function median(values) {
+  const sorted = [...values].sort((a, b) => a - b);
+  const middle = Math.floor(sorted.length / 2);
+  return sorted.length % 2 === 1 ? sorted[middle] : (sorted[middle - 1] + sorted[middle]) / 2;
+}
+
+function round(value, digits = 3) {
+  const scale = 10 ** digits;
+  return Math.round(value * scale) / scale;
+}
+
+function cloneExperiment(experiment, size) {
+  return {
+    ...experiment,
+    id: `single-tone-${size}x${size}`,
+    scan: { ...experiment.scan, columns: size, rows: size },
+    beamforming: { ...experiment.beamforming, backend: "webgpu" },
+  };
+}
+
+function downloadJson(value, filename) {
+  const blob = new Blob([`${JSON.stringify(value, null, 2)}\n`], { type: "application/json" });
+  const url = URL.createObjectURL(blob);
+  const anchor = document.createElement("a");
+  anchor.href = url;
+  anchor.download = filename;
+  anchor.click();
+  setTimeout(() => URL.revokeObjectURL(url), 0);
+}
+
+async function runBenchmark() {
+  runButton.disabled = true;
+  downloadButton.disabled = true;
+  document.documentElement.dataset.benchmarkStatus = "running";
+  delete document.documentElement.dataset.benchmarkError;
+  latestResult = null;
+  globalThis.__micarrsimBenchmarkResult = null;
+  output.textContent = "";
+  try {
+    const [arrayDefinition, scene, baseExperiment] = await Promise.all([
+      loadJson("/assets/array-geometries/rectangular-16x12.array.json"),
+      loadJson("/assets/sample-scenes/single-tone.scene.json"),
+      loadJson("/assets/sample-scenes/single-tone.experiment.json"),
+    ]);
+    status.textContent = "192채널 마이크 신호를 한 번 합성하고 있습니다…";
+    const signalStartedAt = performance.now();
+    const synthesis = synthesizeMicrophoneSignals({ scene, arrayDefinition });
+    const signalGenerationMs = performance.now() - signalStartedAt;
+    let adapterInfo = null;
+
+    status.textContent = "WebGPU pipeline을 예열하고 있습니다…";
+    await computeDasWebGpu({
+      channels: synthesis.channels,
+      sampleRate: synthesis.sampleRate,
+      arrayDefinition,
+      scene,
+      experiment: cloneExperiment(baseExperiment, 16),
+    });
+
+    const resolutions = [];
+    for (const size of [50, 100, 200]) {
+      const samples = [];
+      for (let runIndex = 0; runIndex < runCount; runIndex += 1) {
+        status.textContent = `${size}×${size} · ${runIndex + 1}/${runCount} 실행 중…`;
+        const result = await computeDasWebGpu({
+          channels: synthesis.channels,
+          sampleRate: synthesis.sampleRate,
+          arrayDefinition,
+          scene,
+          experiment: cloneExperiment(baseExperiment, size),
+          onProgress: ({ stage, completedUnits, totalUnits }) => {
+            if (stage === "beamforming" && totalUnits) {
+              status.textContent =
+                `${size}×${size} · ${runIndex + 1}/${runCount} · GPU ${completedUnits}/${totalUnits}`;
+            }
+          },
+        });
+        adapterInfo = result.gpu.adapterInfo;
+        const localizationErrorM = distance3(result.bands[0].peak.position, scene.sources[0].position);
+        samples.push({
+          totalMs: result.timingMs.total,
+          stftMs: result.timingMs.stft,
+          gpuWallMs: result.timingMs.gpuWall,
+          beamformingMs: result.timingMs.beamforming,
+          localizationErrorM,
+          complexMacCount: result.workload.complexMacCount,
+          estimatedGpuWorkingBytes: result.gpu.batchPlan.estimatedGpuWorkingBytes,
+          binBatchSize: result.gpu.batchPlan.binBatchSize,
+          pointBatchSize: result.gpu.batchPlan.pointBatchSize,
+        });
+      }
+      resolutions.push({
+        grid: `${size}x${size}`,
+        pointCount: size * size,
+        runs: runCount,
+        medianTotalMs: round(median(samples.map((sample) => sample.totalMs))),
+        medianStftMs: round(median(samples.map((sample) => sample.stftMs))),
+        medianGpuWallMs: round(median(samples.map((sample) => sample.gpuWallMs))),
+        medianBeamformingMs: round(median(samples.map((sample) => sample.beamformingMs))),
+        localizationErrorM: round(median(samples.map((sample) => sample.localizationErrorM)), 6),
+        complexMacCount: samples[0].complexMacCount,
+        estimatedGpuWorkingMiB: round(samples[0].estimatedGpuWorkingBytes / 1024 ** 2),
+        binBatchSize: samples[0].binBatchSize,
+        pointBatchSize: samples[0].pointBatchSize,
+        samples: samples.map((sample) => ({
+          totalMs: round(sample.totalMs),
+          stftMs: round(sample.stftMs),
+          gpuWallMs: round(sample.gpuWallMs),
+          beamformingMs: round(sample.beamformingMs),
+        })),
+      });
+      output.textContent = JSON.stringify(resolutions, null, 2);
+    }
+
+    const baselineGpuMs = resolutions[0].medianGpuWallMs;
+    for (const resolution of resolutions) {
+      resolution.gpuWallScalingVs50 = round(resolution.medianGpuWallMs / baselineGpuMs, 2);
+    }
+    latestResult = {
+      schemaVersion: "1.0.0",
+      kind: "micarrsim.webgpu-benchmark",
+      createdAt: new Date().toISOString(),
+      fixture: {
+        arrayId: arrayDefinition.id,
+        microphoneCount: arrayDefinition.microphones.length,
+        sceneId: scene.id,
+        durationSeconds: scene.duration,
+        sampleRate: scene.sampleRate,
+        fftSize: baseExperiment.stft.fftSize,
+        hopSize: baseExperiment.stft.hopSize,
+        selectedBandHz: baseExperiment.frequencyBands.map(({ minHz, maxHz }) => [minHz, maxHz]),
+      },
+      environment: {
+        userAgent: navigator.userAgent,
+        adapterInfo,
+        timingKind: "wall-clock",
+        note: "GPU wall time starts after host packing/writeBuffer calls and waits for queued transfers, two compute passes, and readback; total also includes device/pipeline setup and CPU STFT.",
+      },
+      signalGenerationMs: round(signalGenerationMs),
+      warmupGrid: "16x16",
+      resolutions,
+    };
+    output.textContent = JSON.stringify(latestResult, null, 2);
+    globalThis.__micarrsimBenchmarkResult = latestResult;
+    document.documentElement.dataset.benchmarkStatus = "pass";
+    status.textContent = "benchmark 완료";
+    downloadButton.disabled = false;
+  } catch (error) {
+    document.documentElement.dataset.benchmarkStatus = "fail";
+    document.documentElement.dataset.benchmarkError = error.message;
+    status.textContent = `benchmark 실패: ${error.message}`;
+    output.textContent = error instanceof Error ? error.stack : String(error);
+    throw error;
+  } finally {
+    runButton.disabled = false;
+  }
+}
+
+runButton.addEventListener("click", () => runBenchmark().catch(() => {}));
+downloadButton.addEventListener("click", () => {
+  if (latestResult) downloadJson(latestResult, "micarrsim-webgpu-benchmark.json");
+});
+
+if (query.get("autorun") !== "0") runBenchmark().catch(() => {});
diff --git a/docs/benchmarks/README.md b/docs/benchmarks/README.md
new file mode 100644
index 0000000..003193f
--- /dev/null
+++ b/docs/benchmarks/README.md
@@ -0,0 +1,23 @@
+# WebGPU 기준 실측
+
+2026-08-28에 Chrome 151 headless, Intel Gen9 adapter에서 synthetic 1 kHz fixture를 측정했다. 입력은
+48 kHz · 0.25 s, 192 microphones, FFT 512 / hop 256, 900–1100 Hz band이며 16×16 warm-up 뒤 각
+해상도를 3회 실행한 중앙값이다.
+
+| scan grid | complex MAC | GPU wall median | end-to-end median | working set | 50² 대비 GPU wall |
+| --- | ---: | ---: | ---: | ---: | ---: |
+| 50×50 | 44,160,000 | 79.2 ms | 1,818.1 ms | 9.33 MiB | 1.00× |
+| 100×100 | 176,640,000 | 251.4 ms | 2,102.0 ms | 36.91 MiB | 3.17× |
+| 200×200 | 706,560,000 | 891.9 ms | 4,206.7 ms | 80.47 MiB | 11.26× |
+
+`GPU wall`은 host packing/`writeBuffer()` 호출 직후부터 queue에 들어간 transfer 완료, steering pass,
+DAS pass와 readback까지 기다린 시간이다. `end-to-end`에는 이 host 작업과 장치/pipeline 초기화,
+distance 생성, CPU STFT가 포함된다. 200×200은 64 MiB steering buffer 목표를 지키면서 두 FFT bin을
+한 번에 유지하고 scan point를 21,845개씩 나눈다. 따라서 point chunk가 필요해도 CPU FFT를 다시
+계산하지 않는다. 절대 시간은 브라우저, driver와 GPU에 따라 달라지며 자동 성능 통과/실패 기준으로
+쓰지 않는다.
+
+같은 환경의 자동 검증에서 192채널·50×50 GPU raw map은 CPU raw map 대비 peak-normalized RMS
+오차 `3.008e-5`를 기록했고, 모든 37개 테스트를 통과했다. 상세 원자료는
+`webgpu-intel-gen9-2026-08-28.json`에 있다. 사용자는 `/benchmarks/webgpu-scaling.html`에서 자신의
+장치 결과를 다시 측정해 JSON으로 저장할 수 있다.
diff --git a/docs/benchmarks/webgpu-intel-gen9-2026-08-28.json b/docs/benchmarks/webgpu-intel-gen9-2026-08-28.json
new file mode 100644
index 0000000..548e9ef
--- /dev/null
+++ b/docs/benchmarks/webgpu-intel-gen9-2026-08-28.json
@@ -0,0 +1,90 @@
+{
+  "schemaVersion": "1.0.0",
+  "kind": "micarrsim.webgpu-benchmark",
+  "createdAt": "2026-08-28T02:24:00.354Z",
+  "fixture": {
+    "arrayId": "rectangular-16x12",
+    "microphoneCount": 192,
+    "sceneId": "single-tone",
+    "durationSeconds": 0.25,
+    "sampleRate": 48000,
+    "fftSize": 512,
+    "hopSize": 256,
+    "selectedBandHz": [[900, 1100]]
+  },
+  "environment": {
+    "userAgent": "Mozilla/5.0 (Windows NT 10.0; Win64; x64) AppleWebKit/537.36 (KHTML, like Gecko) HeadlessChrome/151.0.0.0 Safari/537.36",
+    "adapterInfo": {
+      "vendor": "intel",
+      "architecture": "gen-9",
+      "device": "",
+      "description": ""
+    },
+    "timingKind": "wall-clock",
+    "note": "GPU wall time starts after host packing/writeBuffer calls and waits for queued transfers, two compute passes, and readback; total also includes device/pipeline setup and CPU STFT."
+  },
+  "signalGenerationMs": 1200.5,
+  "warmupGrid": "16x16",
+  "resolutions": [
+    {
+      "grid": "50x50",
+      "pointCount": 2500,
+      "runs": 3,
+      "medianTotalMs": 1818.1,
+      "medianStftMs": 1669.8,
+      "medianGpuWallMs": 79.2,
+      "medianBeamformingMs": 162.1,
+      "localizationErrorM": 0.008227,
+      "complexMacCount": 44160000,
+      "estimatedGpuWorkingMiB": 9.33,
+      "binBatchSize": 2,
+      "pointBatchSize": 2500,
+      "samples": [
+        { "totalMs": 1907.7, "stftMs": 1727.7, "gpuWallMs": 82.5, "beamformingMs": 180 },
+        { "totalMs": 1795.7, "stftMs": 1633.6, "gpuWallMs": 79.2, "beamformingMs": 162.1 },
+        { "totalMs": 1818.1, "stftMs": 1669.8, "gpuWallMs": 79.1, "beamformingMs": 148.3 }
+      ],
+      "gpuWallScalingVs50": 1
+    },
+    {
+      "grid": "100x100",
+      "pointCount": 10000,
+      "runs": 3,
+      "medianTotalMs": 2102,
+      "medianStftMs": 1681.1,
+      "medianGpuWallMs": 251.4,
+      "medianBeamformingMs": 537.8,
+      "localizationErrorM": 0.001515,
+      "complexMacCount": 176640000,
+      "estimatedGpuWorkingMiB": 36.91,
+      "binBatchSize": 2,
+      "pointBatchSize": 10000,
+      "samples": [
+        { "totalMs": 2102, "stftMs": 1564.2, "gpuWallMs": 251.4, "beamformingMs": 537.8 },
+        { "totalMs": 2220.1, "stftMs": 1681.1, "gpuWallMs": 254.7, "beamformingMs": 539 },
+        { "totalMs": 2098.8, "stftMs": 1685, "gpuWallMs": 220.4, "beamformingMs": 413.8 }
+      ],
+      "gpuWallScalingVs50": 3.17
+    },
+    {
+      "grid": "200x200",
+      "pointCount": 40000,
+      "runs": 3,
+      "medianTotalMs": 4206.7,
+      "medianStftMs": 2057.5,
+      "medianGpuWallMs": 891.9,
+      "medianBeamformingMs": 2039,
+      "localizationErrorM": 0.002099,
+      "complexMacCount": 706560000,
+      "estimatedGpuWorkingMiB": 80.468,
+      "binBatchSize": 2,
+      "pointBatchSize": 21845,
+      "samples": [
+        { "totalMs": 4206.7, "stftMs": 2057.5, "gpuWallMs": 891.9, "beamformingMs": 2149.2 },
+        { "totalMs": 3388.9, "stftMs": 1686.9, "gpuWallMs": 808.5, "beamformingMs": 1702 },
+        { "totalMs": 6227.8, "stftMs": 4188.8, "gpuWallMs": 939.7, "beamformingMs": 2039 }
+      ],
+      "gpuWallScalingVs50": 11.26
+    }
+  ]
+}
diff --git a/packages/acoustics-core/das.js b/packages/acoustics-core/das.js
index 0b191af..81cd978 100644
--- a/packages/acoustics-core/das.js
+++ b/packages/acoustics-core/das.js
@@ -116,6 +116,37 @@ export function estimateDasWorkload(options) {
   };
 }
 
+export function finalizeDasBands(rawMaps, workload, experiment) {
+  return experiment.frequencyBands.map((band, bandIndex) => {
+    const rawMap = rawMaps[bandIndex];
+    for (let pointIndex = 0; pointIndex < workload.pointCount; pointIndex += 1) {
+      rawMap[pointIndex] /= workload.bandBins[bandIndex].length;
+    }
+    const rawPeak = findPeak(rawMap);
+    const displayMap = Float32Array.from(rawMap);
+    let displayScale = "power";
+    if (experiment.beamforming.normalization === "peak" && rawPeak.value > 0) {
+      displayScale = "normalized";
+      for (let pointIndex = 0; pointIndex < workload.pointCount; pointIndex += 1) {
+        displayMap[pointIndex] /= rawPeak.value;
+      }
+    }
+    return {
+      id: band.id,
+      color: band.color,
+      bins: workload.bandBins[bandIndex],
+      rawMap,
+      displayMap,
+      displayScale,
+      peak: {
+        index: rawPeak.index,
+        position: scanIndexToPosition(rawPeak.index, experiment.scan),
+        rawPower: rawPeak.value,
+      },
+    };
+  });
+}
+
 export async function computeDasMap(options) {
   const {
     channels,
@@ -229,34 +260,7 @@ export async function computeDasMap(options) {
     }
   }
 
-  const bands = experiment.frequencyBands.map((band, bandIndex) => {
-    const rawMap = rawMaps[bandIndex];
-    for (let pointIndex = 0; pointIndex < workload.pointCount; pointIndex += 1) {
-      rawMap[pointIndex] /= workload.bandBins[bandIndex].length;
-    }
-    const rawPeak = findPeak(rawMap);
-    const displayMap = Float32Array.from(rawMap);
-    let displayScale = "power";
-    if (experiment.beamforming.normalization === "peak" && rawPeak.value > 0) {
-      displayScale = "normalized";
-      for (let pointIndex = 0; pointIndex < workload.pointCount; pointIndex += 1) {
-        displayMap[pointIndex] /= rawPeak.value;
-      }
-    }
-    return {
-      id: band.id,
-      color: band.color,
-      bins: workload.bandBins[bandIndex],
-      rawMap,
-      displayMap,
-      displayScale,
-      peak: {
-        index: rawPeak.index,
-        position: scanIndexToPosition(rawPeak.index, experiment.scan),
-        rawPower: rawPeak.value,
-      },
-    };
-  });
+  const bands = finalizeDasBands(rawMaps, workload, experiment);
 
   return {
     backend: "cpu",
diff --git a/packages/webgpu-kernels/README.md b/packages/webgpu-kernels/README.md
index 8b70d10..bbe7ba1 100644
--- a/packages/webgpu-kernels/README.md
+++ b/packages/webgpu-kernels/README.md
@@ -1,3 +1,41 @@
 # WebGPU Kernels
 
-WGSL steering vector 및 DAS 연산과 GPU buffer layout, benchmark 코드를 둡니다.
+CPU STFT 결과를 받아 conventional delay-and-sum(DAS)을 WebGPU로 실행합니다. `backend: "auto"`는
+WebGPU unavailable, device limit, validation/OOM, device-loss 같은 운용 오류에서만 CPU 기준 구현으로
+한 번 fallback합니다. `backend: "webgpu"`는 오류를 숨기지 않으며, 취소와 잘못된 입력도 fallback하지
+않습니다.
+
+## 계산 구조
+
+`das.wgsl.js`는 두 compute pass를 제공합니다.
+
+1. `compute_steering`: `(bin, point, microphone)`별 보정된 complex steering vector를 한 번 계산
+2. `compute_das`: 모든 STFT frame에서 steering을 재사용해 frame 평균 power 계산
+
+평면 scan 순서, phase 부호, 감도/위상 보정, complex multiply, bin 평균과 peak normalization은 CPU
+기준 구현과 동일합니다. GPU 결과를 읽은 뒤 band 조합과 정규화는 `finalizeDasBands()`를 함께
+사용하므로 raw map과 display map의 의미도 같습니다.
+
+## 장치 한도와 메모리
+
+실행기는 하드코딩한 GPU 크기를 믿지 않고 `device.limits`를 조회합니다. steering buffer는 기본
+64 MiB, 전체 활성 GPU working set은 기본 256 MiB 이하가 되도록 분할합니다. CPU FFT 재계산을
+피하기 위해 가능한 bin batch를 유지하고 steering이 커지면 point를 먼저 chunk합니다. output storage와
+`MAP_READ` staging buffer는 WebGPU usage 규칙에 따라 분리합니다.
+
+장치 초기화, WGSL `getCompilationInfo()`, async pipeline 생성, validation/OOM/internal error scope,
+`device.lost`, `uncapturederror`, readback 실패를 각각 처리합니다. 관련 기준은 최신
+[WebGPU specification](https://www.w3.org/TR/webgpu/)과
+[WGSL floating-point accuracy](https://www.w3.org/TR/WGSL/#floating-point-accuracy)를 따릅니다.
+
+## 검증과 benchmark
+
+- `/tests/`: CPU/GPU 교차검증 또는 WebGPU가 없는 환경의 실제 CPU fallback
+- `/benchmarks/webgpu-scaling.html`: 192채널 fixture를 예열한 뒤 50²/100²/200² 각각 3회 측정하고
+  JSON 저장
+- `docs/benchmarks/webgpu-intel-gen9-2026-08-28.json`: Intel Gen9에서 기록한 기준 실측
+
+실측의 `gpuWallMs`는 이미 queue에 넣은 input transfer의 완료 대기 + 두 compute pass + readback
+wall-clock입니다. host packing과 `writeBuffer()` 호출은 그 타이머 직전이며, `totalMs`에는 이 작업과
+장치/pipeline 초기화, CPU STFT도 포함됩니다. GPU timestamp query는 필수 기능이 아니므로 지원되지
+않는 장치도 동일한 계산 경로를 쓸 수 있습니다.
diff --git a/packages/webgpu-kernels/compute.js b/packages/webgpu-kernels/compute.js
new file mode 100644
index 0000000..a7f1080
--- /dev/null
+++ b/packages/webgpu-kernels/compute.js
@@ -0,0 +1,58 @@
+import { computeDasMap } from "../acoustics-core/das.js";
+import { computeDasWebGpu } from "./das.js";
+import { isRecoverableWebGpuError } from "./device.js";
+
+function createAbortError(signal) {
+  if (signal?.reason instanceof Error && signal.reason.name === "AbortError") return signal.reason;
+  const error = new Error("DAS 계산이 취소되었습니다.");
+  error.name = "AbortError";
+  return error;
+}
+
+function throwIfAborted(signal) {
+  if (signal?.aborted) throw createAbortError(signal);
+}
+
+/**
+ * Execute the requested DAS backend. `auto` falls back only for typed WebGPU
+ * operational errors; invalid inputs and AbortError remain visible to callers.
+ */
+export async function computeDasWithBackend(options, dependencies = {}) {
+  const computeGpu = dependencies.computeGpu ?? computeDasWebGpu;
+  const computeCpu = dependencies.computeCpu ?? computeDasMap;
+  const requestedBackend = options?.experiment?.beamforming?.backend ?? "auto";
+
+  if (requestedBackend === "cpu") return computeCpu(options);
+  if (requestedBackend === "webgpu") return computeGpu(options);
+  if (requestedBackend !== "auto") {
+    throw new TypeError(`지원하지 않는 DAS backend입니다: ${requestedBackend}`);
+  }
+
+  try {
+    return await computeGpu(options);
+  } catch (gpuError) {
+    if (!isRecoverableWebGpuError(gpuError)) throw gpuError;
+    throwIfAborted(options.signal);
+    options.onFallback?.({ from: "webgpu", to: "cpu", error: gpuError });
+    throwIfAborted(options.signal);
+    try {
+      const cpuResult = await computeCpu(options);
+      return {
+        ...cpuResult,
+        requestedBackend: "auto",
+        fallback: {
+          from: "webgpu",
+          code: gpuError.code,
+          message: gpuError.message,
+        },
+      };
+    } catch (cpuError) {
+      if (cpuError?.name === "AbortError") throw cpuError;
+      throwIfAborted(options.signal);
+      throw new AggregateError(
+        [gpuError, cpuError],
+        `WebGPU 실행 실패 후 CPU fallback도 실패했습니다: ${cpuError.message}`,
+      );
+    }
+  }
+}
diff --git a/packages/webgpu-kernels/das.js b/packages/webgpu-kernels/das.js
new file mode 100644
index 0000000..74fea32
--- /dev/null
+++ b/packages/webgpu-kernels/das.js
@@ -0,0 +1,656 @@
+import {
+  createDistanceMatrix,
+  estimateDasWorkload,
+  finalizeDasBands,
+} from "../acoustics-core/das.js";
+import { extractStftBins } from "../acoustics-core/stft.js";
+import {
+  DAS_PARAMS_BYTE_LENGTH,
+  DAS_WGSL,
+  WORKGROUP_SIZE,
+  packDasParams,
+  packMicrophoneCalibration,
+} from "./das.wgsl.js";
+import {
+  WebGpuError,
+  WebGpuExecutionError,
+  WebGpuLimitError,
+  WebGpuUnavailableError,
+  assertWebGpuContextHealthy,
+  closeWebGpuContext,
+  createCheckedComputePipeline,
+  createWebGpuContext,
+  withWebGpuErrorScopes,
+} from "./device.js";
+
+export {
+  WebGpuError,
+  WebGpuExecutionError,
+  WebGpuLimitError,
+  WebGpuUnavailableError,
+} from "./device.js";
+
+const MEBIBYTE = 1024 ** 2;
+const DEFAULT_MAXIMUM_GPU_WORKING_BYTES = 256 * MEBIBYTE;
+const DEFAULT_MAXIMUM_HOST_WORKING_BYTES = 512 * MEBIBYTE;
+const DEFAULT_MAXIMUM_STEERING_BUFFER_BYTES = 64 * MEBIBYTE;
+export const DEFAULT_MAXIMUM_MACS_PER_INVOCATION = 250_000;
+
+function now() {
+  return globalThis.performance?.now?.() ?? Date.now();
+}
+
+function throwIfAborted(signal) {
+  if (!signal?.aborted) return;
+  const error = new Error("WebGPU DAS 계산이 취소되었습니다.");
+  error.name = "AbortError";
+  throw error;
+}
+
+function yieldToEventLoop() {
+  return new Promise((resolve) => setTimeout(resolve, 0));
+}
+
+function requirePositiveInteger(value, label) {
+  if (!Number.isSafeInteger(value) || value < 1) {
+    throw new WebGpuLimitError(`${label} 값이 유효하지 않습니다: ${value}`);
+  }
+  return value;
+}
+
+function limitValue(limits, name, fallback) {
+  return requirePositiveInteger(limits?.[name] ?? fallback, `device.limits.${name}`);
+}
+
+function minimumBufferSize(byteLength) {
+  return Math.max(4, Math.ceil(byteLength / 4) * 4);
+}
+
+export function packStftBatch(stft, bins) {
+  if (!stft || !Number.isInteger(stft.frameCount) || !Number.isInteger(stft.microphoneCount)) {
+    throw new TypeError("유효한 STFT 결과가 필요합니다.");
+  }
+  if (!Array.isArray(bins) || bins.length === 0) throw new TypeError("하나 이상의 FFT bin이 필요합니다.");
+  const valuesPerBin = stft.frameCount * stft.microphoneCount;
+  const packed = new Float32Array(bins.length * valuesPerBin * 2);
+  bins.forEach((bin, binIndex) => {
+    const spectrum = stft.spectra.get(bin);
+    if (!spectrum || spectrum.real.length !== valuesPerBin || spectrum.imaginary.length !== valuesPerBin) {
+      throw new Error(`FFT bin ${bin}의 spectrum layout이 올바르지 않습니다.`);
+    }
+    const binOffset = binIndex * valuesPerBin * 2;
+    for (let index = 0; index < valuesPerBin; index += 1) {
+      packed[binOffset + index * 2] = spectrum.real[index];
+      packed[binOffset + index * 2 + 1] = spectrum.imaginary[index];
+    }
+  });
+  return packed;
+}
+
+function pointCapacityForBinBatch({
+  workload,
+  binBatchSize,
+  storageBindingLimit,
+  bufferLimit,
+  maximumSteeringBufferBytes,
+  maximumGpuWorkingBytes,
+  maximumHostWorkingBytes,
+  maximumPointDispatch,
+}) {
+  const microphoneCount = workload.microphoneCount;
+  const distanceBytesPerPoint = microphoneCount * Float32Array.BYTES_PER_ELEMENT;
+  const steeringBytesPerPoint =
+    binBatchSize * microphoneCount * Float32Array.BYTES_PER_ELEMENT * 2;
+  const outputBytesPerPoint = binBatchSize * Float32Array.BYTES_PER_ELEMENT;
+  const spectrumBytes = binBatchSize * workload.spectrumBytesPerBin;
+  const binIndexBytes = binBatchSize * Uint32Array.BYTES_PER_ELEMENT;
+  const calibrationBytes = microphoneCount * Float32Array.BYTES_PER_ELEMENT * 2;
+  const storageLimit = Math.min(storageBindingLimit, bufferLimit);
+  const steeringLimit = Math.min(storageLimit, maximumSteeringBufferBytes);
+
+  const gpuFixedBytes = DAS_PARAMS_BYTE_LENGTH + calibrationBytes + spectrumBytes + binIndexBytes;
+  const gpuBytesPerPoint = distanceBytesPerPoint + steeringBytesPerPoint + outputBytesPerPoint * 2;
+  const hostFixedBytes = workload.baseWorkingBytes + spectrumBytes * 2 + binIndexBytes;
+  const hostBytesPerPoint = distanceBytesPerPoint + outputBytesPerPoint;
+
+  return Math.max(
+    0,
+    Math.min(
+      workload.pointCount,
+      maximumPointDispatch,
+      Math.floor(storageLimit / distanceBytesPerPoint),
+      Math.floor(steeringLimit / steeringBytesPerPoint),
+      Math.floor(storageLimit / outputBytesPerPoint),
+      Math.floor(bufferLimit / outputBytesPerPoint),
+      Math.floor((maximumGpuWorkingBytes - gpuFixedBytes) / gpuBytesPerPoint),
+      Math.floor((maximumHostWorkingBytes - hostFixedBytes) / hostBytesPerPoint),
+    ),
+  );
+}
+
+/**
+ * Plan bin and point chunks against the limits of the device that will execute
+ * them. The largest safe spectrum/bin batch is retained when point chunking can
+ * satisfy steering memory; this avoids repeating the CPU FFT for every GPU bin
+ * chunk on wide bands.
+ */
+export function estimateWebGpuBatchPlan(options) {
+  const {
+    workload,
+    limits,
+    maximumGpuWorkingBytes = DEFAULT_MAXIMUM_GPU_WORKING_BYTES,
+    maximumHostWorkingBytes = DEFAULT_MAXIMUM_HOST_WORKING_BYTES,
+    maximumSteeringBufferBytes = DEFAULT_MAXIMUM_STEERING_BUFFER_BYTES,
+    maximumSpectrumBatchBins = Number.MAX_SAFE_INTEGER,
+    maximumMacsPerInvocation = DEFAULT_MAXIMUM_MACS_PER_INVOCATION,
+  } = options;
+  if (!workload || !Array.isArray(workload.selectedBins) || workload.selectedBins.length === 0) {
+    throw new WebGpuLimitError("WebGPU DAS workload에 FFT bin이 없습니다.");
+  }
+  requirePositiveInteger(workload.microphoneCount, "microphoneCount");
+  requirePositiveInteger(workload.pointCount, "pointCount");
+  requirePositiveInteger(workload.frameCount, "frameCount");
+  requirePositiveInteger(workload.spectrumBytesPerBin, "spectrumBytesPerBin");
+  requirePositiveInteger(maximumGpuWorkingBytes, "maximumGpuWorkingBytes");
+  requirePositiveInteger(maximumHostWorkingBytes, "maximumHostWorkingBytes");
+  requirePositiveInteger(maximumSteeringBufferBytes, "maximumSteeringBufferBytes");
+  requirePositiveInteger(maximumMacsPerInvocation, "maximumMacsPerInvocation");
+
+  const macsPerInvocation = workload.frameCount * workload.microphoneCount;
+  if (!Number.isSafeInteger(macsPerInvocation) || macsPerInvocation > maximumMacsPerInvocation) {
+    throw new WebGpuLimitError(
+      `DAS invocation당 직렬 누산 ${macsPerInvocation.toLocaleString()}회가 안전 한도 ` +
+        `${maximumMacsPerInvocation.toLocaleString()}회를 넘습니다. duration을 줄이세요.`,
+    );
+  }
+
+  const storageBindingLimit = limitValue(limits, "maxStorageBufferBindingSize", 128 * MEBIBYTE);
+  const bufferLimit = limitValue(limits, "maxBufferSize", 256 * MEBIBYTE);
+  const maximumDispatch = limitValue(limits, "maxComputeWorkgroupsPerDimension", 65_535);
+  const maximumInvocations = limitValue(limits, "maxComputeInvocationsPerWorkgroup", 256);
+  const maximumWorkgroupX = limitValue(limits, "maxComputeWorkgroupSizeX", 256);
+  const maximumStorageBuffers = limitValue(limits, "maxStorageBuffersPerShaderStage", 8);
+  const maximumUniformBytes = limitValue(limits, "maxUniformBufferBindingSize", 65_536);
+  const storageLimit = Math.min(storageBindingLimit, bufferLimit);
+
+  if (maximumInvocations < WORKGROUP_SIZE || maximumWorkgroupX < WORKGROUP_SIZE) {
+    throw new WebGpuLimitError(
+      `DAS workgroup ${WORKGROUP_SIZE}개 invocation을 지원하지 않는 WebGPU 장치입니다.`,
+    );
+  }
+  if (maximumStorageBuffers < 6) {
+    throw new WebGpuLimitError(`DAS에 storage buffer 6개가 필요하지만 장치 한도는 ${maximumStorageBuffers}개입니다.`);
+  }
+  if (maximumUniformBytes < DAS_PARAMS_BYTE_LENGTH) {
+    throw new WebGpuLimitError(`DAS uniform ${DAS_PARAMS_BYTE_LENGTH} bytes가 장치 한도를 넘습니다.`);
+  }
+  const calibrationBytes = workload.microphoneCount * Float32Array.BYTES_PER_ELEMENT * 2;
+  if (calibrationBytes > storageLimit) {
+    throw new WebGpuLimitError(
+      `마이크 calibration buffer ${calibrationBytes.toLocaleString()} bytes가 storage binding 한도를 넘습니다.`,
+    );
+  }
+  const distanceBytesPerPoint = workload.microphoneCount * Float32Array.BYTES_PER_ELEMENT;
+  if (distanceBytesPerPoint > storageLimit) {
+    throw new WebGpuLimitError("scan point 한 개의 distance data도 storage binding에 넣을 수 없습니다.");
+  }
+  if (workload.spectrumBytesPerBin > storageLimit) {
+    throw new WebGpuLimitError(
+      `FFT bin 한 개의 spectrum ${workload.spectrumBytesPerBin.toLocaleString()} bytes가 storage binding 한도를 넘습니다.`,
+    );
+  }
+
+  const maximumBinsBySpectrum = Math.floor(storageLimit / workload.spectrumBytesPerBin);
+  const maximumBinsByDispatch = maximumDispatch;
+  const requestedBinLimit = Math.max(1, Math.floor(maximumSpectrumBatchBins));
+  const largestCandidate = Math.min(
+    workload.selectedBins.length,
+    maximumBinsBySpectrum,
+    maximumBinsByDispatch,
+    requestedBinLimit,
+  );
+  if (largestCandidate < 1) throw new WebGpuLimitError("FFT bin 한 개도 WebGPU batch에 넣을 수 없습니다.");
+
+  let binBatchSize = largestCandidate;
+  let pointBatchSize = 0;
+  for (; binBatchSize >= 1; binBatchSize -= 1) {
+    pointBatchSize = pointCapacityForBinBatch({
+      workload,
+      binBatchSize,
+      storageBindingLimit,
+      bufferLimit,
+      maximumSteeringBufferBytes,
+      maximumGpuWorkingBytes,
+      maximumHostWorkingBytes,
+      maximumPointDispatch: maximumDispatch,
+    });
+    if (pointBatchSize >= 1) break;
+  }
+  if (pointBatchSize < 1) {
+    throw new WebGpuLimitError(
+      "현재 장치 및 애플리케이션 메모리 한도에서는 scan point 한 개도 처리할 수 없습니다.",
+    );
+  }
+
+  const spectrumBytes = binBatchSize * workload.spectrumBytesPerBin;
+  const distanceBytes = pointBatchSize * distanceBytesPerPoint;
+  const steeringBytes =
+    binBatchSize * pointBatchSize * workload.microphoneCount * Float32Array.BYTES_PER_ELEMENT * 2;
+  const outputBytes = binBatchSize * pointBatchSize * Float32Array.BYTES_PER_ELEMENT;
+  const calibrationAndParamsBytes = calibrationBytes + DAS_PARAMS_BYTE_LENGTH;
+  const estimatedGpuWorkingBytes =
+    calibrationAndParamsBytes + spectrumBytes + binBatchSize * 4 + distanceBytes + steeringBytes + outputBytes * 2;
+
+  return {
+    binBatchSize,
+    pointBatchSize,
+    binBatchCount: Math.ceil(workload.selectedBins.length / binBatchSize),
+    pointBatchCount: Math.ceil(workload.pointCount / pointBatchSize),
+    macsPerInvocation,
+    steeringBytes,
+    estimatedGpuWorkingBytes,
+    maximumGpuWorkingBytes,
+    maximumHostWorkingBytes,
+  };
+}
+
+export function estimateWebGpuBatchSize(options) {
+  return estimateWebGpuBatchPlan(options).binBatchSize;
+}
+
+function gpuGlobals() {
+  const usage = globalThis.GPUBufferUsage;
+  const mapMode = globalThis.GPUMapMode;
+  if (!usage || !mapMode) {
+    throw new WebGpuUnavailableError("WebGPU buffer 상수를 사용할 수 없습니다.");
+  }
+  return { usage, mapMode };
+}
+
+function createBuffer(device, { label, byteLength, usage, data }) {
+  const buffer = device.createBuffer({
+    label,
+    size: minimumBufferSize(byteLength),
+    usage,
+  });
+  if (data && data.byteLength > 0) {
+    device.queue.writeBuffer(buffer, 0, data.buffer, data.byteOffset, data.byteLength);
+  }
+  return buffer;
+}
+
+function destroyBuffers(buffers) {
+  for (const buffer of buffers) {
+    try {
+      buffer?.destroy();
+    } catch {
+      // A lost device can invalidate resources before deterministic cleanup.
+    }
+  }
+}
+
+async function readOutputBuffer(context, readbackBuffer, byteLength, mapMode) {
+  const mappedOutcome = readbackBuffer.mapAsync(mapMode.READ).then(
+    () => ({ kind: "mapped" }),
+    (error) => ({ kind: "error", error }),
+  );
+  const lostOutcome = context.lostPromise.then((info) => ({ kind: "lost", info }));
+  const outcome = await Promise.race([mappedOutcome, lostOutcome]);
+  if (outcome.kind === "lost") {
+    throw new WebGpuExecutionError(`WebGPU 장치가 계산 중 손실되었습니다: ${outcome.info.message || "원인 미상"}`);
+  }
+  if (outcome.kind === "error") {
+    throw new WebGpuExecutionError("WebGPU 결과 buffer를 읽지 못했습니다.", { cause: outcome.error });
+  }
+  const copy = readbackBuffer.getMappedRange(0, byteLength).slice(0);
+  readbackBuffer.unmap();
+  return new Float32Array(copy);
+}
+
+async function executePointBatch(options) {
+  const {
+    context,
+    pipelines,
+    buffers,
+    pointCount,
+    binCount,
+    microphoneCount,
+    outputByteLength,
+    signal,
+    usage,
+    mapMode,
+  } = options;
+  throwIfAborted(signal);
+  assertWebGpuContextHealthy(context);
+  return withWebGpuErrorScopes(context.device, async () => {
+    const { device } = context;
+    const steeringBindGroup = device.createBindGroup({
+      label: "DAS steering bind group",
+      layout: pipelines.steering.getBindGroupLayout(0),
+      entries: [
+        { binding: 0, resource: { buffer: buffers.params } },
+        { binding: 1, resource: { buffer: buffers.bins } },
+        { binding: 3, resource: { buffer: buffers.distances } },
+        { binding: 4, resource: { buffer: buffers.calibration } },
+        { binding: 5, resource: { buffer: buffers.steering } },
+      ],
+    });
+    const dasBindGroup = device.createBindGroup({
+      label: "DAS beamforming bind group",
+      layout: pipelines.das.getBindGroupLayout(0),
+      entries: [
+        { binding: 0, resource: { buffer: buffers.params } },
+        { binding: 2, resource: { buffer: buffers.spectrum } },
+        { binding: 5, resource: { buffer: buffers.steering } },
+        { binding: 6, resource: { buffer: buffers.output } },
+      ],
+    });
+    const encoder = device.createCommandEncoder({ label: "DAS command encoder" });
+    const pass = encoder.beginComputePass({ label: "DAS steering and beamforming pass" });
+    pass.setPipeline(pipelines.steering);
+    pass.setBindGroup(0, steeringBindGroup);
+    pass.dispatchWorkgroups(Math.ceil(microphoneCount / WORKGROUP_SIZE), pointCount, binCount);
+    pass.setPipeline(pipelines.das);
+    pass.setBindGroup(0, dasBindGroup);
+    pass.dispatchWorkgroups(Math.ceil(pointCount / WORKGROUP_SIZE), binCount, 1);
+    pass.end();
+    encoder.copyBufferToBuffer(buffers.output, 0, buffers.readback, 0, outputByteLength);
+    device.queue.submit([encoder.finish()]);
+    const output = await readOutputBuffer(context, buffers.readback, outputByteLength, mapMode);
+    throwIfAborted(signal);
+    assertWebGpuContextHealthy(context);
+    return output;
+  });
+}
+
+function createBandIndexLookup(workload) {
+  const lookup = new Map();
+  workload.bandBins.forEach((bins, bandIndex) => {
+    bins.forEach((bin) => {
+      if (!lookup.has(bin)) lookup.set(bin, []);
+      lookup.get(bin).push(bandIndex);
+    });
+  });
+  return lookup;
+}
+
+export async function computeDasWebGpu(options) {
+  const {
+    channels,
+    sampleRate,
+    arrayDefinition,
+    scene,
+    experiment,
+    signal,
+    onProgress,
+    maximumGpuWorkingBytes = DEFAULT_MAXIMUM_GPU_WORKING_BYTES,
+    maximumHostWorkingBytes = DEFAULT_MAXIMUM_HOST_WORKING_BYTES,
+    maximumSteeringBufferBytes = DEFAULT_MAXIMUM_STEERING_BUFFER_BYTES,
+    maximumSpectrumBatchBins = Number.MAX_SAFE_INTEGER,
+    maximumMacsPerInvocation = DEFAULT_MAXIMUM_MACS_PER_INVOCATION,
+    maximumGpuComplexMacs = 20_000_000_000,
+    maximumFftButterflies = 1_000_000_000,
+    stftTransformBatchSize = 32,
+    gpu,
+    powerPreference,
+  } = options;
+  if (!Array.isArray(channels) || channels.length !== arrayDefinition?.microphones?.length) {
+    throw new Error(`채널 수 ${channels?.length ?? 0}와 마이크 수 ${arrayDefinition?.microphones?.length ?? 0}가 다릅니다.`);
+  }
+  const totalStartedAt = now();
+  const workload = estimateDasWorkload({ channels, sampleRate, arrayDefinition, experiment });
+  if (workload.complexMacCount > maximumGpuComplexMacs) {
+    throw new WebGpuLimitError(
+      `WebGPU DAS 예상 복소 누산 ${workload.complexMacCount.toLocaleString()}회가 안전 한도 ` +
+        `${maximumGpuComplexMacs.toLocaleString()}회를 넘습니다.`,
+    );
+  }
+
+  const context = await createWebGpuContext({ gpu, powerPreference });
+  const resources = [];
+  let stftMs = 0;
+  let gpuWallMs = 0;
+  try {
+    const { usage, mapMode } = gpuGlobals();
+    assertWebGpuContextHealthy(context);
+    const plan = estimateWebGpuBatchPlan({
+      workload,
+      limits: context.device.limits,
+      maximumGpuWorkingBytes,
+      maximumHostWorkingBytes,
+      maximumSteeringBufferBytes,
+      maximumSpectrumBatchBins,
+      maximumMacsPerInvocation,
+    });
+    const fftButterflyCount =
+      plan.binBatchCount *
+      workload.frameCount *
+      workload.microphoneCount *
+      (experiment.stft.fftSize / 2) *
+      Math.log2(experiment.stft.fftSize);
+    if (fftButterflyCount > maximumFftButterflies) {
+      throw new WebGpuLimitError(
+        `CPU STFT 예상 butterfly ${Math.ceil(fftButterflyCount).toLocaleString()}회가 안전 한도 ` +
+          `${maximumFftButterflies.toLocaleString()}회를 넘습니다.`,
+      );
+    }
+
+    const pipelines = await withWebGpuErrorScopes(context.device, async () => ({
+      steering: await createCheckedComputePipeline(context.device, {
+        code: DAS_WGSL,
+        entryPoint: "compute_steering",
+        label: "DAS steering pipeline",
+      }),
+      das: await createCheckedComputePipeline(context.device, {
+        code: DAS_WGSL,
+        entryPoint: "compute_das",
+        label: "DAS beamforming pipeline",
+      }),
+    }));
+    const calibration = packMicrophoneCalibration(arrayDefinition.microphones);
+    const calibrationBuffer = await withWebGpuErrorScopes(context.device, async () =>
+      createBuffer(context.device, {
+        label: "DAS microphone calibration",
+        byteLength: calibration.byteLength,
+        usage: usage.STORAGE | usage.COPY_DST,
+        data: calibration,
+      }),
+    );
+    resources.push(calibrationBuffer);
+
+    const distances = createDistanceMatrix(arrayDefinition, experiment.scan);
+    const rawMaps = experiment.frequencyBands.map(() => new Float32Array(workload.pointCount));
+    const bandIndexesByBin = createBandIndexLookup(workload);
+    const totalDispatchUnits = plan.binBatchCount * plan.pointBatchCount;
+    let completedDispatchUnits = 0;
+    let completedBins = 0;
+
+    for (let binStart = 0; binStart < workload.selectedBins.length; binStart += plan.binBatchSize) {
+      throwIfAborted(signal);
+      await yieldToEventLoop();
+      const batchBins = workload.selectedBins.slice(binStart, binStart + plan.binBatchSize);
+      const stftStartedAt = now();
+      const stft = await extractStftBins(channels, {
+        fftSize: experiment.stft.fftSize,
+        hopSize: experiment.stft.hopSize,
+        bins: batchBins,
+        signal,
+        transformBatchSize: stftTransformBatchSize,
+        onProgress: ({ completedTransforms, totalTransforms }) =>
+          onProgress?.({
+            stage: "stft",
+            backend: "webgpu",
+            completedBins,
+            totalBins: workload.selectedBins.length,
+            completedTransforms,
+            totalTransforms,
+            spectrumBatch: Math.floor(binStart / plan.binBatchSize) + 1,
+            spectrumBatchCount: plan.binBatchCount,
+          }),
+      });
+      const packedSpectrum = packStftBatch(stft, batchBins);
+      stftMs += now() - stftStartedAt;
+      const binIndices = Uint32Array.from(batchBins);
+      const binResources = [];
+      try {
+        const { binsBuffer, spectrumBuffer } = await withWebGpuErrorScopes(context.device, async () => ({
+          binsBuffer: createBuffer(context.device, {
+            label: "DAS FFT bin indices",
+            byteLength: binIndices.byteLength,
+            usage: usage.STORAGE | usage.COPY_DST,
+            data: binIndices,
+          }),
+          spectrumBuffer: createBuffer(context.device, {
+            label: "DAS packed STFT spectrum",
+            byteLength: packedSpectrum.byteLength,
+            usage: usage.STORAGE | usage.COPY_DST,
+            data: packedSpectrum,
+          }),
+        }));
+        binResources.push(binsBuffer, spectrumBuffer);
+
+        for (let pointStart = 0; pointStart < workload.pointCount; pointStart += plan.pointBatchSize) {
+          throwIfAborted(signal);
+          const pointCount = Math.min(plan.pointBatchSize, workload.pointCount - pointStart);
+          const distanceStart = pointStart * workload.microphoneCount;
+          const distanceEnd = distanceStart + pointCount * workload.microphoneCount;
+          const pointDistances = distances.subarray(distanceStart, distanceEnd);
+          const params = packDasParams({
+            pointCount,
+            microphoneCount: workload.microphoneCount,
+            frameCount: workload.frameCount,
+            binCount: batchBins.length,
+            sampleRate,
+            fftSize: experiment.stft.fftSize,
+            speedOfSound: scene.speedOfSound,
+          });
+          const steeringByteLength =
+            batchBins.length * pointCount * workload.microphoneCount * Float32Array.BYTES_PER_ELEMENT * 2;
+          const outputByteLength = batchBins.length * pointCount * Float32Array.BYTES_PER_ELEMENT;
+          const pointResources = [];
+          try {
+            const {
+              paramsBuffer,
+              distancesBuffer,
+              steeringBuffer,
+              outputBuffer,
+              readbackBuffer,
+            } = await withWebGpuErrorScopes(context.device, async () => ({
+              paramsBuffer: createBuffer(context.device, {
+                label: "DAS parameters",
+                byteLength: params.byteLength,
+                usage: usage.UNIFORM | usage.COPY_DST,
+                data: new Uint8Array(params),
+              }),
+              distancesBuffer: createBuffer(context.device, {
+                label: "DAS point-to-microphone distances",
+                byteLength: pointDistances.byteLength,
+                usage: usage.STORAGE | usage.COPY_DST,
+                data: pointDistances,
+              }),
+              steeringBuffer: createBuffer(context.device, {
+                label: "DAS steering vectors",
+                byteLength: steeringByteLength,
+                usage: usage.STORAGE,
+              }),
+              outputBuffer: createBuffer(context.device, {
+                label: "DAS output maps",
+                byteLength: outputByteLength,
+                usage: usage.STORAGE | usage.COPY_SRC,
+              }),
+              readbackBuffer: createBuffer(context.device, {
+                label: "DAS readback staging",
+                byteLength: outputByteLength,
+                usage: usage.MAP_READ | usage.COPY_DST,
+              }),
+            }));
+            pointResources.push(paramsBuffer, distancesBuffer, steeringBuffer, outputBuffer, readbackBuffer);
+            const gpuStartedAt = now();
+            const output = await executePointBatch({
+              context,
+              pipelines,
+              buffers: {
+                params: paramsBuffer,
+                bins: binsBuffer,
+                spectrum: spectrumBuffer,
+                distances: distancesBuffer,
+                calibration: calibrationBuffer,
+                steering: steeringBuffer,
+                output: outputBuffer,
+                readback: readbackBuffer,
+              },
+              pointCount,
+              binCount: batchBins.length,
+              microphoneCount: workload.microphoneCount,
+              outputByteLength,
+              signal,
+              usage,
+              mapMode,
+            });
+            gpuWallMs += now() - gpuStartedAt;
+            for (let outputIndex = 0; outputIndex < output.length; outputIndex += 1) {
+              if (!Number.isFinite(output[outputIndex]) || output[outputIndex] < 0) {
+                throw new WebGpuExecutionError(
+                  `WebGPU DAS output[${outputIndex}]이 유효한 power 값이 아닙니다: ${output[outputIndex]}`,
+                );
+              }
+            }
+            batchBins.forEach((bin, localBinIndex) => {
+              const localMapOffset = localBinIndex * pointCount;
+              for (const bandIndex of bandIndexesByBin.get(bin)) {
+                const rawMap = rawMaps[bandIndex];
+                for (let localPointIndex = 0; localPointIndex < pointCount; localPointIndex += 1) {
+                  rawMap[pointStart + localPointIndex] += output[localMapOffset + localPointIndex];
+                }
+              }
+            });
+          } finally {
+            destroyBuffers(pointResources);
+          }
+          completedDispatchUnits += 1;
+          onProgress?.({
+            stage: "beamforming",
+            backend: "webgpu",
+            completedBins:
+              pointStart + pointCount === workload.pointCount
+                ? completedBins + batchBins.length
+                : completedBins,
+            totalBins: workload.selectedBins.length,
+            completedUnits: completedDispatchUnits,
+            totalUnits: totalDispatchUnits,
+          });
+          await yieldToEventLoop();
+        }
+      } finally {
+        destroyBuffers(binResources);
+      }
+      completedBins += batchBins.length;
+    }
+
+    const bands = finalizeDasBands(rawMaps, workload, experiment);
+    const totalMs = now() - totalStartedAt;
+    return {
+      backend: "webgpu",
+      frameCount: workload.frameCount,
+      microphoneCount: workload.microphoneCount,
+      pointCount: workload.pointCount,
+      workload,
+      fftButterflyCount,
+      bands,
+      timingMs: {
+        total: totalMs,
+        stft: stftMs,
+        gpuWall: gpuWallMs,
+        beamforming: totalMs - stftMs,
+      },
+      gpu: {
+        adapterInfo: context.adapterInfo,
+        timingKind: "wall-clock",
+        batchPlan: plan,
+      },
+    };
+  } finally {
+    destroyBuffers(resources);
+    closeWebGpuContext(context);
+  }
+}
diff --git a/packages/webgpu-kernels/das.wgsl.js b/packages/webgpu-kernels/das.wgsl.js
new file mode 100644
index 0000000..aa6ce87
--- /dev/null
+++ b/packages/webgpu-kernels/das.wgsl.js
@@ -0,0 +1,414 @@
+/**
+ * Two-pass frequency-domain delay-and-sum beamforming kernel.
+ *
+ * Dispatch shapes:
+ *   compute_steering = (ceil(microphoneCount / WORKGROUP_SIZE), pointCount, binCount)
+ *   compute_das      = (ceil(pointCount / WORKGROUP_SIZE), binCount, 1)
+ *
+ * The first pass computes trigonometry once per (bin, point, microphone).
+ * The second pass reuses those steering vectors for every STFT frame. Callers
+ * can bound the intermediate buffer by submitting point/bin batches.
+ *
+ * Flat buffer layouts:
+ *   spectrum:    [bin][frame][microphone][real, imaginary]
+ *   distances:   [point][microphone]
+ *   calibration: [microphone][inverse sensitivity gain, phase offset radians]
+ *   steering:    [bin][point][microphone][real, imaginary]
+ *   output:      [bin][point]
+ */
+
+export const WORKGROUP_SIZE = 64;
+export const DAS_WORKGROUP_SIZE = WORKGROUP_SIZE;
+export const MAX_MICROPHONES = 1024;
+
+// WGSL uniform structs are rounded to a 16-byte boundary. The reserved word
+// keeps the JavaScript and shader representations explicitly 32 bytes.
+export const DAS_PARAMS_BYTE_LENGTH = 32;
+
+export const DAS_PARAMS_LAYOUT = Object.freeze({
+  byteLength: DAS_PARAMS_BYTE_LENGTH,
+  pointCount: Object.freeze({ offset: 0, type: "u32" }),
+  microphoneCount: Object.freeze({ offset: 4, type: "u32" }),
+  frameCount: Object.freeze({ offset: 8, type: "u32" }),
+  binCount: Object.freeze({ offset: 12, type: "u32" }),
+  sampleRate: Object.freeze({ offset: 16, type: "f32" }),
+  fftSize: Object.freeze({ offset: 20, type: "u32" }),
+  speedOfSound: Object.freeze({ offset: 24, type: "f32" }),
+  reserved: Object.freeze({ offset: 28, type: "u32" }),
+});
+
+export const DAS_BUFFER_LAYOUT = Object.freeze({
+  params: Object.freeze({ group: 0, binding: 0, access: "uniform", layout: DAS_PARAMS_LAYOUT }),
+  binIndices: Object.freeze({ group: 0, binding: 1, access: "read", element: "u32" }),
+  spectrum: Object.freeze({
+    group: 0,
+    binding: 2,
+    access: "read",
+    element: "f32",
+    order: "bin-frame-microphone-complex",
+  }),
+  distances: Object.freeze({
+    group: 0,
+    binding: 3,
+    access: "read",
+    element: "f32",
+    order: "point-microphone",
+  }),
+  calibration: Object.freeze({
+    group: 0,
+    binding: 4,
+    access: "read",
+    element: "vec2",
+    order: "microphone-(inverseGain-phaseRadians)",
+  }),
+  steering: Object.freeze({
+    group: 0,
+    binding: 5,
+    access: "read_write",
+    element: "vec2",
+    order: "bin-point-microphone-complex",
+  }),
+  output: Object.freeze({
+    group: 0,
+    binding: 6,
+    access: "read_write",
+    element: "f32",
+    order: "bin-point",
+  }),
+});
+
+export const DAS_ENTRY_POINTS = Object.freeze({
+  steering: "compute_steering",
+  beamforming: "compute_das",
+});
+
+function requireUint32(name, value) {
+  if (!Number.isInteger(value) || value < 0 || value > 0xffff_ffff) {
+    throw new RangeError(`${name} must be an unsigned 32-bit integer.`);
+  }
+  return value;
+}
+
+function requireUint32Product(name, values) {
+  let product = 1;
+  for (const value of values) {
+    if (value !== 0 && product > 0xffff_ffff / value) {
+      throw new RangeError(`${name} exceeds the shader's u32 address space.`);
+    }
+    product *= value;
+  }
+  return product;
+}
+
+function requirePositiveF32(name, value) {
+  const rounded = Math.fround(value);
+  if (!(rounded > 0) || !Number.isFinite(rounded)) {
+    throw new RangeError(`${name} must be a finite positive f32 value.`);
+  }
+  return rounded;
+}
+
+function requireFiniteF32(name, value) {
+  const rounded = Math.fround(value);
+  if (!Number.isFinite(rounded)) {
+    throw new RangeError(`${name} must be a finite f32 value.`);
+  }
+  return rounded;
+}
+
+function requireFiniteNumber(name, value) {
+  if (!Number.isFinite(value)) throw new RangeError(`${name} must be finite.`);
+  return value;
+}
+
+/** Pack the binding-0 uniform using the offsets in DAS_PARAMS_LAYOUT. */
+export function packDasParams(options) {
+  const {
+    pointCount,
+    microphoneCount,
+    frameCount,
+    binCount,
+    sampleRate,
+    fftSize,
+    speedOfSound,
+  } = options;
+  const counts = {
+    pointCount: requireUint32("pointCount", pointCount),
+    microphoneCount: requireUint32("microphoneCount", microphoneCount),
+    frameCount: requireUint32("frameCount", frameCount),
+    binCount: requireUint32("binCount", binCount),
+    fftSize: requireUint32("fftSize", fftSize),
+  };
+  if (counts.microphoneCount > MAX_MICROPHONES) {
+    throw new RangeError(`microphoneCount must not exceed ${MAX_MICROPHONES}.`);
+  }
+  if (counts.fftSize === 0) throw new RangeError("fftSize must be greater than zero.");
+
+  requireUint32Product("distance buffer element count", [
+    counts.pointCount,
+    counts.microphoneCount,
+  ]);
+  requireUint32Product("output buffer element count", [counts.binCount, counts.pointCount]);
+  requireUint32Product("steering buffer element count", [
+    counts.binCount,
+    counts.pointCount,
+    counts.microphoneCount,
+  ]);
+  requireUint32Product("spectrum buffer element count", [
+    counts.binCount,
+    counts.frameCount,
+    counts.microphoneCount,
+    2,
+  ]);
+
+  const packedSampleRate = requirePositiveF32("sampleRate", sampleRate);
+  const packedSpeedOfSound = requirePositiveF32("speedOfSound", speedOfSound);
+  const buffer = new ArrayBuffer(DAS_PARAMS_BYTE_LENGTH);
+  const view = new DataView(buffer);
+  view.setUint32(0, counts.pointCount, true);
+  view.setUint32(4, counts.microphoneCount, true);
+  view.setUint32(8, counts.frameCount, true);
+  view.setUint32(12, counts.binCount, true);
+  view.setFloat32(16, packedSampleRate, true);
+  view.setUint32(20, counts.fftSize, true);
+  view.setFloat32(24, packedSpeedOfSound, true);
+  view.setUint32(28, 0, true);
+  return buffer;
+}
+
+/**
+ * Pack already-linear calibration as tightly packed vec2 entries.
+ * phase is the microphone phase offset in radians. The shader subtracts it
+ * from the propagation phase, exactly as the CPU reference does.
+ */
+export function packDasCalibration(calibration) {
+  if (!Array.isArray(calibration)) throw new TypeError("calibration must be an array.");
+  const packed = new Float32Array(calibration.length * 2);
+  for (let index = 0; index < calibration.length; index += 1) {
+    const entry = calibration[index];
+    const gain = requireFiniteF32(`calibration[${index}].gain`, entry?.gain);
+    if (gain < 0) throw new RangeError(`calibration[${index}].gain must not be negative.`);
+    packed[index * 2] = gain;
+    packed[index * 2 + 1] = requireFiniteF32(
+      `calibration[${index}].phase`,
+      entry?.phase,
+    );
+  }
+  return packed;
+}
+
+/** Convert array-definition microphone metadata to the shader calibration layout. */
+export function packMicrophoneCalibration(microphones) {
+  if (!Array.isArray(microphones)) throw new TypeError("microphones must be an array.");
+  const packed = new Float32Array(microphones.length * 2);
+  for (let index = 0; index < microphones.length; index += 1) {
+    const microphone = microphones[index];
+    const sensitivityDb = requireFiniteNumber(
+      `microphones[${index}].sensitivityDb`,
+      microphone?.sensitivityDb ?? 0,
+    );
+    const phaseOffsetDeg = requireFiniteNumber(
+      `microphones[${index}].phaseOffsetDeg`,
+      microphone?.phaseOffsetDeg ?? 0,
+    );
+    packed[index * 2] = requireFiniteF32(
+      `microphones[${index}] inverse sensitivity gain`,
+      10 ** (-sensitivityDb / 20),
+    );
+    packed[index * 2 + 1] = requireFiniteF32(
+      `microphones[${index}] phase offset`,
+      (phaseOffsetDeg * Math.PI) / 180,
+    );
+  }
+  return packed;
+}
+
+export const DAS_WGSL = /* wgsl */ `
+struct DasParams {
+  point_count: u32,
+  microphone_count: u32,
+  frame_count: u32,
+  bin_count: u32,
+  sample_rate: f32,
+  fft_size: u32,
+  speed_of_sound: f32,
+  reserved: u32,
+}
+
+struct U32Buffer {
+  values: array,
+}
+
+struct F32Buffer {
+  values: array,
+}
+
+struct Vec2Buffer {
+  values: array>,
+}
+
+@group(0) @binding(0) var params: DasParams;
+@group(0) @binding(1) var bin_indices: U32Buffer;
+@group(0) @binding(2) var spectrum: F32Buffer;
+@group(0) @binding(3) var distances: F32Buffer;
+@group(0) @binding(4) var calibration: Vec2Buffer;
+@group(0) @binding(5) var steering_vectors: Vec2Buffer;
+@group(0) @binding(6) var output_maps: F32Buffer;
+
+const TAU: f32 = 6.283185307179586;
+const MAX_MICROPHONES: u32 = ${MAX_MICROPHONES}u;
+const MAX_U32: u32 = 0xffffffffu;
+
+fn multiplication_overflows(left: u32, right: u32) -> bool {
+  return left != 0u && right > MAX_U32 / left;
+}
+
+fn wrap_phase(phase: f32) -> f32 {
+  return phase - round(phase / TAU) * TAU;
+}
+
+@compute @workgroup_size(${WORKGROUP_SIZE}, 1, 1)
+fn compute_steering(@builtin(global_invocation_id) invocation: vec3) {
+  let microphone_index = invocation.x;
+  let point_index = invocation.y;
+  let bin_index = invocation.z;
+
+  if (
+    microphone_index >= params.microphone_count ||
+    point_index >= params.point_count ||
+    bin_index >= params.bin_count
+  ) {
+    return;
+  }
+
+  if (
+    params.microphone_count == 0u ||
+    params.microphone_count > MAX_MICROPHONES ||
+    params.fft_size == 0u ||
+    params.sample_rate <= 0.0 ||
+    params.speed_of_sound <= 0.0 ||
+    multiplication_overflows(params.point_count, params.microphone_count) ||
+    multiplication_overflows(params.bin_count, params.point_count)
+  ) {
+    return;
+  }
+
+  let bin_point_count = params.bin_count * params.point_count;
+  if (multiplication_overflows(bin_point_count, params.microphone_count)) {
+    return;
+  }
+
+  let required_distance_values = params.point_count * params.microphone_count;
+  let required_steering_values = bin_point_count * params.microphone_count;
+  if (
+    arrayLength(&bin_indices.values) < params.bin_count ||
+    arrayLength(&calibration.values) < params.microphone_count ||
+    arrayLength(&distances.values) < required_distance_values ||
+    arrayLength(&steering_vectors.values) < required_steering_values
+  ) {
+    return;
+  }
+
+  let bin = f32(bin_indices.values[bin_index]);
+  let angular_wave_number =
+    (TAU * bin * params.sample_rate) /
+    (f32(params.fft_size) * params.speed_of_sound);
+  let distance_index = point_index * params.microphone_count + microphone_index;
+  let steering_index =
+    (bin_index * params.point_count + point_index) * params.microphone_count +
+    microphone_index;
+  let microphone_calibration = calibration.values[microphone_index];
+  let phase = wrap_phase(
+    angular_wave_number * distances.values[distance_index] -
+    microphone_calibration.y
+  );
+  steering_vectors.values[steering_index] =
+    vec2(cos(phase), sin(phase)) * microphone_calibration.x;
+}
+
+@compute @workgroup_size(${WORKGROUP_SIZE}, 1, 1)
+fn compute_das(@builtin(global_invocation_id) invocation: vec3) {
+  let point_index = invocation.x;
+  let bin_index = invocation.y;
+
+  if (point_index >= params.point_count || bin_index >= params.bin_count) {
+    return;
+  }
+
+  if (
+    params.microphone_count == 0u ||
+    params.microphone_count > MAX_MICROPHONES ||
+    params.frame_count == 0u ||
+    multiplication_overflows(params.bin_count, params.point_count) ||
+    multiplication_overflows(params.bin_count, params.frame_count)
+  ) {
+    return;
+  }
+
+  let bin_point_count = params.bin_count * params.point_count;
+  let bin_frame_count = params.bin_count * params.frame_count;
+  if (
+    multiplication_overflows(bin_point_count, params.microphone_count) ||
+    multiplication_overflows(bin_frame_count, params.microphone_count)
+  ) {
+    return;
+  }
+
+  let spectrum_complex_count = bin_frame_count * params.microphone_count;
+  if (multiplication_overflows(spectrum_complex_count, 2u)) {
+    return;
+  }
+
+  let required_output_values = bin_point_count;
+  let required_steering_values = bin_point_count * params.microphone_count;
+  let required_spectrum_values = spectrum_complex_count * 2u;
+  if (
+    arrayLength(&output_maps.values) < required_output_values ||
+    arrayLength(&steering_vectors.values) < required_steering_values ||
+    arrayLength(&spectrum.values) < required_spectrum_values
+  ) {
+    return;
+  }
+
+  let microphone_scale = 1.0 / f32(params.microphone_count);
+  let steering_base =
+    (bin_index * params.point_count + point_index) * params.microphone_count;
+  var accumulated_power = 0.0;
+
+  for (var frame_index = 0u; frame_index < params.frame_count; frame_index += 1u) {
+    var sum = vec2(0.0, 0.0);
+    for (
+      var microphone_index = 0u;
+      microphone_index < params.microphone_count;
+      microphone_index += 1u
+    ) {
+      let complex_index = (
+        (
+          (bin_index * params.frame_count + frame_index) * params.microphone_count +
+          microphone_index
+        ) * 2u
+      );
+      let input = vec2(
+        spectrum.values[complex_index],
+        spectrum.values[complex_index + 1u]
+      );
+      let steering = steering_vectors.values[steering_base + microphone_index];
+
+      // (input.real + i*input.imaginary) *
+      // (steering.real + i*steering.imaginary)
+      sum += vec2(
+        input.x * steering.x - input.y * steering.y,
+        input.x * steering.y + input.y * steering.x
+      );
+    }
+
+    accumulated_power +=
+      dot(sum, sum) * microphone_scale * microphone_scale;
+  }
+
+  let output_index = bin_index * params.point_count + point_index;
+  output_maps.values[output_index] = accumulated_power / f32(params.frame_count);
+}
+`;
+
+export const DAS_SHADER_SOURCE = DAS_WGSL;
diff --git a/packages/webgpu-kernels/device.js b/packages/webgpu-kernels/device.js
new file mode 100644
index 0000000..3fd9fe7
--- /dev/null
+++ b/packages/webgpu-kernels/device.js
@@ -0,0 +1,170 @@
+/** Errors in this module are deliberately typed so `auto` can fall back to CPU
+ * without hiding invalid experiment input or an explicit user cancellation. */
+export class WebGpuError extends Error {
+  constructor(message, options = {}) {
+    super(message, options.cause ? { cause: options.cause } : undefined);
+    this.name = new.target.name;
+    this.code = options.code ?? "WEBGPU_ERROR";
+  }
+}
+
+export class WebGpuUnavailableError extends WebGpuError {
+  constructor(message = "이 브라우저에서 WebGPU를 사용할 수 없습니다.", options = {}) {
+    super(message, { ...options, code: "WEBGPU_UNAVAILABLE" });
+  }
+}
+
+export class WebGpuLimitError extends WebGpuError {
+  constructor(message, options = {}) {
+    super(message, { ...options, code: "WEBGPU_LIMIT" });
+  }
+}
+
+export class WebGpuExecutionError extends WebGpuError {
+  constructor(message, options = {}) {
+    super(message, { ...options, code: "WEBGPU_EXECUTION" });
+  }
+}
+
+export function isRecoverableWebGpuError(error) {
+  return error instanceof WebGpuError;
+}
+
+function describeAdapter(adapter) {
+  const info = adapter?.info;
+  if (!info) return null;
+  return {
+    vendor: info.vendor ?? "",
+    architecture: info.architecture ?? "",
+    device: info.device ?? "",
+    description: info.description ?? "",
+  };
+}
+
+function normalizeRequestError(error, message) {
+  if (error instanceof WebGpuError) return error;
+  return new WebGpuUnavailableError(message, { cause: error });
+}
+
+/**
+ * Request one short-lived WebGPU context. A new call always requests a new
+ * adapter so a device-loss retry never reuses an invalid adapter/device pair.
+ */
+export async function createWebGpuContext(options = {}) {
+  const gpu = options.gpu ?? globalThis.navigator?.gpu;
+  if (!gpu) throw new WebGpuUnavailableError();
+
+  let adapter;
+  try {
+    adapter = await gpu.requestAdapter({ powerPreference: options.powerPreference ?? "high-performance" });
+  } catch (error) {
+    throw normalizeRequestError(error, "WebGPU 어댑터 요청에 실패했습니다.");
+  }
+  if (!adapter) throw new WebGpuUnavailableError("사용 가능한 WebGPU 어댑터가 없습니다.");
+
+  let device;
+  try {
+    device = await adapter.requestDevice();
+  } catch (error) {
+    throw normalizeRequestError(error, "WebGPU 장치 요청에 실패했습니다.");
+  }
+
+  const state = {
+    lostInfo: null,
+    uncapturedError: null,
+    closedByApplication: false,
+  };
+  const onUncapturedError = (event) => {
+    state.uncapturedError = event.error ?? new Error("알 수 없는 WebGPU 오류");
+  };
+  device.addEventListener?.("uncapturederror", onUncapturedError);
+  const lostPromise = device.lost.then((info) => {
+    state.lostInfo = info;
+    return info;
+  });
+
+  return {
+    adapter,
+    device,
+    adapterInfo: describeAdapter(adapter),
+    state,
+    lostPromise,
+    onUncapturedError,
+  };
+}
+
+export function assertWebGpuContextHealthy(context) {
+  const { state } = context;
+  if (state.uncapturedError) {
+    throw new WebGpuExecutionError(`처리되지 않은 WebGPU 오류: ${state.uncapturedError.message}`, {
+      cause: state.uncapturedError,
+    });
+  }
+  if (state.lostInfo && !(state.closedByApplication && state.lostInfo.reason === "destroyed")) {
+    throw new WebGpuExecutionError(`WebGPU 장치가 손실되었습니다: ${state.lostInfo.message || "원인 미상"}`);
+  }
+}
+
+export function closeWebGpuContext(context) {
+  if (!context?.device) return;
+  context.state.closedByApplication = true;
+  context.device.removeEventListener?.("uncapturederror", context.onUncapturedError);
+  context.device.destroy();
+}
+
+export async function createCheckedComputePipeline(device, { code, entryPoint, label }) {
+  let shaderModule;
+  try {
+    shaderModule = device.createShaderModule({ code, label: `${label} shader` });
+    const compilationInfo = await shaderModule.getCompilationInfo();
+    const compilationErrors = compilationInfo.messages.filter((message) => message.type === "error");
+    if (compilationErrors.length > 0) {
+      const details = compilationErrors
+        .slice(0, 4)
+        .map((message) => `${message.lineNum}:${message.linePos} ${message.message}`)
+        .join("\n");
+      throw new WebGpuExecutionError(`WGSL 컴파일에 실패했습니다.\n${details}`);
+    }
+    return await device.createComputePipelineAsync({
+      label,
+      layout: "auto",
+      compute: { module: shaderModule, entryPoint },
+    });
+  } catch (error) {
+    if (error instanceof WebGpuError) throw error;
+    throw new WebGpuExecutionError("WebGPU compute pipeline 생성에 실패했습니다.", { cause: error });
+  }
+}
+
+/** Capture validation/OOM/internal errors around predictable GPU work. */
+export async function withWebGpuErrorScopes(device, callback) {
+  const filters = ["internal", "out-of-memory", "validation"];
+  for (const filter of filters) device.pushErrorScope(filter);
+  let result;
+  let callbackError;
+  try {
+    result = await callback();
+  } catch (error) {
+    callbackError = error;
+  }
+
+  const captured = [];
+  for (let index = filters.length - 1; index >= 0; index -= 1) {
+    try {
+      const error = await device.popErrorScope();
+      if (error) captured.push(error);
+    } catch (error) {
+      captured.push(error);
+    }
+  }
+  if (callbackError) {
+    if (callbackError instanceof WebGpuError) throw callbackError;
+    throw new WebGpuExecutionError("WebGPU 명령 실행에 실패했습니다.", { cause: callbackError });
+  }
+  if (captured.length > 0) {
+    throw new WebGpuExecutionError(`WebGPU 검증 오류: ${captured.map((error) => error.message).join(" / ")}`, {
+      cause: captured[0],
+    });
+  }
+  return result;
+}
diff --git a/packages/webgpu-kernels/index.js b/packages/webgpu-kernels/index.js
new file mode 100644
index 0000000..d21c1d7
--- /dev/null
+++ b/packages/webgpu-kernels/index.js
@@ -0,0 +1,3 @@
+export * from "./compute.js";
+export * from "./das.js";
+export * from "./das.wgsl.js";
diff --git a/tests/integration/cpu-baseline.test.js b/tests/integration/cpu-baseline.test.js
index 74e1bc0..b1e40ff 100644
--- a/tests/integration/cpu-baseline.test.js
+++ b/tests/integration/cpu-baseline.test.js
@@ -30,5 +30,13 @@ test("192채널 · 50×50 CPU baseline이 ground truth를 찾는다", async () =
   assertEqual(result.pointCount, 2500, "scan point count");
   assert(localizationError < 0.03, `localization error=${localizationError} m`);
   assert(Number.isFinite(result.bands[0].peak.rawPower) && result.bands[0].peak.rawPower > 0, "finite peak power");
+  globalThis.__micarrsimCpuBaseline = {
+    channels: synthesis.channels,
+    sampleRate: synthesis.sampleRate,
+    arrayDefinition,
+    scene,
+    experiment,
+    result,
+  };
   document.documentElement.dataset.cpuBaselineMs = elapsedMs.toFixed(1);
 });
diff --git a/tests/integration/webgpu.test.js b/tests/integration/webgpu.test.js
new file mode 100644
index 0000000..ebdc42e
--- /dev/null
+++ b/tests/integration/webgpu.test.js
@@ -0,0 +1,176 @@
+import { computeDasMap } from "/packages/acoustics-core/das.js";
+import { distance3 } from "/packages/acoustics-core/math.js";
+import { synthesizeMicrophoneSignals } from "/packages/acoustics-core/propagation.js";
+import { computeDasWithBackend } from "/packages/webgpu-kernels/compute.js";
+import { computeDasWebGpu } from "/packages/webgpu-kernels/das.js";
+
+function createArray(columns, rows, width, height) {
+  const microphones = [];
+  for (let row = 0; row < rows; row += 1) {
+    for (let column = 0; column < columns; column += 1) {
+      const index = row * columns + column;
+      microphones.push({
+        id: `gpu-mic-${row}-${column}`,
+        position: [
+          -width / 2 + (column * width) / (columns - 1),
+          -height / 2 + (row * height) / (rows - 1),
+          0,
+        ],
+        sensitivityDb: (index % 3) * 2 - 2,
+        phaseOffsetDeg: (index % 5) * 8 - 16,
+      });
+    }
+  }
+  return { id: "gpu-test-array", aperture: { width, height }, microphones };
+}
+
+function createFixture(backend = "webgpu") {
+  const arrayDefinition = createArray(4, 4, 0.3, 0.3);
+  const scene = {
+    id: "gpu-das-scene",
+    sampleRate: 8_000,
+    duration: 0.064,
+    speedOfSound: 343,
+    seed: 11,
+    mix: { normalization: "none", targetPeak: 0.95 },
+    sources: [{
+      id: "source",
+      name: "source",
+      position: [0.1, 0.1, 1.2],
+      gainDb: -12,
+      startTime: 0,
+      duration: 0.064,
+      signal: { type: "tone", frequency: 1_000, phaseDeg: 0 },
+    }],
+    ambient: [],
+  };
+  const experiment = {
+    stft: { fftSize: 128, hopSize: 64, window: "hann" },
+    scan: { z: 1.2, xMin: -0.4, xMax: 0.4, yMin: -0.4, yMax: 0.4, columns: 9, rows: 9 },
+    frequencyBands: [{ id: "1k", minHz: 990, maxHz: 1_010, color: "#65f2bf" }],
+    beamforming: { backend, algorithm: "das", normalization: "peak" },
+  };
+  const synthesis = synthesizeMicrophoneSignals({ scene, arrayDefinition });
+  return {
+    channels: synthesis.channels,
+    sampleRate: synthesis.sampleRate,
+    arrayDefinition,
+    scene,
+    experiment,
+  };
+}
+
+function normalizedMapError(reference, actual) {
+  const referenceScale = Math.max(...reference);
+  let squaredError = 0;
+  let maximumAbsoluteError = 0;
+  for (let index = 0; index < reference.length; index += 1) {
+    const error = Math.abs(actual[index] - reference[index]) / referenceScale;
+    squaredError += error * error;
+    maximumAbsoluteError = Math.max(maximumAbsoluteError, error);
+  }
+  return {
+    rms: Math.sqrt(squaredError / reference.length),
+    maximumAbsoluteError,
+  };
+}
+
+test("WebGPU DAS가 CPU 수치 기준과 일치하거나 auto CPU fallback한다", async () => {
+  const adapter = await globalThis.navigator?.gpu?.requestAdapter();
+  if (!adapter) {
+    const fixture = createFixture("auto");
+    const result = await computeDasWithBackend(fixture);
+    assertEqual(result.backend, "cpu", "auto fallback backend without adapter");
+    assert(result.bands[0].peak.rawPower > 0, "fallback produced a usable map");
+    document.documentElement.dataset.webgpuIntegration = "cpu-fallback";
+    return;
+  }
+
+  const fixture = createFixture("webgpu");
+  const cpuResult = await computeDasMap(fixture);
+  const gpuResult = await computeDasWebGpu(fixture);
+  const cpuBand = cpuResult.bands[0];
+  const gpuBand = gpuResult.bands[0];
+  const error = normalizedMapError(cpuBand.rawMap, gpuBand.rawMap);
+
+  assertEqual(gpuResult.backend, "webgpu", "actual WebGPU backend");
+  assertEqual(gpuBand.rawMap.length, cpuBand.rawMap.length, "map size");
+  assert(error.rms <= 2e-3, `CPU/GPU normalized RMS error=${error.rms}`);
+  assert(
+    error.maximumAbsoluteError <= 1e-2,
+    `CPU/GPU normalized maximum absolute error=${error.maximumAbsoluteError}`,
+  );
+  const gridDiagonal = Math.hypot(
+    (fixture.experiment.scan.xMax - fixture.experiment.scan.xMin) /
+      (fixture.experiment.scan.columns - 1),
+    (fixture.experiment.scan.yMax - fixture.experiment.scan.yMin) /
+      (fixture.experiment.scan.rows - 1),
+  );
+  assert(
+    distance3(gpuBand.peak.position, cpuBand.peak.position) <= gridDiagonal + 1e-9,
+    "CPU/GPU peak positions are within one grid diagonal",
+  );
+  assert(
+    distance3(gpuBand.peak.position, fixture.scene.sources[0].position) <= 0.071,
+    `GPU localization error=${distance3(gpuBand.peak.position, fixture.scene.sources[0].position)} m`,
+  );
+  document.documentElement.dataset.webgpuIntegration = "gpu";
+  document.documentElement.dataset.webgpuRmsError = error.rms.toExponential(3);
+});
+
+test("192채널 · 50×50 WebGPU map이 CPU baseline 허용 오차를 만족한다", async () => {
+  const adapter = await globalThis.navigator?.gpu?.requestAdapter();
+  if (!adapter) {
+    document.documentElement.dataset.webgpuBaseline = "unavailable";
+    return;
+  }
+  const baseline = globalThis.__micarrsimCpuBaseline;
+  assert(baseline, "CPU baseline fixture is available");
+  const experiment = {
+    ...baseline.experiment,
+    beamforming: { ...baseline.experiment.beamforming, backend: "webgpu" },
+  };
+  const result = await computeDasWebGpu({ ...baseline, experiment });
+  const cpuBand = baseline.result.bands[0];
+  const gpuBand = result.bands[0];
+  const error = normalizedMapError(cpuBand.rawMap, gpuBand.rawMap);
+  const peakPowerRelativeError =
+    Math.abs(gpuBand.peak.rawPower - cpuBand.peak.rawPower) / cpuBand.peak.rawPower;
+  const gridDiagonal = Math.hypot(
+    (experiment.scan.xMax - experiment.scan.xMin) / (experiment.scan.columns - 1),
+    (experiment.scan.yMax - experiment.scan.yMin) / (experiment.scan.rows - 1),
+  );
+
+  assert(gpuBand.rawMap.every((value) => Number.isFinite(value) && value >= 0), "finite nonnegative GPU map");
+  assert(error.rms <= 0.005, `192ch CPU/GPU normalized RMS error=${error.rms}`);
+  assert(error.maximumAbsoluteError <= 0.02, `192ch CPU/GPU normalized max error=${error.maximumAbsoluteError}`);
+  assert(peakPowerRelativeError <= 0.02, `192ch CPU/GPU peak power relative error=${peakPowerRelativeError}`);
+  assert(
+    distance3(gpuBand.peak.position, cpuBand.peak.position) <= gridDiagonal + 1e-9,
+    "192ch CPU/GPU peaks are within one grid diagonal",
+  );
+  document.documentElement.dataset.webgpuBaseline = "gpu";
+  document.documentElement.dataset.webgpuBaselineMs = result.timingMs.total.toFixed(1);
+  document.documentElement.dataset.webgpuBaselineRmsError = error.rms.toExponential(3);
+});
+
+test("WebGPU bin·point 강제 batching 결과가 CPU map과 일치한다", async () => {
+  const adapter = await globalThis.navigator?.gpu?.requestAdapter();
+  if (!adapter) return;
+  const fixture = createFixture("webgpu");
+  fixture.experiment.frequencyBands = [
+    { id: "wide-1k", minHz: 900, maxHz: 1_100, color: "#65f2bf" },
+  ];
+  const cpuResult = await computeDasMap(fixture);
+  const gpuResult = await computeDasWebGpu({
+    ...fixture,
+    maximumSpectrumBatchBins: 1,
+    maximumSteeringBufferBytes: 512,
+  });
+  const error = normalizedMapError(cpuResult.bands[0].rawMap, gpuResult.bands[0].rawMap);
+
+  assertEqual(gpuResult.gpu.batchPlan.binBatchSize, 1, "forced bin batch size");
+  assertEqual(gpuResult.gpu.batchPlan.pointBatchSize, 4, "forced point batch size");
+  assert(error.rms <= 0.005, `chunked CPU/GPU normalized RMS error=${error.rms}`);
+  assert(error.maximumAbsoluteError <= 0.02, `chunked CPU/GPU max error=${error.maximumAbsoluteError}`);
+});
diff --git a/tests/test-runner.js b/tests/test-runner.js
index 622ac0e..17d109e 100644
--- a/tests/test-runner.js
+++ b/tests/test-runner.js
@@ -12,7 +12,9 @@ try {
   await import("/tests/unit/contracts.test.js");
   await import("/tests/unit/acoustics.test.js");
   await import("/tests/unit/builder.test.js");
+  await import("/tests/unit/webgpu.test.js");
   await import("/tests/integration/cpu-baseline.test.js");
+  await import("/tests/integration/webgpu.test.js");
   await import("/tests/integration/builder-ui.test.js");
 } catch (error) {
   registeredTests.push({
diff --git a/tests/unit/webgpu.test.js b/tests/unit/webgpu.test.js
new file mode 100644
index 0000000..d409e3a
--- /dev/null
+++ b/tests/unit/webgpu.test.js
@@ -0,0 +1,252 @@
+import {
+  DAS_PARAMS_BYTE_LENGTH,
+  packDasParams,
+  packMicrophoneCalibration,
+} from "/packages/webgpu-kernels/das.wgsl.js";
+import {
+  WebGpuLimitError,
+  WebGpuUnavailableError,
+  estimateWebGpuBatchPlan,
+  estimateWebGpuBatchSize,
+  packStftBatch,
+} from "/packages/webgpu-kernels/das.js";
+import { computeDasWithBackend } from "/packages/webgpu-kernels/compute.js";
+
+function assertNear(actual, expected, tolerance, label) {
+  assert(
+    Math.abs(actual - expected) <= tolerance,
+    `${label}: expected ${expected} ± ${tolerance}, received ${actual}`,
+  );
+}
+
+test("WebGPU DAS uniform과 microphone calibration을 WGSL layout으로 packing한다", () => {
+  const params = packDasParams({
+    pointCount: 81,
+    microphoneCount: 16,
+    frameCount: 7,
+    binCount: 3,
+    sampleRate: 8_000,
+    fftSize: 256,
+    speedOfSound: 343.25,
+  });
+  const view = new DataView(params);
+
+  assertEqual(params.byteLength, DAS_PARAMS_BYTE_LENGTH, "uniform byte length");
+  assertEqual(view.getUint32(0, true), 81, "point count offset");
+  assertEqual(view.getUint32(4, true), 16, "microphone count offset");
+  assertEqual(view.getUint32(8, true), 7, "frame count offset");
+  assertEqual(view.getUint32(12, true), 3, "bin count offset");
+  assertEqual(view.getFloat32(16, true), Math.fround(8_000), "sample rate offset");
+  assertEqual(view.getUint32(20, true), 256, "FFT size offset");
+  assertEqual(view.getFloat32(24, true), Math.fround(343.25), "speed of sound offset");
+  assertEqual(view.getUint32(28, true), 0, "reserved word");
+
+  const calibration = packMicrophoneCalibration([
+    { sensitivityDb: 0, phaseOffsetDeg: 0 },
+    { sensitivityDb: 6.020599913, phaseOffsetDeg: 90 },
+    { sensitivityDb: -6.020599913, phaseOffsetDeg: -45 },
+  ]);
+  assertEqual(calibration.length, 6, "two f32 values per microphone");
+  assertNear(calibration[0], 1, 1e-7, "unity inverse gain");
+  assertNear(calibration[1], 0, 1e-7, "zero phase");
+  assertNear(calibration[2], 0.5, 1e-6, "+6 dB inverse gain");
+  assertNear(calibration[3], Math.PI / 2, 1e-6, "+90 degree phase");
+  assertNear(calibration[4], 2, 1e-6, "-6 dB inverse gain");
+  assertNear(calibration[5], -Math.PI / 4, 1e-6, "-45 degree phase");
+});
+
+test("STFT batch를 bin-frame-microphone-complex 순서로 packing한다", () => {
+  const stft = {
+    frameCount: 2,
+    microphoneCount: 2,
+    spectra: new Map([
+      [3, {
+        real: Float32Array.of(1, 2, 3, 4),
+        imaginary: Float32Array.of(11, 12, 13, 14),
+      }],
+      [7, {
+        real: Float32Array.of(5, 6, 7, 8),
+        imaginary: Float32Array.of(15, 16, 17, 18),
+      }],
+    ]),
+  };
+
+  const packed = packStftBatch(stft, [7, 3]);
+  assert(packed instanceof Float32Array, "packed STFT type");
+  const expected = [
+    5, 15, 6, 16, 7, 17, 8, 18,
+    1, 11, 2, 12, 3, 13, 4, 14,
+  ];
+  assertEqual(packed.length, expected.length, "packed STFT length");
+  expected.forEach((value, index) => assertEqual(packed[index], value, `packed STFT value ${index}`));
+});
+
+test("WebGPU batch size가 storage, dispatch, memory 한도를 모두 지킨다", () => {
+  const workload = {
+    selectedBins: [4, 5, 6, 7, 8],
+    bandBins: [[4, 5, 6, 7, 8]],
+    frameCount: 2,
+    microphoneCount: 4,
+    pointCount: 8,
+    distanceBytes: 128,
+    outputBytes: 64,
+    spectrumBytesPerBin: 64,
+    baseWorkingBytes: 256,
+  };
+  const limits = {
+    maxStorageBufferBindingSize: 192,
+    maxBufferSize: 192,
+    maxComputeWorkgroupsPerDimension: 2,
+    maxComputeInvocationsPerWorkgroup: 64,
+    maxComputeWorkgroupSizeX: 64,
+  };
+
+  const batchSize = estimateWebGpuBatchSize({
+    workload,
+    limits,
+    maximumGpuWorkingBytes: 4_096,
+    maximumHostWorkingBytes: 4_096,
+  });
+  assertEqual(batchSize, 2, "dispatch y dimension is the tightest batch limit");
+  const plan = estimateWebGpuBatchPlan({
+    workload,
+    limits,
+    maximumGpuWorkingBytes: 4_096,
+    maximumHostWorkingBytes: 4_096,
+  });
+  assertEqual(plan.pointBatchSize, 2, "steering point dispatch batch");
+  assertEqual(plan.steeringBytes, 128, "two-pass steering buffer bytes");
+
+  let limitError;
+  try {
+    estimateWebGpuBatchSize({
+      workload,
+      limits: { ...limits, maxStorageBufferBindingSize: 31 },
+      maximumGpuWorkingBytes: 4_096,
+      maximumHostWorkingBytes: 4_096,
+    });
+  } catch (error) {
+    limitError = error;
+  }
+  assert(limitError instanceof WebGpuLimitError, "static calibration buffer limit error type");
+
+  let invocationLimitError;
+  try {
+    estimateWebGpuBatchSize({
+      workload,
+      limits,
+      maximumGpuWorkingBytes: 4_096,
+      maximumHostWorkingBytes: 4_096,
+      maximumMacsPerInvocation: 7,
+    });
+  } catch (error) {
+    invocationLimitError = error;
+  }
+  assert(invocationLimitError instanceof WebGpuLimitError, "per-invocation watchdog guard error type");
+
+  const scalingWorkload = {
+    selectedBins: [10, 11],
+    bandBins: [[10, 11]],
+    frameCount: 46,
+    microphoneCount: 192,
+    pointCount: 40_000,
+    distanceBytes: 40_000 * 192 * 4,
+    outputBytes: 40_000 * 4 * 2,
+    spectrumBytesPerBin: 46 * 192 * 8,
+    baseWorkingBytes: 40_000 * 192 * 4 + 40_000 * 4 * 3 + 512 * 8 * 2,
+  };
+  const scalingPlan = estimateWebGpuBatchPlan({
+    workload: scalingWorkload,
+    limits: {
+      maxStorageBufferBindingSize: 128 * 1024 ** 2,
+      maxBufferSize: 256 * 1024 ** 2,
+      maxComputeWorkgroupsPerDimension: 65_535,
+      maxComputeInvocationsPerWorkgroup: 256,
+      maxComputeWorkgroupSizeX: 256,
+      maxStorageBuffersPerShaderStage: 8,
+      maxUniformBufferBindingSize: 65_536,
+    },
+  });
+  assertEqual(scalingPlan.binBatchSize, 2, "200x200 preserves one CPU STFT bin batch");
+  assert(scalingPlan.pointBatchSize < 40_000, "200x200 steering uses point chunks");
+  assert(scalingPlan.steeringBytes <= 64 * 1024 ** 2, "steering target is respected");
+});
+
+test("backend auto는 typed WebGPU 운용 오류만 CPU로 fallback한다", async () => {
+  const calls = [];
+  const options = { experiment: { beamforming: { backend: "auto" } } };
+  const cpuResult = { backend: "cpu", bands: [], marker: "fallback-result" };
+  const result = await computeDasWithBackend(options, {
+    computeGpu: async () => {
+      calls.push("gpu");
+      throw new WebGpuUnavailableError("adapter unavailable");
+    },
+    computeCpu: async () => {
+      calls.push("cpu");
+      return cpuResult;
+    },
+  });
+
+  assertEqual(calls.join(","), "gpu,cpu", "auto backend order");
+  assertEqual(result.backend, "cpu", "fallback backend");
+  assertEqual(result.marker, "fallback-result", "CPU result is preserved");
+
+  const limitFallback = await computeDasWithBackend(options, {
+    computeGpu: async () => {
+      throw new WebGpuLimitError("device limit");
+    },
+    computeCpu: async () => cpuResult,
+  });
+  assertEqual(limitFallback.backend, "cpu", "device limit also uses CPU fallback in auto mode");
+
+  let programmingError;
+  try {
+    await computeDasWithBackend(options, {
+      computeGpu: async () => {
+        throw new TypeError("programming error");
+      },
+      computeCpu: async () => {
+        throw new Error("CPU must not run");
+      },
+    });
+  } catch (error) {
+    programmingError = error;
+  }
+  assert(programmingError instanceof TypeError, "unexpected GPU errors are not hidden by fallback");
+
+  const controller = new AbortController();
+  controller.abort();
+  let abortError;
+  try {
+    await computeDasWithBackend(
+      { ...options, signal: controller.signal },
+      {
+        computeGpu: async () => {
+          throw new WebGpuUnavailableError("adapter unavailable after cancellation");
+        },
+        computeCpu: async () => {
+          throw new Error("cancelled auto mode must not start CPU fallback");
+        },
+      },
+    );
+  } catch (error) {
+    abortError = error;
+  }
+  assertEqual(abortError?.name, "AbortError", "cancellation wins over GPU fallback");
+
+  const explicitOptions = { experiment: { beamforming: { backend: "webgpu" } } };
+  let explicitError;
+  try {
+    await computeDasWithBackend(explicitOptions, {
+      computeGpu: async () => {
+        throw new WebGpuLimitError("forced device limit");
+      },
+      computeCpu: async () => {
+        throw new Error("explicit webgpu must not run CPU fallback");
+      },
+    });
+  } catch (error) {
+    explicitError = error;
+  }
+  assert(explicitError instanceof WebGpuLimitError, "explicit WebGPU preserves operational error");
+});