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_typeformat, access;存储纹理的类型有存储纹理类型描述texture_storage_1d1维存储纹理texture_storage_2d2维存储纹理texture_storage_2d_array2维存储纹理数组texture_storage_3d3维存储纹理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_writetextureStore(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上。constcomputeCodegroup(0) binding(0) var in_tex : texture_2df32; group(0) binding(1) var out_tex : texture_storage_2drgba8unorm, 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: vec3u32) { var x i32(global_id.x); var y i32(global_id.y); // 图像锐化的卷积核 var k arrayf32, 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(x1, 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(x1, y), 0) k[6] * textureLoad(in_tex, vec2(x-1, y1), 0) k[7] * textureLoad(in_tex, vec2(x, y1), 0) k[8] * textureLoad(in_tex, vec2(x1, y1), 0); // 将结果写入到存储纹理中 textureStore(out_tex, vec2(x, y), texel); };constrenderCodestruct DataStruct { builtin(position) pos: vec4f, location(0) uvPos: vec2f, } group(0) binding(0) var sam : sampler; group(0) binding(1) var tex : texture_2df32; 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上下文constcontextcanvas.getContext(webgpu);if(!context){thrownewError(无法获取canvas的WebGPU上下文);}// 配置上下文constcanvasFormatnavigator.gpu.getPreferredCanvasFormat();context.configure({device:device,format:canvasFormat,});// 创建指令编码器constencoderdevice.createCommandEncoder();if(!encoder){thrownewError(创建指令编码器失败);}// 创建采样器constsamplerdevice.createSampler({magFilter:linear,minFilter:linear,});// 创建bitmap假设图像为640x640constresponseawaitfetch(example.png);constimageBitmapawaitcreateImageBitmap(awaitresponse.blob());// 创建纹理用于在计算着色器中作为输入采样/读取// RENDER_ATTACHMENT是copyExternalImageToTexture强制要求的usage// 并不代表这张纹理真的会被用作渲染附件constimageTexturedevice.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_BINDINGconststorageTexturedevice.createTexture({size:[imageBitmap.width,imageBitmap.height],format:rgba8unorm,usage:GPUTextureUsage.TEXTURE_BINDING|GPUTextureUsage.COPY_DST|GPUTextureUsage.STORAGE_BINDING,});// 创建计算着色器模块constcomputeModuledevice.createShaderModule({label:Shader module 0,code:computeCode,});// 创建计算通道编码器constcomputePassencoder.beginComputePass({label:Compute Pass 0,});// 定义计算过程constcomputePipelinedevice.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);// 获取绑定组布局constcomputeBindGroupLayoutcomputePipeline.getBindGroupLayout(0);// 创建绑定组绑定纹理letcomputeBindGroupdevice.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();// 创建渲染通道编码器constrenderPassencoder.beginRenderPass({colorAttachments:[{view:context.getCurrentTexture().createView(),loadOp:clear,clearValue:{r:0.9,g:0.9,b:0.9,a:1.0},storeOp:store,},],});// 定义顶点坐标和纹理坐标constvertexDatanewFloat32Array([-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,// 第四个]);// 创建顶点缓冲区constvertexBufferdevice.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},],};// 创建渲染着色器模块constrenderModuledevice.createShaderModule({label:Example shader module,code:renderCode,});// 定义渲染过程constrenderPipelinedevice.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);// 获取渲染管线绑定组布局constrenderBindGroupLayoutrenderPipeline.getBindGroupLayout(0);// 创建绑定组letrenderBindGroupdevice.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的类型是vec3u32无符号整数如果直接用它计算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×16dispatchWorkgroups(40, 40)则调度了40 × 40个工作组两者相乘正好是640 × 640覆盖了图像的每一个像素、且不多不少。如果图像宽高不是工作组大小的整数倍调度的工作组数量通常需要向上取整例如使用Math.ceil(width / group_size_x)并在着色器中加入边界判断避免对超出图像范围的坐标进行textureLoad或textureStore。