Skip to content

Commit 2e2153e

Browse files
vismaychuriwalavismaychuriwala
committed
Adding clustered rendering (Forward+)
1 parent 3d02d67 commit 2e2153e

7 files changed

Lines changed: 486 additions & 6 deletions

File tree

‎src/renderers/forward_plus.ts‎

Lines changed: 128 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,15 +6,143 @@ export class ForwardPlusRenderer extends renderer.Renderer {
66
// TODO-2: add layouts, pipelines, textures, etc. needed for Forward+ here
77
// you may need extra uniforms such as the camera view matrix and the canvas resolution
88

9+
sceneBindGroupLayout: GPUBindGroupLayout;
10+
sceneBindGroup: GPUBindGroup;
11+
12+
depthTexture: GPUTexture;
13+
depthTextureView: GPUTextureView;
14+
15+
pipeline: GPURenderPipeline;
916
constructor(stage: Stage) {
1017
super(stage);
1118

1219
// TODO-2: initialize layouts, pipelines, textures, etc. needed for Forward+ here
20+
this.sceneBindGroupLayout = renderer.device.createBindGroupLayout({
21+
label: "forward+ scene bgl",
22+
entries: [
23+
{
24+
// camera uniforms
25+
binding: 0,
26+
visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT,
27+
buffer: { type: "uniform" },
28+
},
29+
{
30+
// light set
31+
binding: 1,
32+
visibility: GPUShaderStage.FRAGMENT,
33+
buffer: { type: "read-only-storage" },
34+
},
35+
{
36+
// cluster grid
37+
binding: 2,
38+
visibility: GPUShaderStage.FRAGMENT,
39+
buffer: { type: "read-only-storage" },
40+
},
41+
{
42+
// cluster light indices
43+
binding: 3,
44+
visibility: GPUShaderStage.FRAGMENT,
45+
buffer: { type: "read-only-storage" },
46+
},
47+
],
48+
});
49+
this.sceneBindGroup = renderer.device.createBindGroup({
50+
label: "forward+ scene bg",
51+
layout: this.sceneBindGroupLayout,
52+
entries: [
53+
{ binding: 0, resource: { buffer: this.camera.uniformsBuffer } },
54+
{ binding: 1, resource: { buffer: this.lights.lightSetStorageBuffer } },
55+
{ binding: 2, resource: { buffer: this.lights.clusterLightGridBuffer } },
56+
{
57+
binding: 3,
58+
resource: { buffer: this.lights.clusterLightIndexBuffer },
59+
}
60+
],
61+
});
62+
this.depthTexture = renderer.device.createTexture({
63+
size: [renderer.canvas.width, renderer.canvas.height],
64+
format: "depth24plus",
65+
usage: GPUTextureUsage.RENDER_ATTACHMENT
66+
});
67+
this.depthTextureView = this.depthTexture.createView();
68+
this.pipeline = renderer.device.createRenderPipeline({
69+
layout: renderer.device.createPipelineLayout({
70+
label: "Forward+ pipeline layout",
71+
bindGroupLayouts: [
72+
this.sceneBindGroupLayout,
73+
renderer.modelBindGroupLayout,
74+
renderer.materialBindGroupLayout
75+
]
76+
}),
77+
depthStencil: {
78+
depthWriteEnabled: true,
79+
depthCompare: "less",
80+
format: "depth24plus"
81+
},
82+
vertex: {
83+
module: renderer.device.createShaderModule({
84+
label: "Forward+ shader (same as naive vert)",
85+
code: shaders.naiveVertSrc
86+
}),
87+
buffers: [ renderer.vertexBufferLayout ]
88+
},
89+
fragment: {
90+
module: renderer.device.createShaderModule({
91+
label: "naive frag shader",
92+
code: shaders.forwardPlusFragSrc,
93+
}),
94+
targets: [
95+
{
96+
format: renderer.canvasFormat,
97+
}
98+
]
99+
}
100+
});
101+
13102
}
14103

15104
override draw() {
16105
// TODO-2: run the Forward+ rendering pass:
17106
// - run the clustering compute shader
18107
// - run the main rendering pass, using the computed clusters for efficient lighting
108+
const encoder = renderer.device.createCommandEncoder();
109+
const canvasTextureView = renderer.context.getCurrentTexture().createView();
110+
this.lights.doLightClustering(encoder);
111+
112+
113+
const renderPass = encoder.beginRenderPass({
114+
label: "forward+ render pass",
115+
colorAttachments: [
116+
{
117+
view: canvasTextureView,
118+
clearValue: [0, 0, 0, 0],
119+
loadOp: "clear",
120+
storeOp: "store"
121+
}
122+
],
123+
depthStencilAttachment: {
124+
view: this.depthTextureView,
125+
depthClearValue: 1.0,
126+
depthLoadOp: "clear",
127+
depthStoreOp: "store"
128+
}
129+
});
130+
renderPass.setPipeline(this.pipeline);
131+
132+
renderPass.setBindGroup(shaders.constants.bindGroup_scene, this.sceneBindGroup)
133+
134+
this.scene.iterate(node => {
135+
renderPass.setBindGroup(shaders.constants.bindGroup_model, node.modelBindGroup);
136+
}, material => {
137+
renderPass.setBindGroup(shaders.constants.bindGroup_material, material.materialBindGroup);
138+
}, primitive => {
139+
renderPass.setVertexBuffer(0, primitive.vertexBuffer);
140+
renderPass.setIndexBuffer(primitive.indexBuffer, 'uint32');
141+
renderPass.drawIndexed(primitive.numIndices);
142+
});
143+
144+
renderPass.end();
145+
146+
renderer.device.queue.submit([encoder.finish()]);
19147
}
20148
}

‎src/shaders/clustering.cs.wgsl‎

Lines changed: 141 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,3 +21,144 @@
2121
// - Stop adding lights if the maximum number of lights is reached.
2222

2323
// - Store the number of lights assigned to this cluster.
24+
25+
const X_SLICES: u32 = ${xSlices};
26+
const Y_SLICES: u32 = ${ySlices};
27+
const Z_SLICES: u32 = ${zSlices};
28+
const MAX_LIGHTS_PER_CLUSTER: u32 = ${maxLightsPerCluster};
29+
30+
// Workgroup size
31+
const WG_SIZE_X: u32 = ${clusterWorkgroupSizeX};
32+
const WG_SIZE_Y: u32 = ${clusterWorkgroupSizeY};
33+
34+
const LIGHT_RADIUS: f32 = f32(${lightRadius});
35+
const LIGHT_RADIUS2: f32 = LIGHT_RADIUS * LIGHT_RADIUS;
36+
37+
struct AABB {
38+
min: vec3f,
39+
max: vec3f,
40+
};
41+
42+
@group(0) @binding(0) var<uniform> camera_uniforms: CameraUniforms;
43+
@group(0) @binding(1) var<storage, read> lightSet: LightSet;
44+
@group(0) @binding(2) var<storage, read_write> clusterGrid: array<ClusterRecord>;
45+
@group(0) @binding(3) var<storage, read_write> clusterLightIndices: array<u32>;
46+
47+
fn lerp(a: f32, b: f32, t: f32) -> f32 {
48+
return a + (b - a) * t;
49+
}
50+
51+
fn edge_to_ndc(idx: u32, total: u32) -> f32 {
52+
return -1.0 + 2.0 * (f32(idx) / f32(total));
53+
}
54+
55+
fn cluster_index(ix: u32, iy: u32, iz: u32) -> u32 {
56+
return ix + iy * X_SLICES + iz * X_SLICES * Y_SLICES;
57+
}
58+
59+
fn dist2_point_aabb(p: vec3f, bmin: vec3f, bmax: vec3f) -> f32 {
60+
var dx = 0.0;
61+
if (p.x < bmin.x) {
62+
dx = bmin.x - p.x;
63+
} else if (p.x > bmax.x) {
64+
dx = p.x - bmax.x;
65+
}
66+
67+
var dy = 0.0;
68+
if (p.y < bmin.y) {
69+
dy = bmin.y - p.y;
70+
} else if (p.y > bmax.y) {
71+
dy = p.y - bmax.y;
72+
}
73+
74+
var dz = 0.0;
75+
if (p.z < bmin.z) {
76+
dz = bmin.z - p.z;
77+
} else if (p.z > bmax.z) {
78+
dz = p.z - bmax.z;
79+
}
80+
81+
return dx * dx + dy * dy + dz * dz;
82+
}
83+
84+
fn compute_cluster_bounds(ix: u32, iy: u32, iz: u32) -> AABB {
85+
let tanY = camera_uniforms.tanHalfFovY;
86+
let tanX = tanY * camera_uniforms.aspect;
87+
88+
let t0 = f32(iz) / f32(Z_SLICES);
89+
let t1 = f32(iz + 1u) / f32(Z_SLICES);
90+
let z_near_d = lerp(camera_uniforms.nearPlane, camera_uniforms.farPlane, t0);
91+
let z_far_d = lerp(camera_uniforms.nearPlane, camera_uniforms.farPlane, t1);
92+
93+
let z0 = -z_near_d;
94+
let z1 = -z_far_d;
95+
96+
let x0_ndc = edge_to_ndc(ix, X_SLICES);
97+
let x1_ndc = edge_to_ndc(ix + 1u, X_SLICES);
98+
let y0_ndc = edge_to_ndc(iy, Y_SLICES);
99+
let y1_ndc = edge_to_ndc(iy + 1u, Y_SLICES);
100+
101+
let depth_pos_far = z_far_d; // = -z1
102+
let x0_far = x0_ndc * depth_pos_far * tanX;
103+
let x1_far = x1_ndc * depth_pos_far * tanX;
104+
let y0_far = y0_ndc * depth_pos_far * tanY;
105+
let y1_far = y1_ndc * depth_pos_far * tanY;
106+
107+
let x_min = min(x0_far, x1_far);
108+
let x_max = max(x0_far, x1_far);
109+
let y_min = min(y0_far, y1_far);
110+
let y_max = max(y0_far, y1_far);
111+
112+
let z_min = min(z0, z1);
113+
let z_max = max(z0, z1);
114+
115+
return AABB(vec3f(x_min, y_min, z_min), vec3f(x_max, y_max, z_max));
116+
}
117+
118+
fn light_view_pos(world_pos: vec3f) -> vec3f {
119+
let viewPos4 = camera_uniforms.viewMat * vec4f(world_pos, 1.0);
120+
return viewPos4.xyz;
121+
}
122+
123+
@compute @workgroup_size(WG_SIZE_X, WG_SIZE_Y, 1)
124+
fn main(@builtin(global_invocation_id) gid: vec3u) {
125+
if (gid.x >= X_SLICES || gid.y >= Y_SLICES || gid.z >= Z_SLICES) {
126+
return;
127+
}
128+
129+
let idx = cluster_index(gid.x, gid.y, gid.z);
130+
let base = idx * MAX_LIGHTS_PER_CLUSTER;
131+
132+
let tanY = camera_uniforms.tanHalfFovY;
133+
let tanX = tanY * camera_uniforms.aspect;
134+
135+
// Compute cluster view-space AABB
136+
let bounds = compute_cluster_bounds(gid.x, gid.y, gid.z);
137+
let bmin = bounds.min;
138+
let bmax = bounds.max;
139+
140+
// Fill list
141+
var count: u32 = 0u;
142+
143+
for (var i: u32 = 0u; i < lightSet.numLights; i = i + 1u) {
144+
// World-space -> view-space
145+
let p_view = light_view_pos(lightSet.lights[i].pos);
146+
147+
// Quick reject on Z (optional)
148+
if (p_view.z < bmin.z - LIGHT_RADIUS || p_view.z > bmax.z + LIGHT_RADIUS) {
149+
continue;
150+
}
151+
152+
let d2 = dist2_point_aabb(p_view, bmin, bmax);
153+
if (d2 <= LIGHT_RADIUS2) {
154+
if (count < MAX_LIGHTS_PER_CLUSTER) {
155+
clusterLightIndices[base + count] = i;
156+
count = count + 1u;
157+
}
158+
}
159+
}
160+
161+
// Store offset and count
162+
clusterGrid[idx].offset = base;
163+
clusterGrid[idx].count = count;
164+
}

‎src/shaders/common.wgsl‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,9 +14,19 @@ struct LightSet {
1414

1515
struct CameraUniforms {
1616
// TODO-1.3: add an entry for the view proj mat (of type mat4x4f)
17+
viewMat: mat4x4<f32>,
1718
viewProjMat: mat4x4f,
19+
nearPlane : f32,
20+
farPlane : f32,
21+
tanHalfFovY : f32,
22+
aspect : f32,
1823
}
1924

25+
struct ClusterRecord {
26+
offset: u32,
27+
count: u32,
28+
};
29+
2030
// CHECKITOUT: this special attenuation function ensures lights don't affect geometry outside the maximum light radius
2131
fn rangeAttenuation(distance: f32) -> f32 {
2232
return clamp(1.f - pow(distance / ${lightRadius}, 4.f), 0.f, 1.f) / (distance * distance);

0 commit comments

Comments
 (0)