20:21 karolherbst: ehh.. debugging gpu code without a debugger is annoying :(
20:23 imirkin_: s/gpu code without a debugger//
20:23 karolherbst: :p
20:25 karolherbst: at least it's not a codegen bug
20:27 imirkin_: hw bug? :)
20:27 imirkin_: all bugs are codegen bugs, even if they're not
20:28 karolherbst: probably a nir one, blaming the structurizer for now
20:28 imirkin_: i wasn't super-pleased with this one: https://cgit.freedesktop.org/mesa/mesa/commit/?id=3e9aacb139dc4cb101780e235fc5dd45acf860f3
20:28 karolherbst: ohh yeah.. that one
20:28 karolherbst: stuff like that is annoying
20:29 karolherbst: I had a similiar bug in nir once, where a vec1 value was just treated as vec4 in nir. Other driver didn't care, on nv hw the vp just overwrote random stuff :(
20:29 karolherbst: that's also slightly annoying to figure out
20:29 imirkin_: hehe
20:31 karolherbst: actually the structurizer is fine
20:31 karolherbst: vec1 64 ssa_50 = deref_struct &ssa_49->field2 (global uint64_t) /* &(*(struct.Node *)ssa_47)[ssa_9].field2 */
20:31 karolherbst: the ssa_9 indirection turns into 0 somewhere
20:31 imirkin_: ssa_9 = 0? :)
20:32 karolherbst: ssa_9 = get_global_id(0)
20:32 imirkin_: hehe
20:32 imirkin_: s/get_global_id//
20:32 karolherbst: it works if I get rid of the loop in the kernel
20:32 karolherbst: ...
20:32 imirkin_: wait, isn't get_global_id(0) == 0? what is that func?
20:33 imirkin_: oh, i guess not 0, but something constant
20:33 karolherbst: it's the invoc id in glsl terms
20:33 karolherbst: is it even called invoc id in compute shaders?
20:33 imirkin_: there are several
20:34 imirkin_: i think this is the gl_GlobalInvocationID
20:34 karolherbst: yep
20:34 imirkin_: or gl_GlobalInvocationIndex
20:34 karolherbst: ID
20:34 karolherbst: index is the flattened value
20:34 imirkin_: in uvec3 gl_GlobalInvocationID
20:34 imirkin_: right yeah
20:34 imirkin_: The value of gl_GlobalInvocationID is equal to gl_WorkGroupID * gl_WorkGroupSize + gl_LocalInvocationID
20:34 karolherbst: yep
20:34 karolherbst: and that all works
20:35 karolherbst: the kernel is this one: https://github.com/KhronosGroup/OpenCL-CTS/blob/master/test_conformance/SVM/test_cross_buffer_pointers.cpp#L38
20:35 karolherbst: pNode->global_id = i; gets filled correctly
20:36 karolherbst: pNode->pNext = pNew; gets only set for pNodes[0]
20:38 karolherbst: I guess I just check which nir pass messes it up... and I am sure it's something stupid
20:39 karolherbst: btw, segfaults on the GPU aren't funny... I even manage to take down gr completly with this test
20:43 karolherbst: uhhhh
20:43 karolherbst: ufff
20:44 karolherbst: yeah well... if the stride becomes 0 of that type, the indirection goes like get_global_id(0) * 0
20:46 imirkin_: oops!
20:53 karolherbst: ohhh
20:54 karolherbst: vec1 64 ssa_31 = deref_cast (struct.Node *)ssa_3 (global struct.Node) /* ptr_stride=0 */ yeah well
20:54 imirkin_: blame it on the structurizer, why don't ya :p
20:55 karolherbst: yeah.. it's probably something in vtn
20:55 karolherbst: because pre structurizing the stride is 0 as well
20:56 karolherbst: that structurizer that one person came up with is nice btw and super helpful :) I can imagine more pleasent things to work on
21:47 karolherbst: yay, another SVM test passing :)
21:48 lovesegfault: SVM?
21:48 karolherbst: shared virtual memory
21:49 karolherbst: actually, I fixed two tests :)
21:49 karolherbst: and the other fails are pure bs
21:49 karolherbst: imirkin_: int64 atomics?
21:49 lovesegfault: Ah :)
21:49 imirkin_: karolherbst: should work afaik
21:50 karolherbst: imirkin_: ohh, really_
21:50 karolherbst: let me check
21:50 imirkin_: at least at some level
21:50 karolherbst: PIPE_SHADER_CAP_INT64_ATOMICS
21:50 karolherbst: is set to 0
21:50 imirkin_: must have gotten added
21:50 imirkin_: and/or means something intel-specific
21:50 karolherbst: well..
21:50 karolherbst: clover is the only one using that cap
21:51 imirkin_: ah
21:51 imirkin_: anyways, afaik it should work
21:51 karolherbst: okay...
21:51 imirkin_: if you get the right stuff into the ir
21:51 karolherbst: sure
21:53 karolherbst: ehhh...
21:53 karolherbst: the svm test also requires generic pointers :(
21:53 karolherbst: ahh, the atom test hit an assert
21:53 karolherbst: nice
21:54 karolherbst: imirkin_: int64 atomics on shared memory?
21:54 imirkin_: dunno about that
21:55 imirkin_: looks like yes
21:55 imirkin_: https://cgit.freedesktop.org/mesa/mesa/tree/src/gallium/drivers/nouveau/codegen/nv50_ir_emit_gm107.cpp#n2642
21:55 imirkin_: make sure to use U64 with CAS though
21:55 karolherbst: okay... well.. shared memory isn't something I looked into that deeply yet anyway... 64 bit pointers and such
21:58 karolherbst: ohh, cool, it works :)
21:59 karolherbst: well, with global memory that is
21:59 karolherbst: I'll ignore shared for now
22:07 imirkin_: i don't see why shared wouldn't work ... this is on pascal, i presume?
22:08 imirkin_: what's the error?
22:08 karolherbst: ohh, some vtn stuff
22:08 imirkin_: oh. yeah. can't help you with that :p
22:08 karolherbst: nir code asserts on 32 bit pointers
22:08 karolherbst: so...
22:08 imirkin_: nv50_ir should handle it fine :p
22:08 karolherbst: CL is stupid, every pointer is 64 bit if your addressing mode is 64 bit
22:08 imirkin_: 128-bit is iffy.
22:08 karolherbst: even for shared memory
22:09 karolherbst: and I was to lazy to properly handle all of that yet
22:09 imirkin_: sure.
22:09 karolherbst: and with generic pointers everything gets even messier
22:19 karolherbst: anyway, FAILED 2 of 14 tests. \o/
22:20 imirkin_: well, lmk if you need help with nv50_ir thing
22:20 imirkin_: things*
22:20 karolherbst: pmoreau wanted to look into that :p
22:20 karolherbst: but at some point I will anyway
22:20 karolherbst: just really want to finish this SVM stuff first
22:20 karolherbst: then I'll have time for other things again
22:23 karolherbst: imirkin_: ohh, or did you mean codegen in general
22:23 imirkin_: codegen in general.
22:23 imirkin_: i.e. lack of support for something, etc
22:24 karolherbst: yeah... the only annoying bit are still RA fails with wide types... the RA situation when it tries to allocate a phi of movs + merges or stuff.. I think that was the big issue
22:24 imirkin_: did you end up pushing your ifx?
22:24 karolherbst: nope
22:25 imirkin_: o
22:25 edgecase: imirkin_ this old Thinkpad has suspend issues >= 5.3.0 also
22:25 karolherbst: imirkin_: but I meant this fix anyway: https://github.com/karolherbst/mesa/commit/4d96398060f998f337237911c0ef60b2001d122b
22:25 imirkin_: edgecase: bisect?
22:25 edgecase: only 3 -rcX versions to bisect, yeah
22:25 edgecase: 5.2.21 ok -> 5.3.0 fail
22:25 imirkin_: karolherbst: that last hunk is almost certainly inaccurate
22:26 karolherbst: imirkin_: I mean, if you don't mind I can fix the patch, but I thought you still wanted to take a look or something... or did I push it?
22:26 imirkin_: karolherbst: iirc i said give me a week ... that was a month ago
22:26 imirkin_: it's hard to make time to properly review this stuff
22:26 imirkin_: i'd rather have no review than half-review
22:26 karolherbst: well.. I pushed it
22:26 karolherbst: I just forgot
22:26 imirkin_: :)
22:27 edgecase: i'd still like to pursue the freeze at dma fence, what's next step, check existing bugs, file new one?
22:27 karolherbst: ohh wait..
22:27 karolherbst: different commit with similiar messsage
22:27 karolherbst: ehhh
22:29 karolherbst: imirkin_: I'll push it then if you don't mind
22:29 karolherbst: hopefully somebody hits a regressions, but beyond shaders spilling it doesn't affect any, sooo... well
22:30 imirkin_: karolherbst: the second one you pasted, please don't
22:30 imirkin_: the one you had originally to fix use-after-free seemed like it might be fine
22:31 karolherbst: mhhh
22:31 karolherbst: the second one was the wrapper thingy, right
22:32 imirkin_: https://github.com/karolherbst/mesa/commit/4d96398060f998f337237911c0ef60b2001d122b
22:32 imirkin_: this seems questionable.
22:32 imirkin_: that code is SUPER finicky
22:32 karolherbst: ah, right, but that's for some other issue anyway,, not planning to push it
22:32 imirkin_: oh. so what were you asking about then?
22:33 karolherbst: https://github.com/karolherbst/mesa/commits/master those
22:33 karolherbst: (didn't push yet though)
22:33 imirkin_: right, that seems more acceptable
22:33 imirkin_: although i haven't reviewed carefully
22:34 imirkin_: and likely won't be able to
22:34 imirkin_: so use your own judgement :)
22:34 karolherbst: I'll run piglit over them again and see what I decide
22:34 karolherbst: but yeah...
22:34 karolherbst: would be cool to have this annoying bug fixed though
22:34 imirkin_: yes
22:34 karolherbst: and I like this approach more than the previous attempts
22:35 imirkin_: yeah, seems non-horrible in principle