@@ -60,6 +60,9 @@ struct Uniforms {
6060
6161var<workgroup> sharedCenterScreen: array<vec2f, 64>;
6262var<workgroup> sharedCoeffs: array<vec3f, 64>;
63+ #ifdef HEATMAP_MODE
64+ var<workgroup> sharedHeatCount: atomic<u32>;
65+ #endif
6366
6467// Pick mode stores per-splat opacity, ID and depth; color mode stores packed RGBA.
6568// Depth test mode also needs per-splat view depth for occlusion against scene geometry.
@@ -104,6 +107,20 @@ fn evalSplatPick(pixelCoord: vec2f, center: vec2f, coeffX: f32, coeffY: f32, coe
104107}
105108#endif
106109
110+ #ifdef HEATMAP_MODE
111+ fn heatmapColor(v: f32) -> vec3f {
112+ let t = saturate(v / 2000.0);
113+ if (t < 0.2) {
114+ return mix(vec3f(0.0, 0.0, 1.0), vec3f(0.0, 1.0, 1.0), t * 5.0);
115+ } else if (t < 0.4) {
116+ return mix(vec3f(0.0, 1.0, 1.0), vec3f(1.0, 1.0, 0.0), (t - 0.2) * 5.0);
117+ } else if (t < 0.6) {
118+ return mix(vec3f(1.0, 1.0, 0.0), vec3f(1.0, 0.0, 0.0), (t - 0.4) * 5.0);
119+ }
120+ return mix(vec3f(1.0, 0.0, 0.0), vec3f(0.15, 0.0, 0.0), (t - 0.6) * 2.5);
121+ }
122+ #endif
123+
107124@compute @workgroup_size(8, 8)
108125fn main(
109126 @builtin(local_invocation_id) lid: vec3u,
@@ -179,6 +196,12 @@ fn main(
179196
180197 let tileCount = tEnd - tStart;
181198
199+ #ifdef HEATMAP_MODE
200+ if (localIdx == 0u) { atomicStore(&sharedHeatCount, 0u); }
201+ workgroupBarrier();
202+ var processedCount: u32 = 0u;
203+ #endif
204+
182205 let numBatches = (tileCount + BATCH_SIZE - 1u) / BATCH_SIZE;
183206 var threadDone = false;
184207
@@ -288,6 +311,10 @@ fn main(
288311 c11 += splatColor.rgb * weight.w;
289312 T = select(T, newT, valid);
290313
314+ #ifdef HEATMAP_MODE
315+ processedCount += 1u;
316+ #endif
317+
291318 if (all(T < half4(ALPHA_THRESHOLD))) {
292319 threadDone = true;
293320 break;
@@ -299,44 +326,65 @@ fn main(
299326 workgroupBarrier();
300327 }
301328
302- // Write results for the 2x2 pixel quad owned by this thread.
303- // Pick mode: store the front-most pick ID and (accumulated depth, weight) per pixel.
304- // Color mode: convert accumulated gamma-space color to linear via decodeGamma3 and store
305- // to the rgba16float output texture; alpha holds total opacity (1 - transmittance).
306- if (basePixel.x < uniforms.screenWidth && basePixel.y < uniforms.screenHeight) {
307- #ifdef PICK_MODE
308- textureStore(pickIdTexture, basePixel, vec4u(pickId00, 0u, 0u, 0u));
309- textureStore(pickDepthTexture, basePixel, vec4f(dAcc00, wAcc00, 0.0, 0.0));
310- #else
311- textureStore(outputTexture, basePixel, vec4f(decodeGamma3(vec3f(c00)), f32(half(1.0) - T.x)));
312- #endif
313- }
314- if (basePixel.x + 1u < uniforms.screenWidth && basePixel.y < uniforms.screenHeight) {
315- let px10 = vec2u(basePixel.x + 1u, basePixel.y);
316- #ifdef PICK_MODE
317- textureStore(pickIdTexture, px10, vec4u(pickId10, 0u, 0u, 0u));
318- textureStore(pickDepthTexture, px10, vec4f(dAcc10, wAcc10, 0.0, 0.0));
319- #else
320- textureStore(outputTexture, px10, vec4f(decodeGamma3(vec3f(c10)), f32(half(1.0) - T.y)));
321- #endif
322- }
323- if (basePixel.x < uniforms.screenWidth && basePixel.y + 1u < uniforms.screenHeight) {
324- let px01 = vec2u(basePixel.x, basePixel.y + 1u);
325- #ifdef PICK_MODE
326- textureStore(pickIdTexture, px01, vec4u(pickId01, 0u, 0u, 0u));
327- textureStore(pickDepthTexture, px01, vec4f(dAcc01, wAcc01, 0.0, 0.0));
328- #else
329- textureStore(outputTexture, px01, vec4f(decodeGamma3(vec3f(c01)), f32(half(1.0) - T.z)));
330- #endif
331- }
332- if (basePixel.x + 1u < uniforms.screenWidth && basePixel.y + 1u < uniforms.screenHeight) {
333- let px11 = vec2u(basePixel.x + 1u, basePixel.y + 1u);
334- #ifdef PICK_MODE
335- textureStore(pickIdTexture, px11, vec4u(pickId11, 0u, 0u, 0u));
336- textureStore(pickDepthTexture, px11, vec4f(dAcc11, wAcc11, 0.0, 0.0));
337- #else
338- textureStore(outputTexture, px11, vec4f(decodeGamma3(vec3f(c11)), f32(half(1.0) - T.w)));
339- #endif
340- }
329+ #ifdef HEATMAP_MODE
330+ atomicAdd(&sharedHeatCount, processedCount);
331+ workgroupBarrier();
332+ let avgCount = f32(atomicLoad(&sharedHeatCount)) / 64.0;
333+ let heatColor = vec4f(heatmapColor(avgCount), 1.0);
334+ if (basePixel.x < uniforms.screenWidth && basePixel.y < uniforms.screenHeight) {
335+ textureStore(outputTexture, basePixel, heatColor);
336+ }
337+ if (basePixel.x + 1u < uniforms.screenWidth && basePixel.y < uniforms.screenHeight) {
338+ textureStore(outputTexture, vec2u(basePixel.x + 1u, basePixel.y), heatColor);
339+ }
340+ if (basePixel.x < uniforms.screenWidth && basePixel.y + 1u < uniforms.screenHeight) {
341+ textureStore(outputTexture, vec2u(basePixel.x, basePixel.y + 1u), heatColor);
342+ }
343+ if (basePixel.x + 1u < uniforms.screenWidth && basePixel.y + 1u < uniforms.screenHeight) {
344+ textureStore(outputTexture, vec2u(basePixel.x + 1u, basePixel.y + 1u), heatColor);
345+ }
346+ #else
347+
348+ // Write results for the 2x2 pixel quad owned by this thread.
349+ // Pick mode: store the front-most pick ID and (accumulated depth, weight) per pixel.
350+ // Color mode: convert accumulated gamma-space color to linear via decodeGamma3 and store
351+ // to the rgba16float output texture; alpha holds total opacity (1 - transmittance).
352+ if (basePixel.x < uniforms.screenWidth && basePixel.y < uniforms.screenHeight) {
353+ #ifdef PICK_MODE
354+ textureStore(pickIdTexture, basePixel, vec4u(pickId00, 0u, 0u, 0u));
355+ textureStore(pickDepthTexture, basePixel, vec4f(dAcc00, wAcc00, 0.0, 0.0));
356+ #else
357+ textureStore(outputTexture, basePixel, vec4f(decodeGamma3(vec3f(c00)), f32(half(1.0) - T.x)));
358+ #endif
359+ }
360+ if (basePixel.x + 1u < uniforms.screenWidth && basePixel.y < uniforms.screenHeight) {
361+ let px10 = vec2u(basePixel.x + 1u, basePixel.y);
362+ #ifdef PICK_MODE
363+ textureStore(pickIdTexture, px10, vec4u(pickId10, 0u, 0u, 0u));
364+ textureStore(pickDepthTexture, px10, vec4f(dAcc10, wAcc10, 0.0, 0.0));
365+ #else
366+ textureStore(outputTexture, px10, vec4f(decodeGamma3(vec3f(c10)), f32(half(1.0) - T.y)));
367+ #endif
368+ }
369+ if (basePixel.x < uniforms.screenWidth && basePixel.y + 1u < uniforms.screenHeight) {
370+ let px01 = vec2u(basePixel.x, basePixel.y + 1u);
371+ #ifdef PICK_MODE
372+ textureStore(pickIdTexture, px01, vec4u(pickId01, 0u, 0u, 0u));
373+ textureStore(pickDepthTexture, px01, vec4f(dAcc01, wAcc01, 0.0, 0.0));
374+ #else
375+ textureStore(outputTexture, px01, vec4f(decodeGamma3(vec3f(c01)), f32(half(1.0) - T.z)));
376+ #endif
377+ }
378+ if (basePixel.x + 1u < uniforms.screenWidth && basePixel.y + 1u < uniforms.screenHeight) {
379+ let px11 = vec2u(basePixel.x + 1u, basePixel.y + 1u);
380+ #ifdef PICK_MODE
381+ textureStore(pickIdTexture, px11, vec4u(pickId11, 0u, 0u, 0u));
382+ textureStore(pickDepthTexture, px11, vec4f(dAcc11, wAcc11, 0.0, 0.0));
383+ #else
384+ textureStore(outputTexture, px11, vec4f(decodeGamma3(vec3f(c11)), f32(half(1.0) - T.w)));
385+ #endif
386+ }
387+
388+ #endif
341389}
342390` ;
0 commit comments