vello

It's an area of active research. One very promising project is vello [1] which is essentially a software renderer on the GPU (requires GPU compute support). Preliminary results are looking good. Better than anything else out there good. -- [1] https://github.com/linebender/vello

I love these kinds of things and use them in GPU programming, among other things. Things have changed in a variety of ways: population count and count-trailing-zeros are generally available as fast instructions now. Multiply is also now just as fast as other operations, so is not to be avoided. A couple examples. [1] computes the sum of the number of bytes used of four consecutive segments of a bezier path - each segment can be lineto, quadto, or curveto, can be the end of a path or not, and be i16 or f32. 4 tag bytes are packed into a 32 bit word, it computes all these, then sums them together. [2] linearizes a recursive subdivision into an iterative loop. The stack depth is represented as the number of zeros in a word, so pushing the stack is a left shift, and popping is a right shift. It turns out you want to pop multiple levels at onc

Raph Linus has recently has his work posted on HN, which features GPU rasterization [1] Also, Pathfinder [2] has the ability to do most of its rasterization on the GPU. I read somewhere that an experimental renderer for Google’s Fuchsia (Spinel?) [3] also does this. 1: https://github.com/linebender/piet-gpu 2: https://github.com/servo/pathfinder 3: https://www.tdcommons.org/cgi/viewcontent.cgi?article=1580&c...

Thanks for the matmul sizes, but the question I am more interested in is precision. Matrix multiply throughput can be dramatically higher for tensor cores[1] than normal shader ALU, specifically for reduced precision arithmetic. I'm wondering to what extent that's accessible on the M1, and to what extent IREE can address them. Regarding prefix sum, the specific question I'm interested in is that SPIRV OpControlBarrier with device scope gets translated into threadgroup_barrier(mem_device) [2]. That's insufficient to make decoupled look-back work. Conversely, if you're not using decoupled look-back, you're not getting the full throughput on GPUs that do support that barrier. I'm wondering how your infrastructure deals with that. [1]: https://developer.nvidia.com/blog/programming-tensor-cores-c... [2]: https://github.com/linebender/piet-gpu/