[图形渲染]讲透RenderTarget 第九章：Compute Shader 与 RenderTarget

第九章Compute Shader 与 RenderTarget一句话概括不走光栅化也能画——Compute Shader 通过 UAV/Storage Image 直接写纹理。生活类比不用画笔光栅化直接用手指蘸颜料点上去Compute 逐像素写入。⏱ 30 秒概览Compute Shader 不走光栅化管线通过UAVDX/ Storage ImageVulkan/GL随机读写纹理任意位置。HLSL 用RWTexture2DGLSL 用imageStore/imageLoad。Compute 的优势可用Shared MemoryLDS缓存邻域像素降噪/模糊、可一次 Dispatch 生成多级 Mip、可做 Scatter 写入和 Atomic 操作。劣势没有硬件混合、没有 MSAA。写入后必须插入 BarrierCache Flush 执行依赖确保后续 Pass 能安全读取。UAV AtomicInterlockedAdd/CAS打开了新维度OIT 链表构建、体素化原子计数。选择 Fragment 还是 Compute 取决于是否需要 LDS/混合/MSAA——没有绝对优劣需 Profile。9.1 不走光栅化管线也能画——UAV / Storage Image前八章讨论的所有 RT 写入本质上都走同一条路顶点着色 → 光栅化 → 片段着色 → ROP → RT这条路的核心角色是光栅化器——它决定每个三角形覆盖哪些像素然后对每个像素执行片段着色器。最终由 ROP 把结果写入 RenderTarget。但 Compute Shader 打破了这个范式。它不需要三角形、不需要光栅化、不需要 ROP。它可以直接把数据写到任意像素位置。在 DX11/12 中这个能力叫UAVUnordered Access View——对资源的无序读写视图。在 Vulkan/OpenGL 中对应概念叫Storage Image/Image Load Store。它们的本质是一样的让 Shader 可以随机读写一张纹理或 Buffer)的任意位置而不受光栅化管线的约束。Compute Shader → UAV / Storage Image → 纹理≈ RenderTarget这里有一个概念上的微妙之处Compute Shader 写入的纹理还算RenderTarget吗严格说RenderTarget是图形管线的术语——绑定到 ROP 输出端、用OMSetRenderTargets或vkCmdBeginRendering设置的。Compute Shader 写入的是 UAV/Storage Image不走 ROP不算传统意义的 RT。但在实践中这些纹理和 RT 经常混用Compute Shader 写入一张纹理 → 下一个 Pass 当作 RT 的输入Fragment Shader 画到一张 RT → 下一个 Pass 用 Compute Shader 读取并写到另一张 UAV → 再下一个 Pass 当 RT 的输入所以从数据流的角度看UAV/Storage Image 就是 RT 生态的一部分。9.2 imageStoreGLSL/ RWTexture2DHLSLHLSL——RWTexture2D// 声明一个可读写的 2D 纹理 RWTexture2Dfloat4 outputTexture : register(u0); [numthreads(8, 8, 1)] void CSMain(uint3 id : SV_DispatchThreadID) { // id.xy 就是当前线程对应的像素坐标 float2 uv float2(id.xy) / float2(width, height); // 从另一张只读纹理采样 float4 color inputTexture.SampleLevel(linearSampler, uv, 0); // 做一些后处理 color.rgb pow(color.rgb, 2.2); // gamma 校正 // 直接写入目标位置 outputTexture[id.xy] color; }关键点RWTexture2Dfloat4绑定到 UAV slotu0通过[id.xy]像数组一样直接写入也可以读取float4 old outputTexture[id.xy]支持的格式有限——不是所有纹理格式都支持 UAV Typed Load/StoreGLSL——imageStore / imageLoadlayout(binding 0, rgba16f) uniform image2D outputImage; layout(binding 1) uniform sampler2D inputTexture; layout(local_size_x 8, local_size_y 8) in; void main() { ivec2 pixelCoord ivec2(gl_GlobalInvocationID.xy); vec2 uv vec2(pixelCoord) / vec2(imageSize(outputImage)); vec4 color texture(inputTexture, uv); color.rgb pow(color.rgb, vec3(2.2)); imageStore(outputImage, pixelCoord, color); }关键点image2D对应 Vulkan 的 Storage Imagelayout(rgba16f)必须声明格式Vulkan 中也可以通过shaderStorageImageWriteWithoutFormat特性省略imageStore(image, coord, data)写入imageLoad(image, coord)读取Metalkernel void csMain( texture2dfloat, access::read inTex [[texture(0)]], texture2dfloat, access::write outTex [[texture(1)]], uint2 gid [[thread_position_in_grid]]) { float4 color inTex.read(gid); color.rgb pow(color.rgb, 2.2); outTex.write(color, gid); }Metal 中用access::write/access::read_write声明写入权限。UAV Typed Load 的格式限制不是所有格式都支持 UAV Typed Load通过[]读取。DX11 时代只有R32_FLOAT/UINT/SINT三种格式保证支持 Typed UAV Load。DX12 和现代 GPU 扩展了支持列表但仍有限制格式DX11 UAV LoadDX12 UAV LoadVulkan Storage ImageR32_FLOAT✅✅✅R32_UINT✅✅✅RGBA8_UNORM❌✅需检查 Feature Support✅需检查格式能力RGBA16F❌✅部分 GPU✅需检查格式能力R11G11B10F❌✅少数 GPU❌多数 GPU 不支持写入Store的格式支持比读取Load更广泛——如果只写不读限制会少些。9.3 何时用 Compute 替代 Fragment Shader 写 RT这是一个工程上的关键决策。两者都能把数据写到纹理上该用哪个Fragment Shader全屏三角形/四边形的特点通过光栅化管线输出直接走 ROP可以利用硬件的Early-Z如果绑定了 Depth输出像素天然按光栅化顺序排列Cache 友好自动处理Sample CountMSAA混合Blending由 ROP 硬件完成——免费每个像素只执行一次——不会重复计算Compute Shader 的特点不走光栅化线程组Thread Group在 GPU 上自由调度可以使用Shared MemoryLDS/Groupshared——线程组内共享数据可以随机写入任意位置——Fragment Shader 只能写自己那个像素可以做Scatter 写入——一个线程写多个位置可以在同一 Dispatch 中读写同一张纹理的不同区域需要 Barrier没有内置的混合——需要手动实现决策指南场景推荐原因简单全屏后处理Bloom、ToneMapFragment硬件自带混合、Cache 友好需要 Shared Memory 的操作降噪、模糊ComputeLDS 可缓存邻域像素减少重复采样降采样链Mip GenerationCompute一次 Dispatch 生成多级 Mip比多遍 Draw 高效Scatter 写入粒子着色ComputeFragment Shader 无法随机写需要 Atomic 操作OIT、VoxelizeCompute 或 Fragment UAV需要 UAV 才能 AtomicTAA Temporal Filter两者皆可Compute 稍占优可用 LDS 缓存历史帧MSAA 渲染FragmentCompute 不走光栅化没有硬件 MSAA性能陷阱Quad OccupancyFragment Shader 有一个隐藏代价——Quad Overdraw。GPU 以 2×2 像素为最小单位Quad执行 Fragment Shader。如果三角形很小比如只覆盖 1 个像素Quad 中其他 3 个像素的计算就浪费了。对于全屏三角形Quad 效率接近 100%。但如果你在做 G-Buffer Lighting Pass全屏 QuadCompute Shader 的 8×8 线程组也接近 100%——两者差不多。真正的性能差异取决于是否需要 Shared MemoryCompute 大优势是否需要硬件混合Fragment 大优势是否有 MSAAFragment 专属Cache 访问模式取决于具体算法结论没有绝对的优劣需要 Profile。多数引擎对同一效果提供 Fragment 和 Compute 两条路径。9.4 同步与 BarrierCompute 写完后怎么安全读取Compute Shader 写入 UAV/Storage Image 后其他 Pass 要读取这张纹理中间必须插入 Barrier——告诉 GPU “写操作已完成接下来的读操作可以开始了”。为什么需要 BarrierGPU 是高度并行的。Compute Shader 的写入可能还在 L2 Cache 中没刷回显存而下一个 Fragment Shader 已经试图从显存读这张纹理了——读到的是旧数据。Barrier 做两件事Flush Cache确保写入数据对所有后续读者可见Execution Dependency确保所有写操作在读操作开始前完成各 API 的 BarrierDX12// Compute 写入完成后CD3DX12_RESOURCE_BARRIER barrierCD3DX12_RESOURCE_BARRIER::Transition(texture.Get(),D3D12_RESOURCE_STATE_UNORDERED_ACCESS,// Compute 写入状态D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE// Fragment Shader 读取);commandList-ResourceBarrier(1,barrier);VulkanVkImageMemoryBarrier barrier{};barrier.oldLayoutVK_IMAGE_LAYOUT_GENERAL;// Compute 写入barrier.newLayoutVK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL;// 后续采样barrier.srcAccessMaskVK_ACCESS_SHADER_WRITE_BIT;barrier.dstAccessMaskVK_ACCESS_SHADER_READ_BIT;vkCmdPipelineBarrier(commandBuffer,VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT,// 等 Compute 完成VK_PIPELINE_STAGE_FRAGMENT_SHADER_BIT,// 才开始 Fragment0,0,nullptr,0,nullptr,1,barrier);Metal// Metal 中 Compute Encoder 和 Render Encoder 之间自动有隐式 Barrier// 但如果同一个 Compute Encoder 中读写同一张纹理需要computeEncoder.memoryBarrier(resources:[texture])OpenGL// Compute Shader 写入后 glMemoryBarrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT); // 或者如果后续是纹理采样 glMemoryBarrier(GL_TEXTURE_FETCH_BARRIER_BIT);Barrier 粒度的选择Barrier 越精细GPU 可以做的并行优化越多粗粒度简单但慢 等所有 Compute 都完成 细粒度复杂但快 只等 Compute 对这张特定纹理的写入完成DX12 和 Vulkan 支持细粒度 Barrier指定具体资源和 Subresource。OpenGL 的glMemoryBarrier是全局性的粒度较粗。Compute ↔ Compute 之间的 Barrier如果两个连续的 Compute Pass 读写同一张纹理比如 Ping-Pong 模糊也需要 BarrierCompute Pass A: 读 Texture0, 写 Texture1 ← BarrierTexture1: UAV Write → UAV Read Compute Pass B: 读 Texture1, 写 Texture09.5 Compute RT 的典型场景场景 1全屏后处理最常见的用途。用 Compute Shader 替代全屏三角形做后处理[Color RT] → Compute ShaderToneMapping 色彩校正→ [Final RT]优点可以用 Shared Memory 缓存邻域像素减少采样次数。缺点失去硬件混合能力。场景 2降噪Denoising光线追踪产生的 noisy 图像需要空间-时间降噪。降噪算法如 SVGF、ReLAX通常需要读取大邻域5×5 或更大跨帧累积读写历史 RT大量加权计算Compute Shader Shared Memory 是最佳选择。场景 3Mipmap 生成一次 Dispatch 可以读取上一级 Mip、计算下一级 Mip、然后在线程组内继续计算再下一级。用 Fragment Shader 需要多遍 Draw Call。// 伪代码一次 Dispatch 生成 4 级 Mip groupshared float4 shared_mip[8][8]; [numthreads(8, 8, 1)] void MipGen(uint3 id : SV_DispatchThreadID, uint3 gid : SV_GroupThreadID) { // Level 0 → Level 1 shared_mip[gid.y][gid.x] SrcMip0[id.xy * 2]; // 简化实际需要 4 点平均 GroupMemoryBarrierWithGroupSync(); if (gid.x 4 gid.y 4) { // Level 1 → Level 2 float4 avg average of 2×2 block in shared_mip; DstMip1[...] avg; shared_mip[gid.y][gid.x] avg; } GroupMemoryBarrierWithGroupSync(); // ... 继续 Level 2 → Level 3, Level 3 → Level 4 }场景 4Temporal 累积TAA / Temporal FilterTAA 的核心是当前帧和历史帧的混合。Compute Shader 可以用 Shared Memory 缓存 Motion Vector 的邻域用于 Clamp高效处理历史帧的 Rejection[当前 Color RT] [Motion Vector RT] [历史 Color RT] → Compute ShaderTemporal Filter → [新历史 Color RT]场景 5Hi-ZHierarchical Z-Buffer生成SSR 和遮挡剔除需要的 Hi-Z Buffer 是深度纹理的 Mipmap 链每级取最大/最小深度。本质上是 Mipmap 生成的变种Compute Shader 同样适合。9.6 UAV 上的 Atomic 操作——OIT、Voxelization 等结构化写入与普通颜色写入的本质区别普通 RT 写入是只写——ROP 把颜色混合进 RT每个像素的写入是确定性的深度测试决定谁赢。但有些算法需要在同一个像素位置累积多个值而且这些值来自不同的三角形不同的 Draw Call 甚至不同的线程。这就需要Atomic 操作——原子地读-改-写同一个内存位置不会丢失任何笔数据。UAV 上支持的 Atomic 操作InterlockedAdd/atomicAdd原子加InterlockedMin/atomicMin原子取最小InterlockedMax/atomicMax原子取最大InterlockedCompareExchange/atomicCompSwapCASCompare and SwapInterlockedOr/InterlockedAnd/InterlockedXor原子位操作这些操作的关键特性多个线程同时执行时结果是正确的不会出现写丢失race condition。OITOrder Independent Transparency的链表构建OIT 是 Atomic UAV 的经典应用。问题透明物体需要从后往前排序绘制。如果场景复杂排序代价很高甚至不可能精确排序。方案Per-Pixel Linked List——为每个像素维护一个链表存储所有覆盖该像素的透明片段最后一次性排序并混合。数据结构RWTexture2Duint headPointerImage; // 每像素一个头指针UAV RWStructuredBufferFragment fragmentList; // 全局片段链表UAV Buffer RWByteAddressBuffer counter; // 全局原子计数器Fragment Shader 写入过程struct Fragment { float4 color; float depth; uint next; }; void PSMain(PSInput input) { Fragment frag; frag.color computeColor(input); frag.depth input.position.z; // 1. 原子递增计数器拿到新的链表节点索引 uint newIndex; counter.InterlockedAdd(0, 1, newIndex); // 2. 原子交换头指针——把自己插入链表头部 uint oldHead; InterlockedExchange(headPointerImage[pixelCoord], newIndex, oldHead); // 3. 把旧的头指针存为自己的 next frag.next oldHead; // 4. 写入链表 fragmentList[newIndex] frag; }Resolve PassCompute 或全屏 Quad// 读取当前像素的链表 uint index headPointerImage[pixelCoord]; // 遍历链表收集所有片段 // 按深度排序 // 从后往前混合 // 输出最终颜色这里的InterlockedExchange和InterlockedAdd就是 Atomic 操作——确保几百个三角形的片段能正确地插入同一个像素的链表而不丢失。注意整个过程没有用到传统的 RT 混合Blending而是用 UAV Atomic 手动实现了一个更强大的混合逻辑。体素化Voxelization的原子计数实时全局光照如 VXGI需要把场景体素化——把三角形转换为 3D 体素网格。体素化过程中多个三角形可能覆盖同一个体素。需要把每个三角形的颜色/法线贡献累积到对应体素RWTexture3Duint voxelGrid; // 3D UAVuint 格式方便 Atomic void VSMain(Vertex v) { // ... 投影到体素空间 } void PSMain(PSInput input) { uint3 voxelCoord worldToVoxel(input.worldPos); // 把颜色打包为 uint uint packedColor packColor(input.color); // 原子地累积颜色 InterlockedAdd(voxelGrid[voxelCoord], packedColor); // 或更复杂的用 CAS 做加权平均 }这里的RWTexture3D就是一个 3D 的 UAV对应于 3D RT 的概念但不走 ROP。何时需要 Typed UAV / Structured Buffer 而非传统 RT需求传统 RTROPTyped UAVRWTexture2DStructured BufferRWStructuredBuffer简单颜色写入✅ 最佳可以但没必要不适合硬件混合✅ 免费❌ 需手动实现❌ 需手动实现Atomic 操作❌ RT 不支持 Atomic✅ 支持有限 Atomic✅ 支持全部 Atomic链表/自定义结构❌ RT 只能写固定格式❌ 只能写像素格式✅ 任意结构体随机地址写入❌ 只能写当前像素✅ 可以✅ 可以3D 纹理写入❌不走光栅化✅ RWTexture3D用 Buffer 模拟经验法则如果你只需要画到一张 2D 图上→ 传统 RT如果你需要 Atomic 或 Scatter 写入 → UAVTyped 或 Structured如果你需要存复杂结构体如链表节点→ Structured Buffer如果后续需要当纹理采样 → Typed UAVRWTexture2D比 Buffer 更方便本章小结Compute Shader 通过 UAV/Storage Image 直接写纹理——不走光栅化、不经过 ROP但输出的纹理和传统 RT 在数据流中地位相同Fragment vs Compute 没有绝对优劣——需要 Shared Memory 用 Compute需要硬件混合或 MSAA 用 Fragment写入后必须插入 Barrier——确保 Cache 刷新和执行顺序Atomic 操作打开了 RT 写入的新维度——OIT 链表、体素化、无序透明混合等算法依赖 UAV Atomic传统 RTROP、Typed UAV、Structured Buffer 各有适用场景——根据是否需要混合、Atomic、自定义结构选择设计哲学ROP 与 UAV——从管线写到自由写的范式迁移 传统光栅化管线的 RT 写入是线性流水线模型顶点 → 光栅化 → 片段 → ROP → RT。每个像素只能写到自己对应的位置不能越界。这种约束很严格但正因为严格GPU 才能做深度流水线优化乱序执行、Early-Z、Tile-Based 合并等。Compute Shader UAV 打破了这种约束——你可以写任意位置、做原子操作、访问邻域像素。自由度换来了表达力但也失去了 GPU 能替你做的隐式优化混合、MSAA、Tile 合并。这和 CPU 编程中高级语言 vs 汇编的取舍完全类似抽象层级越低能力越强但需要自己管理的细节越多。未来的方向不是Compute 取代 Fragment而是混合管线——简单的逐像素操作用光栅化享受硬件优化邻域计算/稀疏写入/原子操作用 Compute享受灵活性。光线追踪的 UAV 输出第 14.3 节更进一步验证了这个趋势。 思考题为什么 Compute Shader 不能利用硬件 MSAA从 GPU 架构角度解释。如果 GPU 未来支持Compute Shader 中的硬件混合需要怎样的硬件变化有没有现有技术的雏形OIT顺序无关透明为什么需要 Atomic 操作如果没有 Atomic能否用纯光栅化管线实现同样效果下一章讲引擎如何帮你管理这些 RT——Pool、Handle、Frame Graph。当一帧中有 20~50 张 RT 交叉读写时手动管理生命周期和 Barrier 已超出人脑并发上限。你将看到引擎是怎么把人管对象升级为图管对象以及这条路径和编译器优化的惊人相似性。