WebGPU 计算着色器深度实战：浏览器端 GPGPU 通用计算与深度学习推理

第一节：矩阵乘法 GPU 加速实战

2.1 Naive 实现与性能基线

矩阵乘法是 GPU 计算的经典入门案例，也是理解 GPU 并行模式的基石。我们先实现一个"朴素"版本（无共享内存），然后逐步优化：

@group(0) @binding(0) var<storage, read> matA: array<f32>;
@group(0) @binding(1) var<storage, read> matB: array<f32>;
@group(0) @binding(2) var<storage, read_write> matC: array<f32>;

override K: u32;  // A的列数 = B的行数
override N: u32;  // B的列数
override M: u32;  // A的行数

@compute @workgroup_size(TILE_SIZE, TILE_SIZE, 1)
fn naive_matmul(@builtin(global_invocation_id) gid: vec3<u32>) {
    let row = gid.x;
    let col = gid.y;
    if (row >= M || col >= N) { return; }
    var acc: f32 = 0.0;
    for (var k: u32 = 0u; k < K; k = k + 1u) {
        acc = acc + matA[row * K + k] * matB[k * N + col];
    }
    matC[row * N + col] = acc;
}

这个朴素版本的问题：每个元素计算需要 2K 次全局内存读取（A的一行 + B的一列），而 GPU 全局内存延迟约 400-800 个时钟周期。对于 M=N=K=1024 的矩阵，内存墙导致 SM（Streaming Multiprocessor）利用率极低。

2.2 Tiling 优化：共享内存的威力

核心思想：将矩阵分块加载到 Shared Memory（片上 SRAM，延迟约 30 cycles）：

const TILE = 16u;
var<workgroup> tileA: array<f32, 256>; // TILE * TILE
var<workgroup> tileB: array<f32, 256>;

@compute @workgroup_size(TILE, TILE, 1)
fn tiled_matmul(@builtin(global_invocation_id) gid: vec3<u32>,
                 @builtin(local_invocation_id) lid: vec3<u32>) {
    let row = gid.x; let col = gid.y;
    let local_row = lid.x; let local_col = lid.y;
    var acc: f32 = 0.0;
    // 每次迭代加载一个 tile 并对部分积累加
    for (var t: u32 = 0u; t < K; t = t + TILE) {
        // 协作加载 tile 到共享内存
        tileA[local_row * TILE + local_col] = matA[row * K + (t + local_col)];
        tileB[local_row * TILE + local_col] = matB[(t + local_row) * N + col];
        workgroupBarrier();
        // 计算当前 tile 的部分积
        for (var k: u32 = 0u; k < TILE; k = k + 1u) {
            acc = acc + tileA[local_row * TILE + k] * tileB[k * TILE + local_col];
        }
        workgroupBarrier();
    }
    matC[row * N + col] = acc;
}

2.3 性能基准数据

第三节：Parallel Reduction — GPU 求和的艺术

3.1 为什么 Reduction 难？

Reduction（归约）操作——求和、求最大/最小值、求均值——在 GPU 上面临 bank conflict 和warp divergence 两大挑战。一个将 N 个元素求和的操作，最优复杂度为 O(log N) 而非 O(N)。

3.2 经典 Nearest Power-of-2 算法

// 阶段 1：每个 workgroup 将 sdata 数组两两求和，步长 stride 每轮减半
var<workgroup> sdata: array<f32, 256>;
@compute @workgroup_size(256, 1, 1)
fn reduce_sum(@builtin(global_invocation_id) gid: vec3<u32>,
               @builtin(local_invocation_id) lid: vec3<u32>) {
    sdata[lid.x] = input[gid.x];
    workgroupBarrier();
    for (var stride: u32 = 128u; stride > 0u; stride = stride >> 1u) {
        if (lid.x < stride) {
            sdata[lid.x] = sdata[lid.x] + sdata[lid.x + stride];
        }
        workgroupBarrier();
    }
    if (lid.x == 0u) {
        output[gid.y] = sdata[0]; // 每个 workgroup 一个部分和
    }
}

3.3 Bank Conflict 规避

GPU Shared Memory 分为 32 个 Bank，每个 Bank 宽度 4 字节。当同一 Warp 中多个线程访问同一 Bank 的不同地址时，发生 Bank Conflict（串行化）。解法：使用交错寻址模式替代顺序访问模式。

第四节：2D 卷积 GPU 加速与深度学习推理

4.1 im2col + 矩阵乘法

将卷积操作转换为矩阵乘法是 cuDNN 的核心策略。对于输入特征图 (H, W, C_in) 和卷积核 (K, K, C_in, C_out)，先将每个感受野展开为一列，生成 (H'×W', K×K×C_in) 矩阵，然后与 (K×K×C_in, C_out) 权重矩阵相乘。WebGPU 上的实现方式与 Tiled Matrix Multiply 类似。

4.2 Direct Convolution（Winograd/Fast算法）

Winograd 算法将 3×3 卷积的乘法次数从 9 减少到 4（针对 F(2×2, 3×3)），代价是额外的加法变换。WebGPU 实现 Winograd 的核心在于利用 Shared Memory 缓存输入 tile 并预计算变换矩阵。

4.3 移动端 GPU 上的 MobileNet 推理

深度可分离卷积（Depthwise Separable Convolution）是 MobileNet 的核心算子：

• Depthwise：每个输入通道独立卷积，计算量小但缺乏跨通道信息融合
• Pointwise：1×1 卷积实现跨通道混合

在移动 GPU（如 Adreno 650 / Mali-G77）上，WebGPU 推理 MobileNetV2（输入 224×224×3）的延迟约 15–28ms，足以在浏览器端实现 30fps 实时目标检测。

第五节：与竞品技术的横向对比

第六节：生产级优化经验总结

• 减少 CPU-GPU 传输：使用 Mapped Buffer 或 Copy Buffer 预加载数据，避免频繁 writeBuffer
• Pipeline Cache
• Avoid Readback：尽量在 GPU 端完成全链路，只在最终结果时使用 getMappedRange
• FP16 加速：若 GPU 支持，使用 f16（若精度可接受）可获得 1.5–2x 性能提升
• 双缓冲（Double Buffering）：在渲染+计算混合场景下交替使用两套 Buffer，隐藏延迟