大家好~本文使用WebGPU和光線追蹤算法，從0開始實現和逐步優化Demo，展示了從渲染500個2D物體都吃力到流暢渲染4百萬個2D物體的優化過程和思路

大家好~本文使用WebGPU和光線追蹤算法，從0開始實現和逐步優化Demo，展示了從渲染500個2D物體都吃力到流暢渲染4百萬個2D物體的優化過程和思路

完整代碼在這里

需求

我們對Demo提出下面的要求：

• 渲染1百萬個以上的2D物體時達到30FPS以上
  暫時只渲染一種2D物體：圓環
• 內存占用小
• 剔除被遮擋的物體
• 放大物體時無鋸齒，達到矢量渲染的效果

成果

我們最終能夠流暢渲染4百萬個圓環

性能指標：

• 畫布大小為：width=512, height=512
• 45 FPS左右，也就是每幀花費21毫秒
• 內存占用211MB

硬件：

• Win10操作系統
• Canary瀏覽器
• RTX2060s顯卡

下面讓我們從0開始，介紹實現和優化的步驟：

1、選擇渲染的算法

思考如何實現與渲染相關的需求

現在，我們來回顧下與渲染相關的需求，從而確定渲染算法：

放大物體時無鋸齒，達到矢量渲染的效果

要實現該需求，可以用參數化的方法來表示物體，如用圓心坐標、半徑、圓環寬度 來表示一個圓環；
然后可以根據下面的公式判斷一個點是否在圓環上：

let distanceSquare = Math.pow((點-圓心坐標), 2)

let isPointInRing =  distanceSquare >= Math.pow(半徑, 2) && distanceSquare <= Math.pow(半徑 + 圓環寬度, 2)

因為圓環是2D的，是直接繪制在屏幕上的，所以這里的一個點就是屏幕上的一個像素

從上面的分析可知，通過“如果一個像素在圓環上則渲染”的邏輯，就可以實現需求

剔除被遮擋的物體

要實現該需求，首先圓環加上“層”的數據；
然后遍歷所有圓環，判斷像素在哪些圓環上；
最后取出最大“層”的圓環，將它的顏色作為像素的顏色

選擇渲染的算法
有兩種渲染算法可選擇：光柵化渲染、光追渲染

根據我們之前對渲染需求的思考，我們應該選擇光追渲染作為渲染算法

典型的光追渲染會依次執行下面兩個pass：
1、光追pass
從相機向每個像素發射射線，經過多次射線彈射，計算每個像素的顏色
2、屏幕pass
將屏幕光柵化，渲染每個像素的顏色

其中光追pass是在CS（計算著色器）中進行，屏幕pass是在VS（頂點著色器）和FS（片元著色器）中進行

因為這里是2D物體，所以需要對光追算法做一些簡化：
射線不需要彈射，因為不需要計算間接光照；
因為物體直接就在屏幕上，所以不需要計算射線與物體相交，而是按照參數化公式計算像素是否在2D物體上

2、實現內存需求

現在來考慮如何實現“內存占用小”的需求

這里只考慮CPU端內存，它主要存放了場景數據，包括圓環的transform、geometry、material的數據

我們使用ECS架構，將物體建模為gameObject+components
這里的components具體就是transform、geometry、material這些組件

按照Data Oriented的思想，每種組件就是一個ArrayBuffer，連續地存放組件數據（如transform buffer連續地存放圓環的localPositionX, localPostionY, layer（層）數據）

component為對應buffer中數據的索引

通過ECS的設計，我們就將場景數據都放在buffer中了。在初始化時，只根據最大的組件數量，創建一次對應的組件buffer。這樣就實現了內存占用最小的目標

3、渲染1個圓環

現在來渲染出一個圓環，也不考慮剔除

我們依次實現下面兩個pass：

光追pass

我們需要傳入下面的場景數據到CS中，CS中對應的數據結構如下：

struct AccelerationStructure {
  worldMin : vec2<f32>,
  worldMax : vec2<f32>,

  instanceIndex: f32,
  pad_0: f32,
  pad_1: f32,
  pad_2: f32,
}

struct Instance {
  geometryIndex: f32,
  materialIndex: f32,

  localPosition: vec2<f32>,
}


struct Geometry {
  c: vec2<f32>,
  w: f32,
  r: f32,
}

struct Material {
  color: vec3<f32>,
  pad_0: f32,
}

 struct AccelerationStructures {
  accelerationStructures : array<AccelerationStructure>,
}

 struct Instances {
  instances :  array<Instance>,
}

 struct Geometrys {
  geometrys :  array<Geometry>,
}

 struct Materials {
  materials :  array<Material>,
}

 struct Pixels {
  pixels : array<vec4<f32>>
}

 struct ScreenDimension {
  resolution : vec2<f32>
}

@binding(0) @group(0) var<storage, read> sceneAccelerationStructure :  AccelerationStructures;
@binding(1) @group(0) var<storage, read> sceneInstanceData :  Instances;
@binding(2) @group(0) var<storage, read> sceneGeometryData :  Geometrys;
@binding(3) @group(0) var<storage, read> sceneMaterialData :  Materials;

@binding(4) @group(0) var<storage, read_write> pixelBuffer :  Pixels;

@binding(5) @group(0) var<uniform> screenDimension : ScreenDimension;

struct中的pad用于對齊

sceneAccelerationStructure包含圓環的包圍盒數據，作為加速相交檢測的加速結構。其中worldMin、worldMax為世界坐標系中的min、max

sceneInstanceData、sceneGeometryData、sceneMaterialData為場景中所有的組件數據（sceneInstanceData為Transform數據）

instanceIndex, geometryIndex, materialIndex為對應組件數據的索引

pixelBuffer是光追pass的輸出，用來存放像素顏色

screenDimension存放畫布的width、height

這里為了簡單，沒有傳相機數據

在執行CS時，啟動畫布的width*height個work groups，從而讓每個像素對應一個work group，并行執行每個像素
相關代碼為：

	const passEncoder = commandEncoder.beginComputePass();
	...
	passEncoder.dispatchWorkgroups(width, height);

下面我們來看下CS中計算每個像素顏色的相關代碼：

struct RayPayload {
   radiance: vec3<f32>,
}

struct Ray {
   target: vec2<f32>,
}

struct RingIntersect {
  isHit: bool,
  instanceIndex: f32,
}


fn _isIntersectWithAABB2D(ray: Ray, aabb: AABB2D) -> bool {
	var target = ray.target;
	var min = aabb.min;
	var max = aabb.max;

	return target.x > min.x && target.x < max.x && target.y > min.y && target.y < max.y;
}


fn _isIntersectWithRing(ray: Ray, geometry: Geometry) -> bool {
	var target = ray.target;

	var c = geometry.c;
	var w = geometry.w;
	var r = geometry.r;

	var distanceSquare = pow(target.x - c.x, 2.0) + pow(target.y - c.y, 2.0);

	return distanceSquare >= pow(r, 2) && distanceSquare <= pow(r + w, 2);
}


fn _intersectScene(ray: Ray) -> RingIntersect {
	var intersectResult: RingIntersect;

	intersectResult.isHit = false;

	var as: AccelerationStructure;

	//遍歷所有的圓環，判斷點在哪個圓環上
	for (var i: u32 = 0u; i < arrayLength(& sceneAccelerationStructure.accelerationStructures); i = i + 1u) {
		as = sceneAccelerationStructure.accelerationStructures[i];

		if (_isIntersectWithAABB2D(ray, AABB2D(as.worldMin, as.worldMax))) {
			var instance: Instance = sceneInstanceData.instances[u32(as.instanceIndex)];
			var geometryIndex = u32(instance.geometryIndex);


			var geometry: Geometry = sceneGeometryData.geometrys[geometryIndex];

			if (_isIntersectWithRing(ray, geometry)) {
					//這里沒考慮剔除被遮擋的物體，只是直接取與第一個相交圓環的結果

					intersectResult.isHit = true;
					intersectResult.instanceIndex = as.instanceIndex;

					break;
			}
		}
	}

	return intersectResult;
}

//這里使用了WGSL的pointer type，相當于GLSL的inout，實現對payload的引用修改
fn _handleRayHit(payload: ptr<function,RayPayload>, ray: Ray, intersectResult: RingIntersect)->bool {
  var instance: Instance = sceneInstanceData.instances[u32(intersectResult.instanceIndex)];
  var materialIndex = u32(instance.materialIndex);

  var material:Material = sceneMaterialData.materials[materialIndex];

  (*payload).radiance = material.color;

  return false;
}

fn _handleRayMiss(payload: ptr<function,RayPayload>)->bool {
  (*payload).radiance = vec3<f32>(0.0, 0.0, 0.0);

  return false;
}


fn _traceRay(ray: Ray, payload: ptr<function,RayPayload>)->bool {
  var intersectResult: RingIntersect = _intersectScene(ray);

  if (intersectResult.isHit) {
    return _handleRayHit(payload, ray, intersectResult);
  }

  return _handleRayMiss(payload);
}

@compute @workgroup_size(1, 1, 1)
fn main(@builtin(global_invocation_id) GlobalInvocationID: vec3<u32>) {
  var ipos = vec2<u32>(GlobalInvocationID.x, GlobalInvocationID.y);

  var resolution = vec2<f32>(screenDimension.resolution);

  var pixelColor = vec3<f32>(0.0, 0.0, 0.0);

  //取像素中心為像素坐標
  var sampledPixel = vec2<f32>(f32(ipos.x) + 0.5, f32(ipos.y) + 0.5);
  //獲得屏幕坐標系中的像素坐標
  var uv = (sampledPixel / resolution) * 2.0 - 1.0;


  //像素坐標就是點坐標（因為沒有使用相機，所以它就相當于在世界坐標系中）
  var target = uv;
  /* 如果使用了相機的話，需要將其轉換到世界坐標系中：
  var target = 相機的視圖矩陣的逆矩陣 * 相機的投影矩陣的逆矩陣 * uv;
  */

  var payload: RayPayload;
  payload.radiance = vec3<f32>(0.0, 0.0, 0.0);


  var _isContinueBounce = _traceRay(Ray(target.xy), & payload);


  pixelColor = payload.radiance;

  var pixelIndex = ipos.y * u32(resolution.x) + ipos.x;
  pixelBuffer.pixels[pixelIndex] = vec4<f32>(pixelColor, 1.0);
}

看代碼應該就能理解了吧，就不說明了

屏幕pass

在VS中，使用一個大的三角形包含整個屏幕，代碼為：

struct VertexOutput {
	@builtin(position) Position: vec4 < f32 >,
	@location(0) uv: vec2 < f32 >,  
}

@vertex
fn main(
	@builtin(vertex_index) VertexIndex: u32
) -> VertexOutput {
	var output: VertexOutput;
	output.uv = vec2<f32>(f32((VertexIndex << 1) & 2), f32(VertexIndex & 2));
	output.Position = vec4<f32>(output.uv * 2.0 - 1.0, 0.0, 1.0);

	return output;
}

在FS中，讀取光追pass輸出的像素顏色，代碼為：

struct Pixels {
	pixels: array<vec4<f32>>
}

struct ScreenDimension {
	resolution: vec2<f32>
}


@binding(0) @group(0) var<storage, read_write > pixelBuffer : Pixels;
@binding(1) @group(0) var<uniform>screenDimension : ScreenDimension;

@fragment
fn main(
  @location(0) uv: vec2<f32>
) -> @location(0) vec4 < f32 > {
	var resolution = vec2<f32>(screenDimension.resolution);

	var bufferCoord = vec2<u32>(floor(uv * resolution));
	var pixelIndex = bufferCoord.y * u32(resolution.x) + bufferCoord.x;

	var pixelColor = pixelBuffer.pixels[pixelIndex].rgb;

	return vec4<f32>(pixelColor, 1.0);
}

4、測試渲染極限

現在我們來渲染500個圓環，測試下FPS

渲染結果如下圖所示：

發現性能問題：現在FPS為45，低于60！

我們來分析下CS中的性能熱點：
首先在粗粒度上確定熱點范圍

我們將_intersectScene函數中的for循環注釋掉，再次運行Demo，FPS提升為60

這說明_intersectScene的for循環這塊是性能熱點

我們來分析這塊代碼：
現在每次遍歷都會執行arrayLength函數，而它的性能較差。因此將其提出來只執行一次

代碼修改為：

  var length = arrayLength(&sceneAccelerationStructure.accelerationStructures);
  for (var i : u32 = 0u; i < length; i = i + 1u) {

取消for循環的注釋，再次運行Demo，FPS提升為60

現在for循環會遍歷每個圓環，應該要引入BVH來減少遍歷次數。不過因為BVH實現起來比較復雜，所以不忙引入，而是先考慮容易實現的優化

5、嘗試設置workgroup_size

目前work group的size設為(1,1,1)，沒有啟用局部單位

我們介紹下work group和它的局部單位的概念：

如上圖所示，一個紅色的立方體表示一個work group（通過"passEncoder.dispatchWorkgroups(x1,y1,z1)"來指定啟動(x1乘以y1乘以z1)個work group）； 紅色立方體里面又包含多個小的立方體，每個表示一個局部單位（通過"@workgroup_size(x2,y2,z2)"來指定啟動一個work group中(x2乘以y2乘以z2)個局部單位)

那么總共就有(x1乘以y1乘以z1乘以x2乘以y2乘以z2)次著色器調用，它們都是并行執行的

既然都是并行執行，那么只需要用work group就好了，為什么要用到work group內的局部單位呢？
這是因為下面兩個原因：
1、提供共享變量
同一個work group的局部單位之間可以共享變量
如實現圖像降噪時，需要獲得一個像素周圍3*3范圍的像素來一起計算，它們需要共享一些數據。因此可指定為@workgroup_size(3,3,1)
2、提高性能
因為GPU中有多個計算單位，每個計算單位可以看成是一個work group；每個計算單位中一般可以同時執行8乘以8=64個線程。
如果將work group size設為1，則只使用了計算單位中的1個線程而不是64個線程
所以，一般將work group size設為64（如(64,1,1)或者(8,8,1)），則可以使用計算單位的64個線程，而減少了所需的計算單位數量（變為原來的1/64），這樣就提高了性能

這也是為什么我們要使用WebGPU而不是WebGL來實現本Demo：因為只有WebGPU才有CS，而CS可以充分利用計算單位中的線程，提高性能

啟用局部單位

現在我們將work group數量減少為原來的1/64：

	let workgroup_size = {
		width: 8,
		height: 8
	}

	const passEncoder = commandEncoder.beginComputePass();
	...
	passEncoder.dispatchWorkgroups(Math.ceil(width / workgroup_size.width), Math.ceil(height / workgroup_size.height));

同時將work group的size設為(8,8,1)：

@compute @workgroup_size(8, 8, 1)

但是運行Demo后，發現FPS不僅沒有增加，反而更低了！
這是為什么呢？

通過學習下面的資料：
Optimizing GPU occupancy and resource usage with large thread groups

我們知道了一個計算單位中只有64KB的內存用于存儲（我們只用到了64KB的VGPRs (Vector General-Purpose Registers))，并且由于一個計算單位會并行執行兩個線程組，所以每個線程組只有32KB的內存大小。

然而在_intersectScene的for循環中遍歷了所有的圓環。這意味著會將所有圓環的數據都加載到計算單位的內存中，從而超過了32KB的大小！
這導致了計算單位只能使用1個線程組，并且該線程組只能使用1個而不是所有的線程！
所以不僅相當于沒啟用局部單位，反而可能因為試圖啟用局部單位而造成的同步開銷，導致FPS不升反降！

結論

綜上分析，我們需要引入BVH，大幅減少for循環中需要遍歷的數據，使其小于32KB，從而能夠將其載入到計算單位的內存中；
然后我們再啟用8*8的局部單位！

6、實現BVH

BVH是一個用于空間劃分的樹，相關介紹可參考：
場景管理方法之BVH介紹

GAMES101-現代計算機圖形學入門-閆令琪

我們需要在CPU端構造BVH樹，將其傳入到CS；然后在CS->_intersectScene函數中遍歷BVH

6.1、實現構造BVH

構造BVH樹

我們在CPU端使用最簡單暴力的Middle方法構造BVH樹，步驟如下：

1、計算所有圓環的包圍盒AABB
2、構造根節點，以所有AABB形成的整體AABB為該節點的包圍盒
3、從x軸方向，按照AABB的中心點位置排序所有的AABB
4、以中間的AABB為分界線分割，形成兩個子節點，每個子節點以其包含的所有AABB形成的整體AABB為該子節點的包圍盒
5、遞歸地構造這兩個子節點，并且交替地從y軸方向開始（即x軸->y軸->x軸。。。。。。交替）排序該節點包含的所有AABB和分割。。。。。。直到節點包含的AABB個數<=5或達到最大深度時結束遞歸

構造加速結構

因為CS中只能用數組，所以需要將BVH樹拍平成數組，作為加速結構傳送到CS

我們將加速結構設計為兩層：TopLevel和BottomLevel

BottomLevel的數據類型如下：

type worldMinX = number
type worldMinY = number
type worldMaxX = number
type worldMaxY = number

type instanceIndex = number

type bottomLevel = Array<[worldMinX, worldMinY, worldMaxX, worldMaxY, instanceIndex]>

BottomLevel用來保存所有圓環對應的包圍盒、instanceIndex

TopLevel的數據類型如下：

type wholeWorldMinX = number
type wholeWorldMinY = number
type wholeWorldMaxX = number
type wholeWorldMaxY = number

type leafInstanceOffset = number
type leafInstanceCount = number
type child1Index = number
type child2Index = number

type topLevel = Array<[
	wholeWorldMinX, wholeWorldMinY, wholeWorldMaxX, wholeWorldMaxY,
	leafInstanceOffset,
	leafInstanceCount,
	child1Index,
	child2Index
]>

TopLevel用來保存BVH節點的包圍盒、節點包含的AABB在BottomLevel數組中的索引（如果該節點不是葉節點，則leafInstanceCount=0）、子節點在topLevel數組中的索引

然后將加速結構傳到CS中，CS中對應數據結構如下：

struct TopLevel {
  worldMin : vec2<f32>,
  worldMax : vec2<f32>,

  leafInstanceOffset: f32,
  leafInstanceCount: f32,
  child1Index: f32,
  child2Index: f32
}

struct BottomLevel {
  worldMin : vec2<f32>,
  worldMax : vec2<f32>,

  instanceIndex: f32,
  pad_0: f32,
  pad_1: f32,
  pad_2: f32,
}

 struct TopLevels {
  topLevels : array<TopLevel>,
}

 struct BottomLevels {
  bottomLevels : array<BottomLevel>,
}

@binding(0) @group(0) var<storage> topLevel : TopLevels;
@binding(1) @group(0) var<storage> bottomLevel : BottomLevels;

6.2、CPU端實現遍歷BVH

為了方便測試，我們先在CPU端實現遍歷BVH；解決所有bug后，再移植到CS中

可以分別用遞歸和迭代來實現

考慮到WGSL不支持遞歸函數，所以我們只用迭代來實現
相關代碼如下：

type traverseResult = {
	isHit: boolean,
	instanceIndex: instanceIndex | null
}

let _isPointIntersectWithAABB = (
	point,
	wholeWorldMinX, wholeWorldMinY, wholeWorldMaxX, wholeWorldMaxY,
) => {
	return point[0] > wholeWorldMinX && point[0] < wholeWorldMaxX && point[1] > wholeWorldMinY && point[1] < wholeWorldMaxY
}

let _isPointIntersectWithTopLevelNode = (point, node: topLevelNodeData) => {
	let [
		wholeWorldMinX, wholeWorldMinY, wholeWorldMaxX, wholeWorldMaxY,
		leafInstanceOffset,
		leafInstanceCount,
		child1Index,
		child2Index
	] = node

	return _isPointIntersectWithAABB(
		point,
		wholeWorldMinX, wholeWorldMinY, wholeWorldMaxX, wholeWorldMaxY,
	)
}

let _isLeafNode = (node: topLevelNodeData) => {
	let leafInstanceCountOffset = 5

	return node[leafInstanceCountOffset] !== 0
}

let _handleIntersectWithLeafNode = (intersectResult, isIntersectWithInstance, point, node: topLevelNodeData, bottomLevelArr: bottomLevelArr) => {
	let [
		wholeWorldMinX, wholeWorldMinY, wholeWorldMaxX, wholeWorldMaxY,
		leafInstanceOffset,
		leafInstanceCount,
		child1Index,
		child2Index
	] = node

	while (leafInstanceCount > 0) {
		let [worldMinX, worldMinY, worldMaxX, worldMaxY, instanceIndex] = bottomLevelArr[leafInstanceOffset]

		if (_isPointIntersectWithAABB(
			point,
			worldMinX, worldMinY, worldMaxX, worldMaxY
		)) {
			if (isIntersectWithInstance(point, instanceIndex)) {
				intersectResult.isHit = true
				intersectResult.instanceIndex = instanceIndex

				break;
			}
		}

		leafInstanceCount -= 1
		leafInstanceOffset += 1
	}
}

let _hasChild = (node, childIndexOffset) => {
	return node[childIndexOffset] !== 0
}

export let traverse = (isIntersectWithInstance, point, topLevelArr: topLevelArr, bottomLevelArr: bottomLevelArr): traverseResult => {
	let rootNode = topLevelArr[0]

	let child1IndexOffset = 6
	let child2IndexOffset = 7

	let stackContainer = [rootNode]
	let stackSize = 1

	let intersectResult = {
		isHit: false,
		instanceIndex: null
	}

	while (stackSize > 0) {
		let currentNode = stackContainer[stackSize - 1]

		stackSize -= 1

		if (_isPointIntersectWithTopLevelNode(point, currentNode)) {
			if (_isLeafNode(currentNode)) {
				_handleIntersectWithLeafNode(intersectResult, isIntersectWithInstance, point, currentNode, bottomLevelArr)

				if (intersectResult.isHit) {
					break
				}
			}
			else {
				if (_hasChild(currentNode, child1IndexOffset)) {
					stackContainer[stackSize] = topLevelArr[currentNode[child1IndexOffset]]
					stackSize += 1
				}
				if (_hasChild(currentNode, child2IndexOffset)) {
					stackContainer[stackSize] = topLevelArr[currentNode[child2IndexOffset]]
					stackSize += 1
				}
			}
		}
	}

	return intersectResult
}

這里用了棧(stackContainer)來保存需要遍歷的節點

本來我們可以直接通過stackContainer.push方法將節點push到棧中，但考慮到WGSL的數組沒有push操作，所以這里我們就增加了stackSize這個數據，從而能夠通過"stackContainer[stackSize] = 節點"來代替"stackContainer.push(節點)"

6.3、GPU端實現遍歷BVH

CPU端測試通過后，我們將其移植到CS中

這里值得注意的是因為WGSL創建數組時必須定義大小，所以棧的大小必須預先確定且為常數

棧的大小其實就是BVH樹的最大深度，我們可以先暫時指定為20

相關代碼如下：

fn _intersectScene(ray: Ray)->RingIntersect {
const MAX_DEPTH = 20;

var stackContainer:array<TopLevel, MAX_DEPTH>;

...
}

7、測試渲染極限

現在我們將圓環數量增加200倍，渲染500*200=10萬個圓環，測試下FPS

我們將圓環的半徑和圓環寬度縮小為1/10，這樣方便顯示
渲染結果如下圖所示：

FPS沒有變化

結論

通過引入BVH，渲染性能提高了200倍

8、設置workgroup_size

因為引入了BVH，需要遍歷的節點數大大減少了，所以減少了顯存占用

現在再次啟用局部單位：
我們將work group減少為原來的1/64，將work group的size設為(8,8,1)：

運行Demo后，發現FPS變為60了

理論上可以提高64倍的渲染速度
所以我們將圓環數量增加40倍，渲染10萬*40=4百萬個圓環，測試下FPS
結果FPS跟之前10萬個時一樣

結論

通過啟用局部單位，渲染性能提高了40倍

9、測試內存占用

通過Chrome dev tool->Memory->Take heap snapshot，可以看到包含4百萬個圓環的場景在CPU端只占用了211MB左右的內存，說明內存占用確實小

10、使用LBVH算法來構造BVH

現在當圓環數量為4百萬個時，CPU端構造BVH需要花費100秒以上的時間！

因此，我們保持“構造加速結構”的代碼不變，修改“構造BVH樹”的算法為LBVH算法

它的步驟如下：

1、計算所有圓環的包圍盒AABB
2、構造根節點，以所有AABB形成的AABB為該節點的包圍盒
3、根據根節點的包圍盒，在x、y軸方向上將其1024等分，根據AABB的中心點在哪個區域而計算出AABB在x、y軸方向上的格子坐標
4、將格子坐標轉換為Morton Code
5、根據Morton Code將所有的AABB排序
6、對于該有序Morton Code數組，我們利用二分查找出第一個不同的bit位（也即是從0變為1的index），此時我們即可將最高位為0的所有BVs歸入此node（此時是root節點）的左子樹，最高位為1的所有BVs歸入右子樹；同理，我們對左右子樹按照下一個bit位來遞歸的處理，直到遞歸的處理完全部bit位，LBVH即可建立完畢

最后一步如下圖所示：

結論

通過改為LBVH算法，構造時間降低為10秒左右，性能提高了10倍

之所以LBVH算法更快，是因為只排序了一次！

11、實現剔除

之前分析過剔除的實現思路：

首先圓環加上“層”的數據；
然后遍歷所有圓環，判斷像素在哪些圓環上；
最后取出最大“層”的圓環，將它的顏色作為像素的顏色

因此，可以在圓環的transform組件數據中增加layer數據（從1開始的正整數），用來表示“層”；
然后在構造加速結構BottomLevel時，讀取transform組件中的layer數據，將其保存到BottomLevel數據中；
最后修改CS代碼，在遍歷BVH->檢測到像素在圓環上時比較layer，并且不再停止遍歷，而是繼續遍歷棧的其它節點；在遍歷棧的其它節點時，如果找到了像素在其上的圓環，也不再停止

實現剔除后，運行Demo渲染1百萬個圓環時FPS都才15（之前是4百萬個 45 FPS），渲染性能估計降低了10倍

12、改進遍歷BVH

我們從下面幾個方面對剔除進行優化

traverse order優化

現在如果找到了像素在其上的圓環，會繼續遍歷其它節點。
我們需要減少遍歷的節點數量

我們可以在構造BVH樹時，為每個節點增加maxLayer數據，它為該節點包含的所有的圓環中最大的層；
然后在遍歷BVH時：
檢測到像素在圓環上時，將圓環的layer記錄到相交結果intersectResult中；
在遍歷棧中的節點的while循環中，如果該節點的maxLayer <= intersectResult.layer，則說明該節點包含的所有圓環都被遮擋了，直接continue，跳過；
另外，在遍歷葉節點的所有圓環時，如果像素所在的圓環的layer==該節點的maxLayer，則說明已經找到了葉節點包含的所有圓環中的最大層的圓環，則break，停止搜索該葉節點包含的其它圓環。

通過上面的優化，可以大幅降低遍歷的節點數量

合并數據

因為BVH樹的節點需要保存maxLayer數據，而這個數據實際上是保存在TopLevel中的，對應到CS中的數據結構就是：

struct TopLevel {
	worldMin : vec2<f32>,
	worldMax : vec2<f32>,

	leafInstanceOffset: f32,
	leafInstanceCount: f32,
	child1Index: f32,
	child2Index: f32

	maxLayer: f32,
	pad_0: f32,
	pad_1: f32,
	pad_2: f32,
}

我們可以看到，因為增加了maxLayer，需要增加3個pad數據來對齊，這樣浪費了顯存

而占用盡可能少的顯存是非常重要的，因為經過之前的分析，我們知道一個計算單位只有32KB的顯存可用，超過的話就會導致啟用局部單位失效！

仔細觀察后，我們發現leafInstanceCount和maxLayer只需要占用<32位的字節數
我們可以讓這兩個數據各占16位，其中leafInstanceCount在高位，maxLayer在低位；然后將其合成一個32位f32
從而TopLevel修改為：

struct TopLevel {
	worldMin : vec2<f32>,
	worldMax : vec2<f32>,

	leafInstanceOffset: f32,
	leafInstanceCountAndMaxLayer: f32,
	child1Index: f32,
	child2Index: f32
}

這樣就消除了pad數據，減少了顯存占用

但是當渲染的圓環數量超過1百萬個時，會出現“從leafInstanceCountAndMaxLayer中取出的maxLayer為0（應該>=1）”的bug！
這是因為leafInstanceCount（葉節點的圓環個數）過大，占用了超出了16位的字節數，從而影響到maxLayer的值！

所以我們重新分配，讓leafInstanceCount占24位，maxLayer占8位，則解決了bug

注：理論上16位的leafInstanceCount可以最大為65535，但實際上我發現當leafInstanceCount>1000時，就出現了超過16位的情況！我估計是WGLSL可能占用了f32類型的數據的最高幾位，導致leafInstanceCount實際可用位數<16位

13、測試渲染極限

現在我們將圓環數量恢復為4百萬個圓環，FPS又恢復為45左右

當我們嘗試渲染5百萬個圓環時，遇到了“我們BottomLevel Buffer數據超出了Storage Buffer的最大限制：128MB”的問題

關于這個限制，WebGPU官方有相關的討論issue：
Limit for the maximum buffer size

繞過該限制的可能方案是將其拆成多個Storage Buffer

結論

通過traverse order優化，渲染性能提高了10倍左右

當我們嘗試渲染5百萬個圓環時，遇到了超出Storage Buffer最大大小的限制

總結

感謝大家的學習~

在本文中，我們先提出了需求；然后按照需求來設計和選擇算法；然后實現最簡單的版本；接著不斷優化，直到達到Storage Buffer的最大大小限制為止

我們的優化的成果為：

• 通過引入BVH，渲染性能提高了200倍
• 通過啟用局部單位，渲染性能提高了40倍
• 通過改為LBVH算法，構造BVH的性能提高了10倍
• 通過traverse order優化，使得剔除的性能提高了10倍左右

目前我們最多渲染4百萬個圓環，因為再多就會超出Storage Buffer最大大小的限制

后續的改進方向

后面我們希望能夠渲染千萬級2D物體，可以從下面的方向改進：

• 將加速結構拆成多個Storage Buffer
• 優化構造BVH，使葉節點包含的圓環數量盡量少，且節點的包圍盒盡量不重疊，這樣才能提高遍歷BVH的性能
  可考慮使用HLBVH算法
• 優化遍歷BVH：考慮并行遍歷BVH、無棧的遍歷
  參考Ray Tracing學習之Traversal

參考資料

OpenGL4.3 新特性: 計算著色器 Compute Shader
Bad preformance of simple ray trace compute shader
Optimizing GPU occupancy and resource usage with large thread groups
我所理解的DirectX Ray Tracing
并行構建BVH
Build LBVH on GPUs
Ray Tracing學習之Traversal
光線求交加速算法：邊界體積層次結構(Bounding Volume Hierarchies)3-LBVH(Linear Bounding Volume Hierarchies)
WebGPU 計算管線、計算著色器（通用計算）入門案例：2D 物理模擬

歡迎瀏覽下一篇博文：如何用WebGPU流暢渲染千萬級2D物體：基于光追管線