This commit is contained in:
parent
6bc125be1c
commit
ed5a4379db
76 changed files with 1418 additions and 988 deletions
|
|
@ -15,40 +15,53 @@ const clearCompactedTailStride = 4u;
|
|||
@group(1) @binding(2) var<storage, read_write> counters: Counters;
|
||||
@group(1) @binding(3) var<storage, read_write> compactedAgents: array<Agent>;
|
||||
|
||||
var<workgroup> workgroupAliveCount: atomic<u32>;
|
||||
var<workgroup> workgroupCompactedOffset: u32;
|
||||
var<workgroup> scanData: array<u32, agentWorkgroupSize>;
|
||||
var<workgroup> clearAliveAgentCount: u32;
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
@compute @workgroup_size(agentWorkgroupSize)
|
||||
fn main(
|
||||
@builtin(global_invocation_id) global_id: vec3<u32>,
|
||||
@builtin(local_invocation_id) local_id: vec3<u32>
|
||||
) {
|
||||
let id = get_id(global_id);
|
||||
let lid = local_id.x;
|
||||
|
||||
if local_id.x == 0u {
|
||||
atomicStore(&workgroupAliveCount, 0u);
|
||||
}
|
||||
|
||||
workgroupBarrier();
|
||||
|
||||
var localCompactedIndex = 0u;
|
||||
var agent: Agent;
|
||||
var isAlive = false;
|
||||
var agent: Agent;
|
||||
if id < settings.agentCount {
|
||||
isAlive = agents[id].colorIndex >= 0.0;
|
||||
if isAlive {
|
||||
agent = agents[id];
|
||||
localCompactedIndex = atomicAdd(&workgroupAliveCount, 1u);
|
||||
}
|
||||
}
|
||||
|
||||
// Hillis-Steele inclusive prefix sum across the workgroup. Replaces a
|
||||
// per-thread atomicAdd to a workgroup counter, eliminating serialization
|
||||
// on dense workgroups.
|
||||
scanData[lid] = select(0u, 1u, isAlive);
|
||||
workgroupBarrier();
|
||||
|
||||
if local_id.x == 0u {
|
||||
let groupAliveCount = atomicLoad(&workgroupAliveCount);
|
||||
if groupAliveCount > 0u {
|
||||
workgroupCompactedOffset = atomicAdd(&counters.aliveAgentCount, groupAliveCount);
|
||||
var offset: u32 = 1u;
|
||||
while offset < agentWorkgroupSize {
|
||||
let own = scanData[lid];
|
||||
var contribution: u32 = 0u;
|
||||
if lid >= offset {
|
||||
contribution = scanData[lid - offset];
|
||||
}
|
||||
workgroupBarrier();
|
||||
scanData[lid] = own + contribution;
|
||||
workgroupBarrier();
|
||||
offset = offset * 2u;
|
||||
}
|
||||
|
||||
let inclusivePrefix = scanData[lid];
|
||||
let workgroupAliveTotal = scanData[agentWorkgroupSize - 1u];
|
||||
let exclusivePrefix = inclusivePrefix - select(0u, 1u, isAlive);
|
||||
|
||||
if lid == 0u {
|
||||
if workgroupAliveTotal > 0u {
|
||||
workgroupCompactedOffset = atomicAdd(&counters.aliveAgentCount, workgroupAliveTotal);
|
||||
} else {
|
||||
workgroupCompactedOffset = 0u;
|
||||
}
|
||||
|
|
@ -57,11 +70,11 @@ fn main(
|
|||
workgroupBarrier();
|
||||
|
||||
if isAlive {
|
||||
compactedAgents[workgroupCompactedOffset + localCompactedIndex] = agent;
|
||||
compactedAgents[workgroupCompactedOffset + exclusivePrefix] = agent;
|
||||
}
|
||||
}
|
||||
|
||||
@compute @workgroup_size(64)
|
||||
@compute @workgroup_size(agentWorkgroupSize)
|
||||
fn clearCompactedTail(
|
||||
@builtin(global_invocation_id) global_id: vec3<u32>,
|
||||
@builtin(local_invocation_id) local_id: vec3<u32>
|
||||
|
|
|
|||
Loading…
Reference in a new issue