01:07 mohamexiety[d]: phomes_[d]: got distracted but the main mr should be fixed now. thanks a lot for your help! ❤️
13:20 zmike[d]: it seems like descriptor heap doesn't improve perf all that much vs descriptor buffer, which leaves me even less motivated to consider a zink implementation :/
14:33 karolherbst[d]: ohh right.. I wanted to look into the imnmx.u32 r8 ur6 pT + r2ur pattern...
14:48 karolherbst[d]: seeing patterns like this in a few shaders:
14:48 karolherbst[d]: /*0000*/ BMOV.32.CLEAR RZ, B0 ;
14:48 karolherbst[d]: /*0010*/ BSSY B0, 0x30 ;
14:48 karolherbst[d]: /*0020*/ BSYNC B0 ;
14:48 karolherbst[d]: /*0030*/ NOP ;
15:46 mhenning[d]: karolherbst[d]: how do we end up with that?
15:46 karolherbst[d]: I wished I'd knew
15:47 karolherbst[d]: the shader is massive with impressive CFG
15:47 karolherbst[d]:but
15:47 karolherbst[d]: the label NAK referenced doesn't exist
15:48 mhenning[d]: hmm weird
15:48 mhenning[d]: zmike[d]: "doesn't improve perf" on what workload, out of curiosity?
15:53 karolherbst[d]: mhenning[d]: https://gist.githubusercontent.com/karolherbst/b43a50ef16ec7fcf0322a9b809e726b3/raw/cfd673d96fdda25a2e0f854e72b5499423f1d5cb/gistfile1.txt
15:53 karolherbst[d]: just straight up there after from_nir
15:53 karolherbst[d]: and no idea where it's coming from
15:57 mhenning[d]: oh, that's the shared memory initialization thing
15:58 mhenning[d]: src/nouveau/compiler/nak/from_nir.rs line 4086
15:59 karolherbst[d]: ohh the stuff inside `parse_block`? My line numbers are too much off
15:59 karolherbst[d]: mhhhhhh
15:59 mhenning[d]: yes
15:59 karolherbst[d]: weird...
16:17 karolherbst[d]: I should probably try to figure this out and see what's going on there.. but also check how much of a perf problem that is
16:18 karolherbst[d]: but at least it's only in the first block..
16:19 karolherbst[d]: anyway.. seems super unimportant 🙃
16:29 zmike[d]: mhenning[d]: from vkd3d-proton apparently
20:57 karolherbst[d]: mhenning[d]: any further high-level comments on https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40392 ? I kinda want to start CTSing it, but that also includes checking on older GPUs and I'd rather do it once rather than again after more review 😄
21:04 karolherbst[d]: mhh
21:04 karolherbst[d]: r80 = ld.global.a64.constant.u8 [r80+ur8..10], pT // delay=1 wr:0
21:04 karolherbst[d]: r80 = prmt r80 [0x4440] rZ // delay=1 wt=000001
21:04 karolherbst[d]: r80 = shf.r.w.u32.hi rZ r80 r84 // delay=1
21:04 karolherbst[d]: I _wonder_ if we should assume that from loads the high bits to be properly set...
21:06 karolherbst[d]: like not really sure why we are so aggressive with prmt here..
21:07 karolherbst[d]: ohh because of the `u2u32` mhh
21:07 karolherbst[d]: mhhhhhh
21:09 karolherbst[d]: I wanted to play around with the "dest_type" idea of loads, right? mhhh
21:36 karolherbst[d]: or maybe I should do that in copy_prop? mhh
21:48 karolherbst[d]: okay.. those prmts are annoying me 😄