WebGPU 性能榨汁机：利用 Workgroup 共享内存极限加速双边滤波

2026/7/7 10:32:38 2 0 0 0

在 Web 端运行高画质、实时的图像后处理算法，过去常常受限于 WebGL 的性能瓶颈。随着 WebGPU 的正式到来，Web 端开发者终于拥有了直接掌控 GPU 计算管线（Compute Pipeline）的能力。

双边滤波（Bilateral Filter）作为一种保边去噪（Edge-preserving smoothing）的经典算法，广泛应用于图像降噪、人脸美颜、景深渲染（Depth of Field）等后处理场景。然而，双边滤波由于其非可分性（Non-separable），无法像高斯模糊那样拆分为横向和纵向两次一维滤波，必须进行昂贵的二维窗口采样。

在 $16 \times 16$ 的工作组中，如果滤波半径 $R=4$，每个像素需要采样 $(2R+1)^2 = 81$ 个邻域像素。如果采用最直观的朴素实现，每个线程直接进行 81 次全局纹理采样（Global Texture Read），GPU 的显存带宽（VRAM Bandwidth）将瞬间被榨干，导致严重的延迟。

本文将深入探讨如何利用 WebGPU 中的 Workgroup Shared Memory（工作组共享内存），将显存访问次数降低至原来的几十分之一，实现极限加速。

为什么需要 Workgroup 共享内存？

在 GPU 的存储金字塔中：

全局内存（Global Memory / Texture）：容量大，但延迟高（通常需要数百个时钟周期），带宽有限。
工作组共享内存（Workgroup Shared Memory）：位于芯片内部（On-chip），延迟极低，带宽极大，其访问速度接近 L1 缓存。它由同一个 Workgroup 内的所有线程（Thread）共享。

双边滤波中，相邻像素在计算时，其采样窗口有极大的重叠。如下图所示：

  [ 线程 A 采样窗口 ]
  ┌───┬───┬───┬───┐
  │   │   │   │   │
  ├───┼───┼───┼───┤
  │   │ o │ x │   │  <-- 像素 'o' 和 'x' 都在重复读取这些重叠的格点
  └───┴───┴───┴───┘
      [ 线程 B 采样窗口 ]

如果我们能让 Workgroup 内的所有线程协作，一次性把该工作组所需的所有纹理像素全部加载到高速共享内存中，后续的 81 次窗口采样全部在共享内存中进行，就能彻底解放显存带宽。

核心挑战：围栏（Apron / Padding）区域的处理

如果工作组大小（Workgroup Size）设置为 $16 \times 16$，那么这个工作组内部有 256 个线程，负责输出 $16 \times 16$ 区域的滤波结果。

但是，边缘的线程在计算邻域时，需要超出这 $16 \times 16$ 边界之外的像素。如果滤波半径 $R=4$，我们需要向外扩展 4 个像素的“围栏”区域。

因此，我们需要在共享内存中开辟一个大小为 $(16 + 2 \times 4) \times (16 + 2 \times 4) = 24 \times 24$ 的缓冲区：

 ┌───────────────────────────┐ ──┐
 │       Apron (围栏区)      │   │ FILTER_RADIUS (4px)
 │   ┌───────────────────┐   │ ──┼
 │   │                   │   │   │
 │   │  Active Workgroup │   │   │ WORKGROUP_SIZE (16px)
 │   │    (16x16 线程)   │   │   │
 │   └───────────────────┘   │ ──┼
 │                           │   │ FILTER_RADIUS (4px)
 └───────────────────────────┘ ──┘
 ◄───────── 24px ────────────►

协作式加载（Cooperative Loading）

共享内存的大小是 $24 \times 24 = 576$ 个像素，而我们的工作组只有 256 个线程。
这意味着，每个线程需要协同加载约 $576 \div 256 \approx 2.25$ 个像素。

我们需要在代码中，让这 256 个线程通过循环，将这 576 个像素分批从全局纹理加载到共享内存中，并处理好越界采样（Texture Clamping）。

WGSL 实现：高效双边滤波计算着色器

下面是完整的 WebGPU 着色器语言 (WGSL) 实现。代码中详细展示了协作加载、屏障同步、以及基于共享内存的双边滤波计算。

// 绑定资源
@group(0) @binding(0) var inputTex: texture_2d<f32>;
@group(0) @binding(1) var outputTex: texture_storage_2d<rgba8unorm, write>;

// 算法常量定义
const WORKGROUP_SIZE: u32 = 16u;
const FILTER_RADIUS: u32 = 4u;
// 共享内存尺寸 = 工作组尺寸 + 2 * 半径
const SHARED_SIZE: u32 = 24u; // 16 + 2 * 4

// 声明工作组共享内存
var<workgroup> shared_pixels: array<array<vec4<f32>, SHARED_SIZE>, SHARED_SIZE>;

// 双边滤波参数
const SIGMA_SPACE: f32 = 4.0;
const SIGMA_RANGE: f32 = 0.15;

@compute @workgroup_size(16, 16)
fn main(
    @builtin(local_invocation_id) local_id: vec3<u32>,
    @builtin(global_invocation_id) global_id: vec3<u32>,
    @builtin(workgroup_id) workgroup_id: vec3<u32>
) {
    let thread_idx = local_id.y * WORKGROUP_SIZE + local_id.x; // 当前线程在一维中的索引 (0 ~ 255)
    let total_threads = WORKGROUP_SIZE * WORKGROUP_SIZE;       // 256
    let total_pixels_to_load = SHARED_SIZE * SHARED_SIZE;      // 576

    // 1. 协作式加载：将像素从全局纹理载入共享内存
    // 576 个像素分派给 256 个线程，每个线程最多循环 3 次
    for (var i = 0u; i < 3u; i = i + 1u) {
        let linear_idx = thread_idx + i * total_threads;
        
        if (linear_idx < total_pixels_to_load) {
            let shared_y = linear_idx / SHARED_SIZE;
            let shared_x = linear_idx % SHARED_SIZE;

            // 计算当前要载入的像素在全局纹理中的坐标
            // 以当前 Workgroup 的左上角为基准，向左上方偏移 FILTER_RADIUS 个像素
            let global_tex_coords = vec2<i32>(workgroup_id.xy * WORKGROUP_SIZE) 
                                  + vec2<i32>(i32(shared_x), i32(shared_y)) 
                                  - vec2<i32>(i32(FILTER_RADIUS));

            // 安全边界处理：钳制到纹理合法尺寸内
            let dims = textureDimensions(inputTex);
            let clamped_coords = clamp(
                global_tex_coords, 
                vec2<i32>(0, 0), 
                vec2<i32>(dims) - vec2<i32>(1, 1)
            );

            // 从 VRAM 读取并写入 Share Group Memory
            shared_pixels[shared_y][shared_x] = textureLoad(inputTex, clamped_coords, 0);
        }
    }

    // 2. 内存屏障（执行同步）
    // 必须等待所有线程都完成了共享内存的写入，才能开始后续的滤波计算
    workgroupBarrier();

    // 3. 滤波计算
    let img_dims = textureDimensions(outputTex);
    // 仅对未越界的像素进行滤波写入
    if (global_id.x < img_dims.x && global_id.y < img_dims.y) {
        
        // 当前线程在共享内存中对应的中心位置
        let center_x = local_id.x + FILTER_RADIUS;
        let center_y = local_id.y + FILTER_RADIUS;
        let center_color = shared_pixels[center_y][center_x];

        var sum_color = vec4<f32>(0.0);
        var sum_weight = 0.0;

        // 在共享内存中进行 2D 窗口采样
        for (var dy = -i32(FILTER_RADIUS); dy <= i32(FILTER_RADIUS); dy = dy + 1) {
            for (var dx = -i32(FILTER_RADIUS); dx <= i32(FILTER_RADIUS); dx = dx + 1) {
                
                let sample_x = u32(i32(center_x) + dx);
                let sample_y = u32(i32(center_y) + dy);
                let sample_color = shared_pixels[sample_y][sample_x];

                // 空间域权重 (Gaussian Space Weight)
                let dist_sq = f32(dx * dx + dy * dy);
                let space_weight = exp(-dist_sq / (2.0 * SIGMA_SPACE * SIGMA_SPACE));

                // 值域权重 (Gaussian Range Weight)
                let color_diff = sample_color - center_color;
                let color_dist_sq = dot(color_diff, color_diff);
                let range_weight = exp(-color_dist_sq / (2.0 * SIGMA_RANGE * SIGMA_RANGE));

                // 联合权重
                let weight = space_weight * range_weight;
                
                sum_color = sum_color + sample_color * weight;
                sum_weight = sum_weight + weight;
            }
        }

        let filtered_color = sum_color / max(sum_weight, 0.0001);
        
        // 将最终结果写入输出纹理
        textureStore(outputTex, global_id.xy, vec4<f32>(filtered_color.rgb, center_color.a));
    }
}

性能收益深度剖析

为什么上面的做法能带来飞跃性的性能提升？我们来算一笔账。

假设我们要处理一张 $1920 \times 1080$（1080P，约 200 万像素）的图像，滤波半径 $R=4$。

1. 朴素实现（无共享内存）

每个像素执行一次 Shader，每次都要读取 81 个邻域像素。

总显存读取次数：$2,000,000 \times 81 \approx \mathbf{1.62}$ 亿次。
每一次读取都是全局内存访问，会引起剧烈的 L2 缓存抖动（Cache Thrashing）。

2. Workgroup 共享内存优化实现

我们以 $16 \times 16$ 为一个工作组，处理 256 个像素。
由于采用了协作式加载，每个工作组只需加载 $24 \times 24 = 576$ 个像素到共享内存中。

总工作组数量：$1920/16 \times 1080/16 = 120 \times 68 = 8160$ 个。
总显存读取次数：$8160 \times 576 \approx \mathbf{470}$ 万次。
对 VRAM 读取次数减少了约 97.1%！

后续的 1.62 亿次邻域采样全部发生在 Workgroup Shared Memory 中，耗时几乎可以忽略不计。

避坑与调优指南

在实践 WebGPU Compute Shader 优化时，以下几点需要特别注意：

Workgroup Barrier 的性能开销
workgroupBarrier() 是一个硬件级别的同步指令，会阻塞当前工作组中的线程，直到所有线程到达这一行。不恰当的屏障调用会降低 GPU 占用率（Occupancy）。在加载完共享内存后，只需调用一次同步即可，绝对不要将 workgroupBarrier() 写入任何条件分支或循环内部。
共享内存的大小限制（Limits）
WebGPU 规范规定，单个工作组可用的最大共享内存容量（maxComputeWorkgroupStorageSize）标准设备上至少为 16384 字节（16 KB）。
在上面的例子中，我们申请了 array<array<vec4<f32>, 24>, 24>。
大小为：$24 \times 24 \times 16 \text{ bytes (vec4)} = 9216 \text{ 字节} \approx 9 \text{ KB}$。
这完全在安全限制内。如果你的滤波半径 $R$ 极大（例如 $R \ge 8$），则需要适当减小工作组大小（例如改为 $8 \times 8$ 线程组），以防共享内存溢出。
内存对齐与 Padding
WGSL 中的二维数组在映射到硬件时是连续存放的。使用 array<array<T, N>, M> 格式能很好地利用硬件合并合并访问（Coalesced Memory Access），避免因不规则步长导致的 Bank Conflict（银行冲突）。

总结

通过 WebGPU 的 Workgroup Shared Memory，我们成功将双边滤波这种高复杂度的图像算法，在 Web 端跑出了接近原生的性能。这种“空间换时间、共享减带宽”的优化思路，不仅适用于双边滤波，也是实现高效 Web 卷积、粒子系统、流体模拟等 GPGPU 场景的必经之路。

如果你正在开发下一代 Web 图像编辑器、3D 渲染器或 H5 游戏，不妨立刻重构你的 Compute Shader，榨干 WebGPU 的最后一滴硬件红利！

零拷贝极客 WebGPU WGSL 双边滤波