export const computeGsplatLocalPlaceEntriesSource: "\n\n@group(0) @binding(0) var<storage, read> pairBuffer: array<u32>;\n@group(0) @binding(1) var<storage, read> splatPairStart: array<u32>;\n@group(0) @binding(2) var<storage, read> splatPairCount: array<u32>;\n@group(0) @binding(3) var<storage, read> tileSplatCounts: array<u32>;\n@group(0) @binding(4) var<storage, read_write> tileEntries: array<u32>;\n@group(0) @binding(5) var<storage, read> sortElementCount: array<u32>;\n\n@compute @workgroup_size(256)\nfn main(@builtin(global_invocation_id) gid: vec3u, @builtin(num_workgroups) numWorkgroups: vec3u) {\n    let threadIdx = gid.y * (numWorkgroups.x * 256u) + gid.x;\n    let numVisible = sortElementCount[0];\n    if (threadIdx >= numVisible) {\n        return;\n    }\n\n    let rawCount = splatPairCount[threadIdx];\n    // High bit marks large splats handled by the cooperative LargePlaceEntries pass\n    if (rawCount == 0u || (rawCount & 0x80000000u) != 0u) {\n        return;\n    }\n    let count = rawCount;\n\n    let start = splatPairStart[threadIdx];\n    let pairLen = arrayLength(&pairBuffer);\n    let tileEntriesLen = arrayLength(&tileEntries);\n\n    for (var j: u32 = 0u; j < count; j++) {\n        let pairIdx = start + j;\n        if (pairIdx >= pairLen) { break; }\n\n        let packed = pairBuffer[pairIdx];\n        let tileIdx = packed >> 16u;\n        let localOff = packed & 0xFFFFu;\n\n        // tileSplatCounts has been prefix-summed, so it holds the start offset for each tile.\n        // localOff is the within-tile position assigned by atomicAdd during the count pass.\n        let entryIdx = tileSplatCounts[tileIdx] + localOff;\n        if (entryIdx < tileEntriesLen) {\n            tileEntries[entryIdx] = threadIdx;\n        }\n    }\n}\n";