矩阵乘法的性能表很容易写得漂亮，最难看出来的错误却可能只有一个：CPU 按行存数据，着色器按另一种步长读取。
先别比 M 系列和 RTX。拿一个非方阵、边缘不满块的输入，把布局与边界验收做完。本文只讨论浏览器里的 JS/WGSL 数据合同，不讨论 Rust 的异步映射和所有权，也不声称任何设备的跑分。
用 3×5 乘 5×2，而不是一直用方阵
固定 A 为 M×K、B 为 K×N、C 为 M×N，三者都按行存储：
A[row * K + k]、B[k * N + col]、C[row * N + col]。
方阵容易掩盖把 K 与 N 写反的错误。下面这段参考实现可以直接在 Node.js 运行，期望值是六个手算结果：
import assert from 'node:assert/strict'; function reference(A, B, M, K, N) { if (![M, K, N].every(x => Number.isSafeInteger(x) && x > 0)) { throw new RangeError('dimensions must be positive integers'); } if (A.length !== M * K || B.length !== K * N) { throw new RangeError('buffer length does not match shape'); } const C = new Float32Array(M * N); for (let row = 0; row < M; row++) { for (let col = 0; col < N; col++) { let sum = 0; for (let k = 0; k < K; k++) { sum += A[row * K + k] * B[k * N + col]; } C[row * N + col] = sum; } } return C; } const A = Float32Array.from({ length: 15 }, (_, i) => i + 1); const B = Float32Array.from({ length: 10 }, (_, i) => i + 1); assert.deepEqual([...reference(A, B, 3, 5, 2)], [95, 110, 220, 260, 345, 410]); assert.throws(() => reference(A, B, 5, 3, 2)); assert.throws(() => reference(A, B, 0, 5, 2)); console.log('layout reference checks passed');这个参考实现用于检查索引和形状，不是高精度数值计算库。换成一般浮点数据时，GPU 的累加顺序和中间精度可能不同，应事先规定绝对误差与相对误差阈值，不能要求逐位相等。
边缘线程不能提前越过同步点
下面用 16×16 tile。每个工作组有 256 个 invocation，两块工作组数组共占 2048 字节。最后一块可能只有一部分元素有效：读越界时补零，写越界时跳过，但所有线程仍按相同次数经过 barrier。
图中的灰格仅示意计算越界，不按格数表示16×16 tile，也不是实际分配的输出元素；输出仍只有 M×N 项。
把这段 WGSL 与后面的浏览器脚本放进同一个 HTML 文件，使用下面给出的骨架即可，不需要框架：
<!doctype html> <meta charset="utf-8"> <pre id="result">等待计算</pre> <script id="kernel" type="text/wgsl"> struct Shape { M: u32, N: u32, K: u32, padding: u32 } @group(0) @binding(0) var<uniform> shape: Shape; @group(0) @binding(1) var<storage, read> a: array<f32>; @group(0) @binding(2) var<storage, read> b: array<f32>; @group(0) @binding(3) var<storage, read_write> c: array<f32>; var<workgroup> ta: array<array<f32, 16>, 16>; var<workgroup> tb: array<array<f32, 16>, 16>; @compute @workgroup_size(16, 16) fn main(@builtin(global_invocation_id) g: vec3<u32>, @builtin(local_invocation_id) l: vec3<u32>) { var sum = 0.0; let tiles = (shape.K + 15u) / 16u; for (var t = 0u; t < tiles; t++) { let ak = t * 16u + l.x; let bk = t * 16u + l.y; ta[l.y][l.x] = 0.0; tb[l.y][l.x] = 0.0; if (g.y < shape.M && ak < shape.K) { ta[l.y][l.x] = a[g.y * shape.K + ak]; } if (bk < shape.K && g.x < shape.N) { tb[l.y][l.x] = b[bk * shape.N + g.x]; } workgroupBarrier(); for (var k = 0u; k < 16u; k++) { sum += ta[l.y][k] * tb[k][l.x]; } workgroupBarrier(); } if (g.y < shape.M && g.x < shape.N) { c[g.y * shape.N + g.x] = sum; } } </script> <script type="module"> const output = document.querySelector('#result'); async function run() { if (!navigator.gpu) throw Error('当前浏览器没有 WebGPU'); const adapter = await navigator.gpu.requestAdapter(); if (!adapter) throw Error('没有可用 adapter'); const device = await adapter.requestDevice(); const M = 3, K = 5, N = 2, bytes = M * N * 4; const limits = device.limits; if (limits.maxComputeInvocationsPerWorkgroup < 256 || limits.maxComputeWorkgroupSizeX < 16 || limits.maxComputeWorkgroupSizeY < 16 || limits.maxComputeWorkgroupStorageSize < 2048) { device.destroy(); throw Error('此设备限额不满足本示例'); } const dataA = Float32Array.from({ length: M*K }, (_, i) => i+1); const dataB = Float32Array.from({ length: K*N }, (_, i) => i+1); const storage = GPUBufferUsage.STORAGE; function input(data, usage) { if (data.byteLength > limits.maxBufferSize || (usage & storage && data.byteLength > limits.maxStorageBufferBindingSize)) { throw Error('输入缓冲超限'); } const buffer = device.createBuffer({ size: data.byteLength, usage: usage | GPUBufferUsage.COPY_DST }); device.queue.writeBuffer(buffer, 0, data); return buffer; } const shape = input(new Uint32Array([M,N,K,0]), GPUBufferUsage.UNIFORM); const a = input(dataA, storage), b = input(dataB, storage); const c = device.createBuffer({ size: bytes, usage: storage | GPUBufferUsage.COPY_SRC }); const read = device.createBuffer({ size: bytes, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST }); try { const module = device.createShaderModule({ code: document.querySelector('#kernel').textContent }); const info = await module.getCompilationInfo(); const errors = info.messages.filter(m => m.type === 'error'); if (errors.length) throw Error(errors.map(m => m.message).join('\n')); const pipeline = await device.createComputePipelineAsync({ layout: 'auto', compute: { module, entryPoint: 'main' } }); const group = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [shape,a,b,c].map((buffer,binding) => ({ binding, resource: { buffer } })) }); const encoder = device.createCommandEncoder(); const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, group); pass.dispatchWorkgroups(Math.ceil(N/16), Math.ceil(M/16)); pass.end(); encoder.copyBufferToBuffer(c,0,read,0,bytes); device.queue.submit([encoder.finish()]); await read.mapAsync(GPUMapMode.READ); const got = [...new Float32Array(read.getMappedRange())]; read.unmap(); const expected = [95,110,220,260,345,410]; if (!got.every((x,i) => Math.abs(x-expected[i]) < 1e-5)) { throw Error('结果不匹配：' + JSON.stringify(got)); } output.textContent = JSON.stringify({ status: 'PASS', M,K,N, got },null,2); } finally { [shape,a,b,c,read].forEach(x => x.destroy()); device.destroy(); } } run().catch(error => { output.textContent = '未通过：' + error.message; }); </script>它是一个小输入验收程序，不是允许任意矩阵尺寸的通用 GEMM API。若要开放输入，还必须检查乘积是否安全、输出缓冲和 dispatch 是否超限，以及内存分配失败后的清理。这里不能因为样例只分配几十字节，就顺手写成“支持所有大模型”。
限额属于当前 device，不属于显卡宣传页
物理机器有多少统一内存，不能直接回答浏览器允许一个 storage binding 有多大。MDN 说明浏览器还可能按档位暴露 limits。因此要记录实际 adapter、requestDevice 条件以及最终 device.limits，不能拿显卡内存容量代替。GPUSupportedLimits
WGSL 的结构体、数组和向量也有自己的布局约束。本文把四个 u32 放入一个16字节 uniform，避免将 JS 对象字段顺序误当字节布局。更复杂的数据合同应逐字段计算 offset 和 stride。WGSL 内存布局规范
算对之后，再讨论怎样计时
至少分开记录数据上传、管线创建、执行与回读。拿 await mapAsync 前后的总时间报“GPU纯算力”，会把提交、等待和回读都算进去。没有 timestamp-query 或没有正确的测量条件，就报告包含哪些阶段的墙钟耗时，别补一个 TFLOPS 数字。
本文不提供 M 系列、RTX 或手机对比表。CPU 参考、着色器编译、实际 GPU 执行是三个不同的验证层级；没有运行的层级必须标明。浏览器例子需要安全上下文与可用 WebGPU，不能预设所有设备可运行。管线创建文档
WGSL链接为持续更新的编辑草案；硬件性能与跨设备覆盖仍无实测结论。
本地核验：Node.js v26.3.0 执行 CPU 参考断言通过；浏览器宿主 JavaScript 已用node --check做语法检查。WGSL 编译和 GPU 执行尚未实测，因此没有硬件正确性或性能结论。
