00:00 NSA: karolherbst: just tested with the cheap adapter with a cable that is definitely hdmi 2.0
00:00 NSA: exactly the same result
00:01 karolherbst: well might be that nouveau is doing something stupid, but DP should work
00:01 karolherbst: did you try other resolutions?
00:02 NSA: with the adapter?
00:02 NSA: i guess i could try booting to windows and check if that works with the adapter
00:02 NSA: i only tested 4k60hz with the adapter
00:05 pmoreau: Ah, that test uses an OpControlBarrier: that might be the issue, that I do not handle it properly.
00:08 NSA: karolherbst: seems like it might actually be nouveau :(
00:08 NSA: windows seems to be fine 4k60hz with that cheap adapter and an hdmi 2.0 cable
00:21 pmoreau: karolherbst: If you have the blob nearby, could you trace local_arg_def and barrier for me please? I’d like to have a look at which barriers it generates.
00:24 NSA: karolherbst: so even if it doesn't work with nouveau it means the adapter works well with hdmi 2.0 4k@60hz for just €14 i think it was
02:32 pmoreau: karolherbst: I made some improvements to the execution barriers, and added some support for memory barriers. At least it no longer locks up (I think it was due to me not saying I was using a barrier: info->numBarriers), but I get some OOR_ADDR, probably not due to the barrier.
20:46 can705: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! yuwrohercy: Tomin mslusarz sarnex karolherbst Kabouik- FLHerne kb9vqf APTX_ calim towo` damke_ nyef`` schmidtm nd theMaze geaaru smithjd prg ivanich IdleGand
20:50 annadane: i know it's not in the christmas spirit, but i would like to violently injure people like that
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! bgrlxgew: theMaze snico kmshanah NolanSyKinsley paulk-gagarine mooch3 bonbons aboll friedrich whompy glisse geaaru nd prg RSpliet Q-Master damke_ IdleGandalf airlied robclark Kabouik- schmidtm Micha
21:02 mupuf: annadane: indeed... these people really have deep insecurity
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! kshwc: Asu snico towo` RSpliet MrRooks karolherbst theMaze paulk-gagarine schmidtm MichaelLong sarnex nd IdleGandalf salamanderrake Tomin mooch3 robclark udoprog APTX_ infinity0 hussam mslusarz Kabouik- damke
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! tagtkkja: towo` fdo-vcs mslusarz jkucia kmshanah MichaelLong FLHerne NanoSector perfinion prg salamanderrake udoprog MrRooks Asu nd smithjd geaaru sarnex robclark annadane airlied aboll whompy hussa
21:02 mupuf: notdoom758: get out
21:02 mooch3: FUUUUUUUCK
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! oollzrlu: aboll theMaze snico kmshanah paulk-gagarine jkucia infinity0 thc202 NolanSyKinsley smithjd whompy glisse docmax RSpliet ravior IdleGandalf prg O
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! wbhifuay: damke_ docmax paulk-gagarine theMaze thc202 towo` bonbons ivanich salamanderrake udoprog MrRooks dupondje xerpi geaaru ravior APTX_ Q-Master nd Omar007 Tomin robclark aboll annad
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! nhmbofkqpj: udoprog dupondje rellig prg MichaelLong Tomin xerpi mslusarz smithjd hussam notdoom758 aboll SuperHeron annadane dri-logger jkucia paulk-gagarine salamanderrake s
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! tknojgp: glisse bonbons thc202 schmidtm MichaelLong Q-Master FLHerne karolherbst ravior prg NolanSyKinsley perfinion rellig nyef`` udoprog paulk-gagarine whompy smithjd
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! sbkuotaa: ravior APTX_ salamanderrake prg nd theMaze thc202 MrRooks kmshanah notdoom758 whompy RSpliet glisse ivanich bonbons aboll calim damke_ jkucia IdleGandalf
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! kqdghdkxy: docmax Tomin jkucia salamanderrake ravior bonbons sarnex IdleGandalf Q-Master MichaelLong paulk-gagarine km
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! nczulefow: dri-logger NolanSyKinsley paulk-gagarine mooch3 smithjd hussam fdo-vcs Kabouik- kmshanah notdoom758 geaaru theMaze salamand
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! rsymhxttz: snico damke_ Kabouik- nyef`` paulk-gagarine glisse prg jkucia kmshanah salamanderrake smithjd NolanSyKinsley xerpi NanoSector MichaelLong bonbons sarnex mooch3 friedrich aboll t
21:02 mooch3: they've been hitting channels all over freenode
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! eikhgpafj: RSpliet ravior whompy annadane udoprog Q-Master towo` NanoSector rellig xerpi SuperHeron Asu FLHerne jkucia dupondje salama
21:02 notdoom758: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! adyjwwx: mslusarz NolanSyKinsley hussam Kabouik- SuperHeron MrRooks xerpi nd thc202 ivanich jkucia RSpliet geaaru snico udoprog robclark Tomin dupondje karolherbst bonb
21:02 mupuf: imirkin: I need your wizzardry
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! ehjkj: glisse Tomin kmshanah friedrich ravior NanoSector jkucia karolherbst annadane nd aboll calim salamanderrake
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! wronrito: whompy kmshanah snico dupondje geaaru mslusarz SuperHeron Kabouik- thc202 mooch3 hussam annada
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! dsmgshmtyf: prg robclark glisse Tomin towo` SuperHeron paulk-gagarine hussam sarnex ravior kmshanah MichaelLong NanoSector m
21:02 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! bixygmzklu: thc202 theMaze smithjd aboll robclark dri-logger prg rellig perfinion Tomin bonbons salamanderrake ravior Omar00
21:02 orbea:uses ignore
21:03 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! lpxtsgcj: RSpliet paulk-gagarine Tomin dupondje xerpi aboll geaaru IdleGandalf whompy infinity0 NanoSect
21:03 ketami60: ▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! sdjvdn: salamanderrake whompy fdo-vcs RSpliet Asu APTX_ FLHerne friedrich notdoom758 mooch3 glisse theMaze aboll robclark ravior Tomin
21:03 NanoSector: i'd wish those spammers would get a life
21:03 ketami60: ▄▄▄▄▄▄▄▄▄▄▄ HAPPY NIGGER MAS!! IF YOU WANT JOIN A CELEBRATION THAT IS NIGGER FREE PLEASE JOIN #/JOIN RIGHT HERE ON THIS NETWORK!! tvvlfuwgyn: notdoom758 NolanSyKinsley udoprog SuperHeron glisse airlied NanoSector annadane APTX_ whompy mooch3 theMaze Kabouik- docmax karolherbst RSpliet fdo-vcs nyef`` dupondje xerpi towo` Id
21:19 karolherbst: mupuf: :)
21:19 karolherbst: orr wait
21:19 karolherbst: do it yourself :p
21:20 mupuf: karolherbst: thx
23:11 karolherbst: duh.. tex args ordering is indeed annoying
23:24 mupuf:is sorry for the noise, but he had to review his chansrv basics :p
23:25 pmoreau: I need to check by running the code, but it looks like ldc.u8 does not have to be 0x4 aligned.
23:25 pmoreau: mupuf: No worries!
23:30 pmoreau: And if I have something like `cvt.s8.s32 r2, r1; st.global.s8 [dest], r2;` it can be simplified to `st.global.s8 [dest], r1;`. It doesn’t look like that small optimisation is currently done by Nouveau.
23:32 karolherbst: pmoreau: interesting
23:34 karolherbst: pmoreau: so we could actually do ld $r0 l[0x1]?
23:34 pmoreau: Hum, I haven’t tried from lmem.
23:35 karolherbst: well
23:35 karolherbst: it most likely works if it also works for global
23:35 pmoreau: But gmem does allow it, and it looks like cmem does as well.
23:35 karolherbst: shared and const might be more interesting
23:36 karolherbst: pmoreau: you are aware that this only helps with quite a few shaders ;) I doubt it is really worth the effort to detect shifts/cvts for this
23:36 karolherbst: and then
23:36 karolherbst: are just 8 bits written or all 32?
23:37 karolherbst: are the bits 0-7 written or according to the offset from 0x4 alignnment?
23:37 pmoreau: Only 8 should be written, but I haven’t checked
23:38 karolherbst: pmoreau: ohh did you hear about the volta ISA?
23:38 pmoreau: Depends, but I think so
23:39 karolherbst: mhh, did you also hear that it is a 128 bit encoding?
23:39 pmoreau: Yup, I was there when you looked into it, after realising they had sched for each instruction rather than every third. ;-)
23:40 karolherbst: right, didn't know you read the messages :)
23:41 pmoreau: I mean, I was in the channel chatting with you, giving you an explanation about why they would have sched for every instruction, so yes, I was reading the messages :-p
23:41 karolherbst: anywany, for fun I asked within the llvm channel what the status is on support for a 128 encoded ISA, sadly nobody answered, but it kind of was also a bad timing.
23:41 karolherbst: ahhh
23:41 karolherbst: really
23:41 karolherbst: I am getting old
23:41 pmoreau: But generally, I always read the logs
23:41 karolherbst: my memory
23:41 pmoreau: No worries
23:42 pmoreau: Yeah, Christmas/New Year isn’t the best time
23:43 pmoreau: So every instruction would be 128-bit wide?
23:44 pmoreau: Dang, I made him quit
23:44 karolherbst: :p
23:44 karolherbst: no
23:44 karolherbst: nouveau made me
23:44 pmoreau: :-D
23:44 karolherbst: rarely when running piglit, secboot fails in the middle
23:44 karolherbst: and only rmmod, suspend, load helps
23:44 pmoreau: OK
23:44 karolherbst: well, I run it in concurrent mode so...
23:45 karolherbst: I am actually surised the chance of messing up is below 10%
23:45 karolherbst: *surprised
23:45 karolherbst: but yeah
23:45 karolherbst: each instruction would be 128 bit
23:45 karolherbst: I am sure they reduced the amount of forms though
23:45 karolherbst: and that the sched infos are at the same place for all instructions
23:45 karolherbst: generally I think they made it much simplier to parse the instructions
23:46 pmoreau: That’s weird, cause dumping the SASS of code for SM 7.0 still seems to return 64-bit instructions
23:46 karolherbst: no
23:46 karolherbst: it should print two lines for each
23:46 pmoreau: Oh, if you count the sched as well, then sure
23:46 karolherbst: uhm
23:47 karolherbst: the scheds aren't 64 bit
23:47 karolherbst: bits 64-91 contain important stuff
23:47 karolherbst: like predicate sources for some insturcitons
23:47 karolherbst: and other things
23:47 karolherbst: I am not 100% sure how big the sched part is though
23:47 pmoreau: Ah OK, I see
23:48 karolherbst: I think they might be bigger than maxwell
23:48 karolherbst: but maxwell didn't used that much place per instruction
23:48 karolherbst: did maxwell use all 21 bits per instruction?
23:48 pmoreau: I could imagine that if they used 64-bit for 3 instructions, it seems quite extreme that they tripled the number of sched bits per instruction from ~21 to 64.
23:49 karolherbst: if I would guess I would say 18 or 19
23:49 karolherbst: yeah
23:49 karolherbst: I am sure it is as big as in maxwell
23:49 karolherbst: 32 at most
23:49 karolherbst: well
23:49 karolherbst: a 96/32 bit split would make sense
23:50 karolherbst: you could even make the hw so smart, that you have two parts parsing those
23:50 karolherbst: and the first three 32 bit parts go into the instruction parser and the other 32 bits into the sched parser or something
23:53 pmoreau: Looking at maxas, it seems that the sched was taking 17-bit wide for each instruction, and then you had 4 bits per instruction for reuse flags
23:55 pmoreau: So: 51 bits for scheduling, 12 bits for reuse, and only 1 bit unused
23:56 karolherbst: okay
23:56 karolherbst: so at least 21 bits :)
23:58 pmoreau: Yup :-)