Part 11: WebGPU中如何使用存储纹理

WebGPU中如何使用存储纹理

摘要：本文详细介绍了WebGPU中存储纹理（storage texture）的使用方法。存储纹理允许计算着色器直接读写图像数据，为GPU图像处理（如卷积、滤波、色彩变换等）提供了高效途径。文章首先讲解存储纹理的基本概念，包括创建方法、WGSL声明语法、支持的16种核心格式以及三种访问模式（write、read、read_write）。随后通过一个完整的图像锐化示例，展示如何在实际项目中结合计算着色器和渲染管线使用存储纹理，并分析了坐标处理、边界问题、纹理使用标志等关键细节。

计算着色器除了可以访问存储缓冲区进行数据处理外，也可以访问存储纹理（storage texture），因此可以直接对图像数据进行读写，这让GPU上的图像处理（卷积、滤波、色彩变换等）变得高效而直接。

存储纹理

与图像应用中常用的常规纹理不同，存储纹理允许着色器直接向其中写入数据，而不需要先把结果整理成绘制到渲染附件的形式。在调用createTexture创建存储纹理时，usage需包含GPUTextureUsage.STORAGE_BINDING；如果这张纹理之后还要被当作普通纹理采样（例如在渲染通道中显示），还需要同时包含GPUTextureUsage.TEXTURE_BINDING。

要想在计算着色器中访问存储纹理，需要声明合适的存储纹理类型，如：

@group(…) @binding(…) var name: storage_type<format, access>;

存储纹理的类型有：

format必须是WebGPU核心规范支持用作存储纹理的格式之一，一共16种：

r32uint、r32sint、r32float、 rgba8unorm、rgba8snorm、rgba8uint、rgba8sint、 rgba16uint、rgba16sint、rgba16float、 rgba32uint、rgba32sint、rgba32float、 rg32uint、rg32sint、rg32float

这些格式都只含有一个（r）、两个（rg）或四个（rgba）颜色通道；存储纹理不支持深度/模板值访问，也不支持多重采样。需要特别说明的是，bgra8unorm——虽然它常被用作canvas的首选格式——并不在上述核心16种存储纹理格式之列，只有当适配器支持bgra8unorm-storage这一可选特性、并在requestDevice时显式申请之后，才能把bgra8unorm格式的纹理用作存储纹理。

access（访问模式）用来说明着色器如何使用这张存储纹理，在WGSL类型声明中的关键字是write、read或read_write（注意是下划线，不是连字符）。这三者的可用性并不相同：

• write是核心（core）功能，任何支持WebGPU的设备都可以使用，也是texture_storage_*最初、也是最常见的用法。
• read和read_write是后来才加入的扩展能力，需要先检测navigator.gpu.wgslLanguageFeatures.has("readonly_and_readwrite_storage_textures")是否为true，并在着色器代码开头加上requires readonly_and_readwrite_storage_textures;声明，才能使用；对应地，创建绑定组布局时storageTexture.access属性要设置为字符串"read-only"或"read-write"（这里用的是连字符，与WGSL关键字的下划线写法不同，注意区分）。若只需要只读或只写，尽量选择read或write而不是read_write，因为可读写存储纹理对硬件的要求更高，支持的格式也更有限。

声明了存储纹理之后，就可以调用textureStore向其中写入数据（要求该纹理的access是write或read_write）：

textureStore(tex, vec2u(3, 3), vec4f(1.0, 1.0, 1.0, 1.0));

它的第一个参数是进行写入的纹理，第二个参数是写入的位置（整数坐标，不需要归一化），第三个是写入的值。如果access是read或read_write，则还可以调用textureLoad(tex, coords)从存储纹理中读取数据（不需要像普通纹理那样传入mip等级，因为存储纹理没有mipmap）。

例子：图像锐化

下面这个例子先用计算着色器对一张640×640的图片做锐化处理，把结果写入一张存储纹理，再用渲染管线把这张存储纹理绘制到canvas上。

constcomputeCode=`@group(0) @binding(0) var in_tex : texture_2d<f32>; @group(0) @binding(1) var out_tex : texture_storage_2d<rgba8unorm, write>; override group_size_x: u32; override group_size_y: u32; @compute @workgroup_size(group_size_x, group_size_y) fn cs_main( @builtin(global_invocation_id) global_id: vec3<u32>) { var x = i32(global_id.x); var y = i32(global_id.y); // 图像锐化的卷积核 var k = array<f32, 9>(-1.0, -1.0, -1.0, -1.0, 9.0, -1.0, -1.0, -1.0, -1.0); // 读取指定位置的纹素 var texel: vec4f = k[0] * textureLoad(in_tex, vec2(x-1, y-1), 0) + k[1] * textureLoad(in_tex, vec2(x, y-1), 0) + k[2] * textureLoad(in_tex, vec2(x+1, y-1), 0) + k[3] * textureLoad(in_tex, vec2(x-1, y), 0) + k[4] * textureLoad(in_tex, vec2(x, y), 0) + k[5] * textureLoad(in_tex, vec2(x+1, y), 0) + k[6] * textureLoad(in_tex, vec2(x-1, y+1), 0) + k[7] * textureLoad(in_tex, vec2(x, y+1), 0) + k[8] * textureLoad(in_tex, vec2(x+1, y+1), 0); // 将结果写入到存储纹理中 textureStore(out_tex, vec2(x, y), texel); }`;constrenderCode=`struct DataStruct { @builtin(position) pos: vec4f, @location(0) uvPos: vec2f, } @group(0) @binding(0) var sam : sampler; @group(0) @binding(1) var tex : texture_2d<f32>; @vertex fn vs_main(@location(0) coords: vec2f, @location(1) uvCoords: vec2f) -> DataStruct { var outData: DataStruct; outData.pos = vec4f(coords, 0.0, 1.0); outData.uvPos = uvCoords; return outData; } @fragment fn fs_main(fragData: DataStruct) -> @location(0) vec4f { return textureSample(tex, sam, fragData.uvPos); }`;asyncfunctionrunExample(){// 获取逻辑设备和canvas元素// 获取WebGPU上下文constcontext=canvas.getContext("webgpu");if(!context){thrownewError("无法获取canvas的WebGPU上下文");}// 配置上下文constcanvasFormat=navigator.gpu.getPreferredCanvasFormat();context.configure({device:device,format:canvasFormat,});// 创建指令编码器constencoder=device.createCommandEncoder();if(!encoder){thrownewError("创建指令编码器失败");}// 创建采样器constsampler=device.createSampler({magFilter:"linear",minFilter:"linear",});// 创建bitmap，假设图像为640x640constresponse=awaitfetch("example.png");constimageBitmap=awaitcreateImageBitmap(awaitresponse.blob());// 创建纹理，用于在计算着色器中作为输入采样/读取// RENDER_ATTACHMENT是copyExternalImageToTexture强制要求的usage，// 并不代表这张纹理真的会被用作渲染附件constimageTexture=device.createTexture({size:[imageBitmap.width,imageBitmap.height],format:"rgba8unorm",usage:GPUTextureUsage.TEXTURE_BINDING|GPUTextureUsage.COPY_DST|GPUTextureUsage.RENDER_ATTACHMENT,});// 将数据写入到纹理中device.queue.copyExternalImageToTexture({source:imageBitmap},{texture:imageTexture},[imageBitmap.width,imageBitmap.height],);// 创建存储纹理以保存结果：既要能被计算着色器写入（STORAGE_BINDING），// 又要能在渲染通道中被采样（TEXTURE_BINDING）conststorageTexture=device.createTexture({size:[imageBitmap.width,imageBitmap.height],format:"rgba8unorm",usage:GPUTextureUsage.TEXTURE_BINDING|GPUTextureUsage.COPY_DST|GPUTextureUsage.STORAGE_BINDING,});// 创建计算着色器模块constcomputeModule=device.createShaderModule({label:"Shader module 0",code:computeCode,});// 创建计算通道编码器constcomputePass=encoder.beginComputePass({label:"Compute Pass 0",});// 定义计算过程constcomputePipeline=device.createComputePipeline({label:"Compute Pipeline 0",layout:"auto",compute:{module:computeModule,entryPoint:"cs_main",constants:{group_size_x:16,group_size_y:16,},},});computePass.setPipeline(computePipeline);// 获取绑定组布局constcomputeBindGroupLayout=computePipeline.getBindGroupLayout(0);// 创建绑定组绑定纹理letcomputeBindGroup=device.createBindGroup({layout:computeBindGroupLayout,entries:[{binding:0,resource:imageTexture.createView({dimension:"2d",}),},{binding:1,resource:storageTexture.createView({dimension:"2d",}),},],});computePass.setBindGroup(0,computeBindGroup);// 编码计算指令：图像为640x640，工作组大小为16x16，// 因此每个维度需要 640 / 16 = 40 个工作组，恰好覆盖每一个像素computePass.dispatchWorkgroups(40,40);// 结束计算通道编码computePass.end();// 创建渲染通道编码器constrenderPass=encoder.beginRenderPass({colorAttachments:[{view:context.getCurrentTexture().createView(),loadOp:"clear",clearValue:{r:0.9,g:0.9,b:0.9,a:1.0},storeOp:"store",},],});// 定义顶点坐标和纹理坐标constvertexData=newFloat32Array([-1.0,1.0,0.0,0.0,// 第一个顶点-1.0,-1.0,0.0,1.0,// 第二个1.0,1.0,1.0,0.0,// 第三个1.0,-1.0,1.0,1.0,// 第四个]);// 创建顶点缓冲区constvertexBuffer=device.createBuffer({label:"Example vertex buffer",size:vertexData.byteLength,usage:GPUBufferUsage.VERTEX|GPUBufferUsage.COPY_DST,});// 将顶点数据写入顶点缓冲区device.queue.writeBuffer(vertexBuffer,0,vertexData);renderPass.setVertexBuffer(0,vertexBuffer);// 定义缓冲区数据布局constbufferLayout={arrayStride:16,attributes:[{format:"float32x2",offset:0,shaderLocation:0},{format:"float32x2",offset:8,shaderLocation:1},],};// 创建渲染着色器模块constrenderModule=device.createShaderModule({label:"Example shader module",code:renderCode,});// 定义渲染过程constrenderPipeline=device.createRenderPipeline({layout:"auto",vertex:{module:renderModule,entryPoint:"vs_main",buffers:[bufferLayout],},fragment:{module:renderModule,entryPoint:"fs_main",targets:[{format:canvasFormat,},],},primitive:{topology:"triangle-strip",},});renderPass.setPipeline(renderPipeline);// 获取渲染管线绑定组布局constrenderBindGroupLayout=renderPipeline.getBindGroupLayout(0);// 创建绑定组letrenderBindGroup=device.createBindGroup({layout:renderBindGroupLayout,entries:[{binding:0,resource:sampler,},{binding:1,resource:storageTexture.createView(),},],});// 关联绑定组renderPass.setBindGroup(0,renderBindGroup);// 绘制顶点renderPass.draw(4);renderPass.end();// 向GPU提交指令device.queue.submit([encoder.finish()]);}runExample();

几点值得注意的细节：

• 片段着色器与计算着色器访问纹理的方式不同。片段着色器需要借助采样器调用textureSample来访问纹素，使用的是归一化到[0, 1]区间的浮点坐标，并支持双线性等滤波方式；计算着色器（以及access不为write的存储纹理）则使用textureLoad直接按整数纹素坐标读取，不做任何滤波或插值。

• 坐标类型与图像边界。由于global_invocation_id的类型是vec3<u32>（无符号整数），如果直接用它计算x-1、y-1，当x或y等于0时会发生无符号整数下溢（变成一个非常大的正数），导致对in_tex的越界访问。WGSL规范保证越界的textureLoad不会导致程序崩溃，但返回值是实现相关的，通常并不是期望的颜色。上面的代码在读取前先把x、y转换成有符号整数i32，这样x-1、y-1在边界处会得到负数而不是回绕的巨大整数，textureLoad仍然会把它当作越界处理（返回值同样是实现相关的），图像最外圈的一圈像素可能出现结果不准确的情况。这是几乎所有卷积类算法都会遇到的边界问题，实际项目中通常需要额外处理边界（例如夹取坐标到图像范围内，或是提前为图像加一圈"padding"），这里为了让示例简洁而没有加入这部分逻辑。

• imageTexture为什么要包含RENDER_ATTACHMENT。这张纹理本身并不会被当作渲染附件使用，但copyExternalImageToTexture这个API在校验时要求目标纹理的usage必须同时包含COPY_DST和RENDER_ATTACHMENT（部分浏览器实现内部借助一次隐式的渲染操作来完成颜色空间转换、翻转Y轴等工作），所以即便用不到渲染附件的语义，也必须声明这个标志，否则创建纹理或调用拷贝方法时会失败。

• 两个通道共用同一张存储纹理是安全的。这段代码里，计算通道把锐化结果写入storageTexture，渲染通道又把它当作普通纹理采样、绘制到canvas上。因为这两个通道的编码指令都提交自同一个encoder、并在同一次device.queue.submit中一起提交，WebGPU会保证它们按编码顺序执行，并自动插入必要的同步，所以渲染通道读到的一定是计算通道写入完成之后的数据，不需要开发者手动处理同步。

• 工作组数量与图像大小的对应关系。着色器里@workgroup_size(group_size_x, group_size_y)通过可覆盖常量设置为16×16，dispatchWorkgroups(40, 40)则调度了40 × 40个工作组，两者相乘正好是640 × 640，覆盖了图像的每一个像素、且不多不少。如果图像宽高不是工作组大小的整数倍，调度的工作组数量通常需要向上取整（例如使用Math.ceil(width / group_size_x)），并在着色器中加入边界判断，避免对超出图像范围的坐标进行textureLoad或textureStore。