1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
-- Basic compute shader test
-- Exercises: shared memory, barrier(), compute builtins, workgroup_size, buffer binding
struct Params
count: uint
scale: float
end
struct DataBuffer
values: [64]float
end
@group(0) @binding(0) buffer params: Params
@group(0) @binding(1) buffer data: DataBuffer
shared scratch: [64]float
@entry(compute)
@workgroup_size(64, 1, 1)
function main(
@builtin(global_invocation_id) global_id: uvec3,
@builtin(local_invocation_id) local_id: uvec3
)
let idx = local_id.x
-- Load data into shared memory
scratch[idx] = data.values[idx] * params.scale
barrier()
-- Simple parallel reduction step: add neighbor
if idx < 32 then
scratch[idx] = scratch[idx] + scratch[idx + 32]
end
barrier()
-- Write result back
data.values[idx] = scratch[idx]
end