基于WebGPU的gpu-driven rendering探索

1,141 阅读18分钟

前言:作为一个使用了多年 WebGL 的前端引擎开发人员,在得知 WebGPU 标准制定和发布的时候是带着兴奋的心情的。不只是早已厌倦了用了多年的(状态机风格的) GL 系 API ,而且还能将最新游戏引擎所使用的现代图形管线的优势和特点带给到 web,可谓是带给了 web 端图形渲染带来了更广阔的可能性。本文通过对 webgpu API 的初探来尝试验证 gpu-driven pipeline 的可能性,希望能够抛砖引玉,带来一些思考和讨论。

什么是 gpu-driven rendering

随着硬件的发展,我们可以看到GPU的运算能力远远的把CPU抛在后面,实际上现代的3d应用可以设法利用GPU强大的算力来计算、控制、决定绘制的流程,在传统应用中这些都是由CPU控制的,GPU只能作为指令的接收和执行方。

在经典的引擎渲染管线中,场景物体的剔除、DrawCall发起都由CPU来完成,CPU和图形驱动承担了很多指令生成、状态验证等工作,只有最终的指令执行、光栅化、渲染是由GPU完成的。既然最终的渲染数据来自GPU显存,DrawCall的最终执行也在GPU上完成,并且compute shader也是可以访问到这部分内存的,那么可否将更多的前置工作也转移至GPU?GPU能否完成裁剪、剔除和发起DrawCall的工作?这个就是gpu-driven rendering的思想。

总体来看,gpu-driven包含两个部分:

  • 剔除 - 将原本CPU端的剔除工作交给GPU去做,不仅节约了CPU端的算力,还可以利用GPU compute进行更激进的剔除和裁剪

  • DrawCall - 基于GPU自身计算的结果,使用GPU直接控制绘制指令的发起或执行,解放CPU端重复的DrawCall指令生成和验证开销

  • 以上的两步可以做的非常精细和复杂,例如Siggraph15上面Ubi分享的刺客信条大革命案例,实现了相当极致的gpu-driven:

    • 将Mesh的几何数据重新组织&拆分,形成Clusters,每个Cluster固定包含64个三角形

    • 渲染前经历不同粒度的剔除阶段:CPU负责粗粒度的剔除,GPU负责Cluster级别&更加精细的剔除,Cluster、三角形级别的剔除,HiZ遮挡剔除等等

    • GPU决定使用何种LOD级别的几何数据

    • 有了固定顶点和索引数量的Cluster,那么每个DrawCall所需的数据就可以从一个大的VertexBuffer和IndexBuffer中根据索引取得,这样理论上可以将整个场景用一次DrawCall就绘制出来。但是这个方式需要在渲染之前对模型数据做非常大的改动,拆分不定数量的Clusters,并且还要考虑不同材质和不同的纹理、uniforms应当如何存储和获得

在本文的demo中演示的是相对简单的一种gpu-driven rendering,没有将所有的DrawCall视为相同VertexCount和IndexCount,而是保持了每个DrawCall的独立,以及每个DrawCall所对应的不同材质和渲染参数等。实现了使用GPU进行物体级别的culling,以剔除完全不在视椎体内的物体DrawCall。配合WebGPU的RenderBundle特性,也可以实现对整个场景中的不透明物体仅用一次绘制指令 + 1~2次GPU compute就可以绘制出来。

WebGPU中的gpu-driven相关特性

  • RenderBundle - 将一系列渲染指令(set material, set uniforms, set VB, set IB, draw, drawIndexed, ...)提前录制好,生成一个command bundle,这个bundle在后续的渲染中可以直接被重放,实现降低每帧CPU上的重复指令验证和生成的开销。该特性运用得当可以极大减少CPU生成渲染指令的overhead,缺点是缺乏灵活性,bundle一旦生成便无法修改,除非重新录制;

  • DrawIndirect / DrawIndexedIndirect 间接绘制 - 间接绘制是gpu-driven模式的核心API,它的作用是将以往从CPU发起的draw指令的参数存储方式变成显存内的buffer,因此DrawCall执行时的实际参数值来自于渲染时GPU从buffer内读取的值,该指令的好处是可以在GPU中决定渲染时的参数以及是否发起真正的渲染,所以命名为间接(indirect)绘制,通常使用compute shader来计算实际渲染参数并写入到buffer;

  • Compute Pipeline (Compute Shader) 计算管线 - webgpu支持了compute pipeline,可以利用gpu的算力执行一些更加通用的计算和修改buffer等操作,例如gpu剔除等,从而解放cpu端的算力;

一个 web gpu-driven demo 分享

demo共有两种模式,以及一些可开关的选项:

  • 普通渲染模式:可开启/关闭视锥剔除、可开启/关闭间接绘制

  • gpu-driven(bundleRender)渲染模式:可开启/关闭视锥剔除,不可关闭间接绘制

  • 不启用任何instancing绘制方法

  • 可观测数据:

NUM:物体总数量(裁剪+未裁剪)

drawCount:实际执行的draw指令数量(裁剪后)

computeCount:实际执行的compute pipeline次数

jsTime:每帧中js执行的时长(ms),可代表CPU执行时间

drawTime:每帧中js耗费在生成和执行图形API指令上的时长(ms)

开启视锥剔除(使用首帧相机参数):

关闭视锥剔除:

  • 测试配置

普通模式

  • 渲染物体时使用普通的 draw 和 drawIndexed 接口

  • 采用GPU更新所有物体的MVP(减少CPU的计算时间干扰)

  • 采用CPU计算视锥剔除,剔除结果直接影响CPU发起的DrawCall数量

  • 为了实现实验的单一变量比对,额外设置了使用间接绘制接口的开关

  • 主渲染流程代码:

    function drawNormalPass( device: GPUDevice, context: GPUCanvasContext, pipelineObj: PipelineObj, frustum: Frustum, time: number, modelData: Float32Array, // 视锥剔除计算所需的模型数据 ) { const { pipeline, depthView, modelVB, modelIB, vpBindGroup, bindGroups, indirectBuffer, indirectData, updatePipeline, updateBindGroup, } = pipelineObj;

    const commandEncoder = device.createCommandEncoder()
    
    // 1. 使用compute shader更新物体的位置和MVP、BoundingBox
    const computePass = commandEncoder.beginComputePass()
    computePass.setPipeline(updatePipeline)
    computePass.setBindGroup(0, updateBindGroup) // Buffer: modelData + time + duration
    computePass.dispatchWorkgroups(Math.ceil(NUM / 128))
    computePass.end()
    infoRef.computeCount++
    
    const t1 = performance.now()
    
    // 2. 启动渲染Pass
    const renderPassDescriptor: GPURenderPassDescriptor = {
        colorAttachments: [
            {
                view: context.getCurrentTexture().createView(),
                clearValue: { r: 0, g: 0, b: 0, a: 1.0 },
                loadOp: 'clear',
                storeOp: 'store'
            }
        ],
        depthStencilAttachment: {
            view: depthView,
            depthClearValue: 1.0,
            depthLoadOp: 'clear',
            depthStoreOp: 'store',
        }
    }
    const passEncoder = commandEncoder.beginRenderPass(renderPassDescriptor)
    passEncoder.setPipeline(pipeline)
    passEncoder.setBindGroup(0, vpBindGroup) // 相机View、Projection矩阵
    if (infoRef.indirectDraw) {
        // 根据CPU剔除的结果,更新 indirect buffer 数据
        let offset = 0
        if (infoRef.culling) {
            for (let i = 0; i < NUM; i++) {
                if (frustumCulling(frustum, time, modelData, i)) {
                    indirectData[offset + 1] = 1
                } else {
                    indirectData[offset + 1] = 0
                }
                offset += 5
            }
        } else {
            for (let i = 0; i < NUM; i++) {
                indirectData[offset + 1] = 1
                offset += 5
            }
        }
        // 写入更新后的 indirect buffer data
        device.queue.writeBuffer(indirectBuffer, 0, indirectData, 0)
    
        // 3. 使用间接渲染绘制
        offset = 0
        const step = 5 * Uint32Array.BYTES_PER_ELEMENT
        for (let i = 0; i < NUM; i++) {
            if (indirectData[5 * i + 1] > 0) {
                passEncoder.setVertexBuffer(0, modelVB)
                passEncoder.setIndexBuffer(modelIB, "uint16")
                passEncoder.setBindGroup(1, bindGroups[i])
                passEncoder.drawIndexedIndirect(indirectBuffer, offset)
                infoRef.drawCount++
            }
            offset += step
        }
    } else {
        // 3. 使用普通渲染绘制
        if (infoRef.culling) {
            for (let i = 0; i < NUM; i++) {
                if (frustumCulling(frustum, time, modelData, i)) {
                    passEncoder.setVertexBuffer(0, modelVB)
                    passEncoder.setIndexBuffer(modelIB, "uint16")
                    passEncoder.setBindGroup(1, bindGroups[i])
                    passEncoder.drawIndexed(model.indexCount, 1)
                    infoRef.drawCount++
                }
            }
        } else {
            for (let i = 0; i < NUM; i++) {
                passEncoder.setVertexBuffer(0, modelVB)
                passEncoder.setIndexBuffer(modelIB, "uint16")
                passEncoder.setBindGroup(1, bindGroups[i])
                passEncoder.drawIndexed(model.indexCount, 1)
                infoRef.drawCount++
            }
        }
    }
    
    // 4. 结束RenderPass,提交指令
    passEncoder.end()
    device.queue.submit([commandEncoder.finish()])
    
    infoRef.drawTime = (performance.now() - t1).toFixed(3)
    

    }

gpu-driven 模式

  • 渲染物体时使用 drawIndirect 和 drawIndexedIndirect 接口

  • 采用GPU更新所有物体的MVP(减少CPU的计算时间干扰)

  • 采用GPU计算视锥剔除,根据剔除的结果,更新 indirect buffer 内的绘制参数

  • 首帧渲染时开启RenderBundle录制,后续渲染时直接重放录制好的RenderBundle

  • 主渲染流程代码:

    function drawRenderBundlePass( device: GPUDevice, context: GPUCanvasContext, format: GPUTextureFormat, pipelineObj: PipelineObj, frustum: Frustum // 视锥对象 ) { const { pipeline, depthView, modelVB, modelIB, vpBindGroup, bindGroups, indirectBuffer, updatePipeline, updateBindGroup, cullingPipeline, cullingBindGroup, frustumBuffer, } = pipelineObj;

    // 使用 render bundle 必须开启间接渲染
    if (!infoRef.indirectDraw) {
        infoRef.indirectDraw = true
    }
    
    const commandEncoder = device.createCommandEncoder()
    
    // 1. 使用compute shader更新物体的位置和MVP、BoundingBox
    const computePass = commandEncoder.beginComputePass()
    computePass.setPipeline(updatePipeline)
    computePass.setBindGroup(0, updateBindGroup) // Buffer: modelData + time + duration
    computePass.dispatchWorkgroups(Math.ceil(NUM / 128))
    infoRef.computeCount++
    
    // 2. 使用compute shader进行视锥剔除和更新 indirect buffer
    if (infoRef.culling) {
        // 更新GPU端视锥体数据
        device.queue.writeBuffer(frustumBuffer, 0, frustum.array)
        computePass.setPipeline(cullingPipeline)
        computePass.setBindGroup(0, cullingBindGroup) // Buffer: modelData + frustum + indirectData
        computePass.dispatchWorkgroups(Math.ceil(NUM / 128))
        infoRef.computeCount++
    }
    computePass.end()
    
    const t1 = performance.now()
    
    // 3. 初次渲染时录制 renderBundle
    if (!renderBundles) {
        const bundleEncoder = device.createRenderBundleEncoder({
            colorFormats: [format],
            depthStencilFormat: 'depth24plus'
        })
        bundleEncoder.setPipeline(pipeline)
        bundleEncoder.setBindGroup(0, vpBindGroup)
        for (let i = 0; i < NUM; i++) {
            bundleEncoder.setVertexBuffer(0, modelVB)
            bundleEncoder.setIndexBuffer(modelIB, "uint16")
            bundleEncoder.setBindGroup(1, bindGroups[i])
            bundleEncoder.drawIndexedIndirect(indirectBuffer, 5 * i * Uint32Array.BYTES_PER_ELEMENT)
            infoRef.drawCount++
        }
        renderBundles = [bundleEncoder.finish()]
    }
    
    // 4. 重放录制的renderBundle
    const renderPassDescriptor: GPURenderPassDescriptor = {
        colorAttachments: [
            {
                view: context.getCurrentTexture().createView(),
                clearValue: { r: 0, g: 0, b: 0, a: 1.0 },
                loadOp: 'clear',
                storeOp: 'store'
            }
        ],
        depthStencilAttachment: {
            view: depthView,
            depthClearValue: 1.0,
            depthLoadOp: 'clear',
            depthStoreOp: 'store',
        }
    }
    const passEncoder = commandEncoder.beginRenderPass(renderPassDescriptor)
    passEncoder.executeBundles(renderBundles)
    passEncoder.end()
    
    // 5. 结束RenderPass,提交指令
    device.queue.submit([commandEncoder.finish()])
    
    infoRef.drawTime = (performance.now() - t1).toFixed(3)
    infoRef.drawCount++
    

    }

物体位置信息更新shader(WGSL):

包括读取velocity,更新position、boundingBox,写入更新后数据到GRAM

const NUM = u32($$NUM$$);

@group(0) @binding(0) var<storage, read_write> modelData: array<f32>;
@group(0) @binding(1) var<uniform> time: u32;
@group(0) @binding(2) var<uniform> duration: u32;

const groupSize = u32(128);
@compute @workgroup_size(groupSize)
fn main(
    @builtin(global_invocation_id) GlobalInvocationID : vec3<u32>
) {
    var index: u32 = GlobalInvocationID.x;
    if(index >= NUM){
        return;
    }

    var offset = index * 64;

    var pos = vec3<f32>(modelData[offset + 12], modelData[offset + 13], modelData[offset + 14]);
    var boxMin = vec3<f32>(modelData[offset + 16], modelData[offset + 17], modelData[offset + 18]);
    var boxMax = vec3<f32>(modelData[offset + 19], modelData[offset + 20], modelData[offset + 21]);
    var vel = vec3<f32>(modelData[offset + 22], modelData[offset + 23], modelData[offset + 24]);

    var round = time / duration;
    var scale = 1.0;
    if (round % 2 == 0) {
        // forward
        scale = 1.0;
    } else {
        // backward
        scale = -1.0;
    }

    vel *= vec3(scale);

    pos.x += vel.x;
    pos.y += vel.y;
    pos.z += vel.z;

    boxMin.x += vel.x;
    boxMin.y += vel.y;
    boxMin.z += vel.z;

    boxMax.x += vel.x;
    boxMax.y += vel.y;
    boxMax.z += vel.z;

    modelData[offset + 12] = pos.x;
    modelData[offset + 13] = pos.y;
    modelData[offset + 14] = pos.z;

    modelData[offset + 16] = boxMin.x;
    modelData[offset + 17] = boxMin.y;
    modelData[offset + 18] = boxMin.z;

    modelData[offset + 19] = boxMax.x;
    modelData[offset + 20] = boxMax.y;
    modelData[offset + 21] = boxMax.z;
}

GPU culling shader:

包括读取每个物体的boundingBox,视椎体检测以及写入draw instance count(写入0代表跳过这个物体的绘制)

const NUM = u32($$NUM$$);

struct Frustum {
    p0: vec4<f32>,
    p1: vec4<f32>,
    p2: vec4<f32>,
    p3: vec4<f32>,
    p4: vec4<f32>,
    p5: vec4<f32>,
}

@group(0) @binding(0) var<storage, read_write> modelData: array<f32>;
@group(0) @binding(1) var<uniform> frustum: Frustum;
@group(0) @binding(2) var<storage, read_write> indirectData: array<u32>;

const groupSize = u32(128);
@compute @workgroup_size(groupSize)
fn main(
    @builtin(global_invocation_id) GlobalInvocationID : vec3<u32>
) {
    var index: u32 = GlobalInvocationID.x;
    if(index >= NUM){
        return;
    }

    var offset = index * 64;

    var boxMin = vec3<f32>(modelData[offset + 16], modelData[offset + 17], modelData[offset + 18]);
    var boxMax = vec3<f32>(modelData[offset + 19], modelData[offset + 20], modelData[offset + 21]);

    if (frustumIntersectsBox(frustum, boxMin, boxMax) > 0) {
        indirectData[5 * index + 1] = 1; // instance count set to 1, draw object
    } else {
        indirectData[5 * index + 1] = 0; // instance count set to 0, skip draw
    }
}

两种渲染模式对比

  1. 在CPU成为瓶颈的情况下,关闭culling,两种模式的GPU端执行DrawCall都为60000个,比较相同DrawCall数量下使用RenderBundle的效果

普通模式:

帧数大约25fps,每帧CPU时间为62.7ms左右,js执行执行渲染指令的耗时为49ms左右,GPU实际渲染时间为39ms左右

gpu-driven模式:

帧数大约52-54fps,每帧CPU耗时0.6ms左右,js执行渲染指令的耗时为0.1ms左右,GPU实际渲染时间为37ms左右

  1. 在CPU成为瓶颈的情况下,开启culling,普通模式使用直接绘制接口draw/drawIndexed,对比gpu-driven模式下使用compute shader计算culling和控制drawCall

普通模式:

帧率43fps左右,由于开启culling,实际drawCall数量减少到不到20000次(1/3),CPU端执行时间减少到34ms左右(原62.7),js执行执行渲染指令的耗时为21ms左右(原49),GPU实际渲染时间为18.36ms左右(原39)

gpu-driven模式:

帧率58fps左右,开启culling后虽然无法在CPU端看到实际drawCall数量,但是可以估计和普通模式应该一致(culling算法一样),CPU端执行时间为0.55ms左右(原0.6),js执行执行渲染指令的耗时为0.1ms左右(原0.1),GPU实际渲染时间为31.09ms左右(原37)

  • 普通模式下由于culling的作用减少了drawCall数量,得到了直接的CPU和GPU执行时间的降低;

  • 开启gpu-driven模式后,CPU端js执行时间可以忽略不计,但是GPU时间比普通模式更多,原因在于增加的一次compute shader culling的计算时间从CPU端转移到了GPU端,使得gpu-driven模式下的GPU耗时比普通模式要多13ms左右,但即便如此增加的GPU culling运算时间也小于CPU culling的计算时间;

  • 在CPU成为瓶颈的情况下,使用gpu-driven模式渲染对帧率的提升是显著的,不仅可以大幅减少CPU端的耗时,在GPU端增加的workload也相对正常和合理;

  • 当indirect buffer中的instance count数量被设置为0时,即使CPU端发起了渲染指令,在GPU端也会被跳过,从两次gpu-driven模式的GPU执行时间差异可以看出;

  1. 在GPU成为瓶颈的情况下(换成顶点更多面数更多的模型),比较两种模式对性能的影响

普通模式:35-38fps,CPU执行时间16ms,GPU执行时间38ms

gpu-driven模式:帧率30-32fps,CPU执行时间0.67ms,GPU执行时间47ms

  • 在GPU成为瓶颈的情况下,使用gpu-driven模式对帧率没有帮助,甚至因为compute shader culling占用了更多的GPU算力而使得帧率达不到普通模式,但是同样也极大减少了CPU端的图形指令耗时;

  • 使用compute shader culling的GPU端耗时和前两次对比基本一致,30000个物体culling的耗时约为9ms;

draw与drawIndirect指令耗时对比

使用drawIndirect:CPU指令耗时51ms,GPU耗时56ms,帧率21fps

使用draw:CPU指令耗时41ms,GPU耗时45ms,帧率25fps

关闭culling(每帧60000次完整DrawCall),使用draw比使用drawIndirect的帧数高2-4fps,CPU绘制指令时间draw比drawIndirect快10.6ms,平均下来6000次draw和drawIndirect的CPU端指令时间会相差1ms,GPU执行时间差(60000次DrawCall)在1ms范围之内。

gpu-driven rendering 的应用初探

  • 通过上一章节的数据可知,gpu-driven对释放CPU端的图形指令压力和耗时具有较大的帮助,这比较契合web端应用的诉求,通常web端应用由于js运行效率以及其他部分对算力的占用,瓶颈会集中在CPU端,此时若能够利用好gpu-driven的渲染方式则可以比较好的提升web3d应用的性能;

  • drawIndirect的CPU端性能损耗在10000次以下DrawCall级别时为1ms左右,也在可接受的范围内;

简单版

  • 通过人为将一系列的物体渲染指令打包到一起,在渲染时重放事先录制好的RenderBundles

  • 不同渲染物体的matrices、uniforms(BindGroups)、VertexBuffer、IndexBuffer独立存储,对应独立的GPUBuffer,因此不能使用compute shader进行统一的数据更新或culling

  • 此版本的实现最简单,因无需对物体的数据结构、材质、参数等做额外的处理,只是将多个渲染指令bundle在一起,但是无法做到控制不同物体的渲染与否(对应引擎中mesh的visibility设置),每次改变了渲染指令或渲染顺序则必须重新录制整个bundle

进阶版(demo版本)

  • 与demo的实现类似,首先还是将一批物体的渲染指令实现录制成RenderBundles

  • 将不同渲染物体的matrices、bbox等与空间相关的信息pack到一个大的GPUBuffer中(demo中的modelData buffer),录制每个物体的渲染指令时,通过给BindGroup(uniforms)设置偏移来绑定这些数据用于渲染

    webgpu

    规范中规定BindGroup的偏移值不能小于256bytes,因此如果每个物体的数据少于256bytes则会存在一定的buffer空间浪费

  • 可利用compute shader,在GPU中直接读取和修改modelData buffer,达到利用GPU算力来执行model数据更新、视锥剔除、遮挡剔除等工作

  • 通过在GPU端用cs写入indirect buffer(同样使用偏移方式)来控制未通过裁剪和剔除的物体的渲染与否,GPU也会自动根据indirect参数来决定是否跳过物体的draw

  • 该版本有一定的灵活性,可以在不重新录制RenderBundle的情况下,通过GPU控制每个DrawCall的执行或跳过,但是无法调整物体的渲染顺序,因为不同物体的材质、BindGroups、VB、IB的绑定顺序是录制在Bundle内的,即使调换了draw的顺序,也无法调换这些绑定和状态的顺序

终极版(UE5/寒霜引擎等商业引擎版本)

在现代的引擎中,为了减少GPU在像素着色器的重复绘制,一般会每帧对物体的绘制顺序进行排序,一个典型的不透明物体绘制排序算法:

  1. 首先根据物体的renderOrder排序

  2. 相同renderOrder的物体根据材质排序

  3. 相同材质的物体根据所使用的几何数据排序

  4. 相同的几何数据物体根据与相机的距离进行排序

其中a-c的排序结果在场景物体没有增减或变化的情况下是固定的,但是d的排序基本上是每帧都会改变,此时渲染的顺序就无法事先确定,若使用gpu-driven则需要使用cs来计算每个(a-c都相同的)物体到相机的距离,随后在indirect buffer中调换相应DrawCall的参数。

如何在GPU端落地这种排序方法?

希望能用尽可能少的compute shader次数来完成所有DrawCall的排序,因此如果a-c步产生的不同种类的DrawCall越多,就会需要很多次的cs调用来对每种DrawCall分别排序,按照一个UE生成的大型游戏材质shader的种类可能会超过10000,这样数量的不同DrawCall种类就会失去gpu-driven的意义

那么是否可以减少不同种类的DrawCall?

  • 如果能够将模型数据、渲染数据(状态)、shader这些渲染所需的必要条件进行统一,保证每次DrawCall的所需条件能够只根据物体的编号(或数据索引、偏移等)而得到,那么就可以在cs中更进一步,达到控制DrawCall渲染数据、渲染顺序的效果;

  • 【几何】如第一章介绍,一些现代的游戏引擎已经实现将场景内所有物体的几何信息分解成Clusters,因此可以将场景的所有几何信息pack到一个大的geometry buffer,每个DrawCall获取顶点和索引的方式都从buffer中通过固定偏移(cluster包含固定64个三角形)来获取,此外还能获得更加细致的culling和遮挡剔除效果,在UE5中这种方式就是Virtual Geometry(Nanite)

刺客信条大革命中的几何结构 Mesh-Chunk-Cluster

  • 【材质Shader】方面,如果可以将整个场景的材质减少到个位数(甚至只有一种),也可以大幅减少渲染的种类,因此现代游戏引擎已经逐步切换到使用Uber材质,一种材质内实现场景中几乎所有的渲染效果,结合Deferred Shading 延迟渲染,可以做到整个场景只用到5个以下的材质(包括sprite、particles等特殊材质),大幅减少需要测距和排序的cs次数;

  • 【渲染参数】uniforms(BindGroups)和纹理绑定方面,通过减少全局材质的种类,可以将数值数据类uniforms也pack到一个大的buffer中,通过固定的偏移值在渲染时实时取值,这样可以在cs中根据DrawCall的排序进行uniforms数据的调换;纹理方面,现代游戏引擎例如UE5使用了Virtual Texture技术,类似于从头至尾只用一张texture或一个TextureArray,将整个场景当前用到的像素值stream到这个texture上,并在cs中同步更新每个cluster用到的uv,达到整个场景渲染只需绑定一次纹理;DX12还提供了更加激进的BindlessTexture特性使得渲染材质时无需再绑定任何纹理,这些技术我们后续再进行深入探索;

Ubi RedLynx引擎中的Virtual Texture技术

  • 通过以上这些技术,统一了几何、材质、渲染参数后基本上整个场景的绘制可以做到不切换VB、IB,只切换几次材质shader和Uniforms,最终通过少数几次的cs指令(一次完整并行排序算法也需要多次cs调用),实现在gpu中控制DrawCall的剔除+排序;

上面的每一项技术的实现都包含了很多具体的细节,基本上每个都可以扩展成多篇技术文章,这里只是非常粗略的将这些技术列举出来,很多大型游戏引擎因为一些历史原因或者代码框架改动幅度的原因,也并没有一开始就用上完全的gpu-driven rendering,很多引擎都从地形、建筑、树木草地等容易改造的模块开始逐步使用Virtual Texture、Virtual Geometry,随后逐步应用了gpu-driven renderin。在刺客信条大革命中也未一开始就把材质缩减到个位数,而是根据不同材质hash来分批渲染。UE一直到UE5才在Nanite上将整个几何系统改造成Virtual Geometry,带来的收益也相当可观,其本质不仅是减少DrawCall和CPU overhead,而是彻底摆脱传统的几何、材质、贴图等的限制,一次DrawCall解决整个场景的绘制,实现渲染超大或无上限的虚拟世界的目标。

适合Web3D的gpu-driven方式设想

  • MeshCluster、VT、VG的方式对于传统的引擎管线改造非常大,并且这种模式需要较多的渲染前处理工作,例如分解几何数据、集合所有的贴图数据,这些前置处理如果在runtime时候做会比较耗时,适合游戏引擎的编译期或数据打包期间执行。并且这种管线对动态增添和删除Mesh(entity)不太友好,比较适合渲染管线非常固定的大型游戏或3d应用;

  • 在初期可以使用本例子中的进阶版方式,由于这个方式无法进行排序,pixel的overdraw会相对较高,可以同时使用Deferred Shading来减少每个像素的重复着色;

  • 将场景中的待渲染物体归类为static、dynamic两类,static物体对应几何数据结构不变(数据可以变)、材质和shader不变、Uniforms(BindGroups)结构不变,将这些物体(不透明与半透明分开)打包到一起后,可以使用本demo中的RenderBundle方式进行渲染,可以同时使用cs来进行物体级别的culling和遮挡剔除。

    • 对于不透明物体,使用RenderBundle和cs culling之后会造成渲染顺序与摄像机的距离无法成比例,但是依靠z buffer和deferred shading依然可以得到渲染正确的结果;

    • 对于半透明物体,可以被分类为dynamic然后使用普通模式每帧排序后渲染,也可以分类为static后使用Order Independent Transparency (OIT) 技术来实现渲染顺序无关的半透明效果;

    • 对于static的物体前期也可以依靠CPU来进行culling,将culling结果从CPU写入到GPUm

  • 后期可在引擎内逐步实现VG和VT等特性,为完全版gpu-driven打下基础,直到可以将场景数据和渲染数据都以统一buffer和streaming的方式给到GPU,此时可以将整个渲染管线切换成gpu-driven的完全版

References

www.w3.org/TR/webgpu/

zhuanlan.zhihu.com/p/409244895

zhuanlan.zhihu.com/p/37084925

zhuanlan.zhihu.com/p/376267968

zhuanlan.zhihu.com/p/138484024

zhuanlan.zhihu.com/p/368065919

aaltonenhaar_siggraph2015_combined_final_footer_220dpi.pdf