struct LightData { position : vec4; color : vec3; radius : f32; }; [[block]] struct LightsBuffer { lights : array; }; [[group(0), binding(0)]] var lightsBuffer : LightsBuffer; struct TileLightIdData { count : atomic; lightId : array; }; [[block]] struct Tiles { data : array; }; [[group(1), binding(0)]] var tileLightId : Tiles; [[block]] struct Config { numLights : u32; numTiles : u32; tileCountX : u32; tileCountY : u32; numTileLightSlot : u32; tileSize : u32; }; [[group(2), binding(0)]] var config : Config; [[block]] struct Uniforms { min : vec4; max : vec4; viewMatrix : mat4x4; projectionMatrix : mat4x4; fullScreenSize : vec4; }; [[group(3), binding(0)]] var uniforms : Uniforms; [[stage(compute), workgroup_size(64, 1, 1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { var index = GlobalInvocationID.x; if ((index >= config.numLights)) { return; } lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.100000001) + (0.001 * (f32(index) - (64.0 * floor((f32(index) / 64.0)))))); if ((lightsBuffer.lights[index].position.y < uniforms.min.y)) { lightsBuffer.lights[index].position.y = uniforms.max.y; } var M : mat4x4 = uniforms.projectionMatrix; var viewNear : f32 = (-(M[3][2]) / (-1.0 + M[2][2])); var viewFar : f32 = (-(M[3][2]) / (1.0 + M[2][2])); var lightPos = lightsBuffer.lights[index].position; lightPos = (uniforms.viewMatrix * lightPos); lightPos = (lightPos / lightPos.w); var lightRadius : f32 = lightsBuffer.lights[index].radius; var boxMin : vec4 = (lightPos - vec4(vec3(lightRadius), 0.0)); var boxMax : vec4 = (lightPos + vec4(vec3(lightRadius), 0.0)); var frustumPlanes : array, 6>; frustumPlanes[4] = vec4(0.0, 0.0, -1.0, viewNear); frustumPlanes[5] = vec4(0.0, 0.0, 1.0, -(viewFar)); let TILE_SIZE : i32 = 16; let TILE_COUNT_X : i32 = 2; let TILE_COUNT_Y : i32 = 2; for(var y : i32 = 0; (y < TILE_COUNT_Y); y = (y + 1)) { for(var x : i32 = 0; (x < TILE_COUNT_X); x = (x + 1)) { var tilePixel0Idx : vec2 = vec2((x * TILE_SIZE), (y * TILE_SIZE)); var floorCoord : vec2 = (((2.0 * vec2(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - vec2(1.0)); var ceilCoord : vec2 = (((2.0 * vec2((tilePixel0Idx + vec2(TILE_SIZE)))) / uniforms.fullScreenSize.xy) - vec2(1.0)); var viewFloorCoord : vec2 = vec2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1])); var viewCeilCoord : vec2 = vec2((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1])); frustumPlanes[0] = vec4(1.0, 0.0, (-(viewFloorCoord.x) / viewNear), 0.0); frustumPlanes[1] = vec4(-1.0, 0.0, (viewCeilCoord.x / viewNear), 0.0); frustumPlanes[2] = vec4(0.0, 1.0, (-(viewFloorCoord.y) / viewNear), 0.0); frustumPlanes[3] = vec4(0.0, -1.0, (viewCeilCoord.y / viewNear), 0.0); var dp : f32 = 0.0; for(var i : u32 = 0u; (i < 6u); i = (i + 1u)) { var p : vec4; if ((frustumPlanes[i].x > 0.0)) { p.x = boxMax.x; } else { p.x = boxMin.x; } if ((frustumPlanes[i].y > 0.0)) { p.y = boxMax.y; } else { p.y = boxMin.y; } if ((frustumPlanes[i].z > 0.0)) { p.z = boxMax.z; } else { p.z = boxMin.z; } p.w = 1.0; dp = (dp + min(0.0, dot(p, frustumPlanes[i]))); } if ((dp >= 0.0)) { var tileId : u32 = u32((x + (y * TILE_COUNT_X))); if (((tileId < 0u) || (tileId >= config.numTiles))) { continue; } var offset : u32 = atomicAdd(&(tileLightId.data[tileId].count), 1u); if ((offset >= config.numTileLightSlot)) { continue; } tileLightId.data[tileId].lightId[offset] = GlobalInvocationID.x; } } } }