15:30 alyssa: i am once again wondering how aco gets away without spilling linear vgprs
15:34 karolherbst: anyway, if nobody wants to review the CL builtin edge case things, I can probably also merge it anyway if nobody has any complains. It's going to be important for 3.1 conformance so I kinda want to merge it soonish
15:34 karolherbst: https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41726
15:35 karolherbst: but also if anybody has suggestions on how to make the tanpi handling less terrible..
15:40 pendingchaos: alyssa: we limit the number of them that ac_nir_lower_tex_coords() creates, and hope the shader doesn't spill too many sgprs
15:44 alyssa: pendingchaos: yeah, it's the second part of that I'm confused about
15:44 zmike: karolherbst: just merge
15:44 alyssa: I mean that was my plan for jay too but there's a CTS test with sgpr pressure over 1024
15:44 pendingchaos: spilling linear vgprs around function calls is also a bit special, where it's done in spill_preserved() after normal spilling and RA by temporarily setting exec to -1
15:44 pendingchaos: 1024 sgprs divided by 32 lanes is just 32 linear vgprs
15:45 alyssa: I guess you have a lot more than that
16:45 karolherbst: I think my is_entrypoint patch for opt_dead_write_vars also made the intel compiler really fast in my CL CTS runs :')
16:46 karolherbst: (was 22 minutes, now it's back to ~9)
16:52 Kayden: yikes
17:55 alyssa: Lol
18:23 glehmann: https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41953 looking for someone to review some atan2 fp control stuff for blender
20:36 karolherbst: do we want to have fma emulation code for Vulkan drivers?
20:37 karolherbst: I'm still wondering how I want to lower SPIR-V fma in rusticl, and could either call into a libclc function, just as a lowering pass outside of vtn, make the core parts of vtn call into a function to do that, or make a copy of the libclc fma emulation and add it into mesa as a ... nir helper? Not quite sure yet, the emulation is pretty huge
20:38 karolherbst: here is the libclc emulation code: https://github.com/llvm/llvm-project/blob/main/libclc/clc/lib/spirv/math/clc_sw_fma.cl
20:38 karolherbst: but I'm kinda wondering if there isn't a simpler way of doing that...
20:39 airlied: now that mesa can compile CL to NIR can you just pull in the cl file?
20:40 karolherbst: have you seen the includes?
20:40 karolherbst: though I did consider pulling all of libclc into mesa 🙃
20:43 glehmann: karolherbst: imo no to emulated VK_KHR_shader_fma
20:44 karolherbst: yeah... that's fair. I wonder _if_ we ever need it, e.g. games actually requiring to run, but we can discuss this once we hit that...
20:45 glehmann: d3d12 has no fma, and for the few native vk games we can cross that bridge when we get there
20:45 glehmann: but it doesn't make sense to expose something that we don't want people to use
20:45 karolherbst: right
20:46 karolherbst: I'll just need emulation for CL anyway, but I'm fine with using libclc as I'm already doing
23:09 karolherbst: I'm going to commit nir crimes: "workgroup_size: 3, 5, 7 (variable)"