Source
- URL: https://x.com/JesseStojan/status/2058768097939472658
- Author: Jesse S (@JesseStojan)
- Posted: 2026-05-25 04:32:19
Branch
1/ @JesseStojan
@AgileJebrim
Using a hierarchical/structured layers on-device parallel reduction of cascading (multi-pass) kernels.
E.g. hip/cub::DeviceReduce::Sum(...)/sycl::Reduction(...) with more steps.
(I saw you say 16M × uint8 in the replies, but it was one of the first questions in my head lol)
2/ @AgileJebrim
@JesseStojan
Cool. Provide more details please.
3/ @AgileJebrim
@JesseStojan
You’re implementing the low level details, no reusing an existing function.
4/ @JesseStojan
@AgileJebrim
Well in that case, 1st kernel pass, each workgroup strides through the global array, sums the values, tree reduction in shared memory, writes the partial sum to a (uint32/64) partials[blockN] buffer. 2nd kernel pass reduces the partials[] down to the final sum.
5/ @AgileJebrim
@JesseStojan
Tell me more about this step labeled “sums the values”
6/ @JesseStojan
@AgileJebrim
Each thread sums many values using a stride (num_workgroups × threads_per_workgroup). Within each iteration, neighboring threads in a (subgroup/warp/wavefront) access consecutive addresses, enabling coalesced wide bursts (one transaction per subgroup).
7/ @AgileJebrim
@JesseStojan
Do you think you can write a short for loop snippet for this in a shader language?
8/ @JesseStojan
@AgileJebrim
Yeah, though just fair heads up, I’m not seeking employment lol, I’m just a nerd for rendering/CGI and massively parallel computing.