04:30 haigedkinomehed: My understanding is that gpgpu-sim has an assembler in 1.9 SM of cuda SASS, which is forward revision compatible, more than enough to implement real code for NVIDIA graphics cards. so mattst88 please take your wrinkle meds! You are rotting from a face. And turning to an asshole hence. There must be some way you get this fungus off your face and forget bullying, so take your meds ok!!?
04:35 haigedkinomehed: mattst88: also for your two way diarhhea there are solid meds for, please don't speak before you cure yourself ok?
04:40 haigedkinomehed: none do understand why anyone should hire a wrinkly fecalist you see, slight curing may bring you into big game, and even earn you couple of dollars!
07:10 sjh: does this call trace look like a bug or am I possibly missing something from my kernel config? https://bpa.st/M4ZQ
07:29 pmoreau: sjh: This does look like a bug.
07:30 pmoreau: karolherbst: Yeah, I’m thinking the host is not computing a valid block size for that amount of register usage, possibly due to us returning the wrong data. I’ll look into that.
07:31 pmoreau: Cause that’s probably the №1 issue I’m running into at the moment with the CTS tests.
07:37 pmoreau: Lyude: Do you think your series about DP fixes that just got merged could help sjh? (see https://bpa.st/M4ZQ)
07:38 pmoreau: The issue does look different, but it is regarding DP and GV100 so maybe something you ran into.
07:50 sjh: just got the full kernel log
07:53 sjh: https://pastebin.com/raw/nGmMqqda around 100.107391 is where i # modprobe nouveau
07:55 sjh: I can put this on the bug tracker, i think it moved to gitlab now?
07:56 pmoreau: Ah, TU104 and not GV100. And getting some timeout: “i2c: aux 0007: timeout 01119000” and “disp: error 00020000”.
07:58 sjh: i didn't build any firmware into the kernel, i was hoping nouveau would complain in dmesg which files it needed but it hangs like this when i load it
07:58 pmoreau: The bug tracker on gitlab is mostly for Mesa bugs, not kernel ones.
07:59 sjh: ok because bugzilla says I cant post new bugs
07:59 pmoreau: Right
07:59 sjh: github issues then?
08:00 pmoreau: Probably the best is to send on the mailing list.
08:01 pmoreau: And we need to figure out how to track those bugs, and update the wiki.
08:04 sjh: mm okay. i'm going to mess with it a bit more. nouveau works on this card with other distributions
08:05 sjh: but something i've done on gentoo is not playing nice
08:06 sjh: actually no I think I did use proprietary drivers previously too
08:09 pmoreau: If you have a chance, try https://github.com/skeggsb/nouveau/commits/master: it contains the DP patches I was referring to earlier.
08:10 pmoreau: You will need to build it against 0e7e6198af28c1573267aba1be33dd0b7fb35691 from the drm-next branch.
08:13 pmoreau: karolherbst: So, clEnqueueNDRangeKernel() is being called with no block size, leaving it to clover to pick the most appropriate block size.
08:14 pmoreau: However clover ends up using the non-kernel dependent max_threads_per_block and max_block_size for that. :-/ We need to revive robclark’s series about querying those values for a given kernel.
08:54 karolherbst: pmoreau: yeah
12:07 pmoreau: Mmh, looking at the patches again, I think I might just start from scratch; a few things have change since those patches were last written.
12:23 karolherbst: pmoreau: the problem is just, that we need to change how compilation works in clover
12:24 karolherbst: pmoreau: for clGetKernelWorkGroupInfo we need to fully compile before launching the kernel
12:24 pmoreau: Arf, nv50_program_translate & co are only called when launching the kernel
12:25 karolherbst: let me verify that though
12:25 pmoreau: I had forgotten that part and was just thinking about the translation to NIR that happens during program building and no longer during kernel creation or launching as it used to.
12:26 karolherbst: pmoreau: https://gist.githubusercontent.com/karolherbst/8a408ce3e72c9a4d34c1fd73d7ae6299/raw/48c00076f86c62e408c52d80ab06a75b0101151b/gistfile1.txt
12:26 pmoreau: Yeah :-/
12:26 karolherbst: it's not hard to do that earlier though
12:27 karolherbst: it's just clover creating the cso too late
12:28 karolherbst: {ir_type = PIPE_SHADER_IR_NIR_SERIALIZED, prog = 0x4637d0, req_local_mem = 0, req_private_mem = 0, req_input_mem = 40}
12:28 karolherbst: are the input args
12:28 karolherbst: so...
12:28 karolherbst: only that req_local_mem part is tricky
12:28 karolherbst: and maybe req_private_mem?
12:28 karolherbst: dunno
12:29 karolherbst: pmoreau: anyway, I'd like to have the cso created when creating the kernel object, but that might require some clover changes actually
12:29 karolherbst: and then an API to query information about the cso, like reg count, etc...
12:30 pmoreau: clover currently always return 0 when the user queries for private_mem usage.
12:30 karolherbst: but I don't think drivers need to know any of that at compile time actually
12:31 karolherbst: pmoreau: it's useless information anyway
12:31 karolherbst: CL_KERNEL_LOCAL_MEM_SIZE is more relevant
12:31 pmoreau: It remains information that can be queried by the user
12:31 pmoreau: And yes, that other one is more relevant
12:32 karolherbst: returning always 0 is spec conformant :p
12:32 karolherbst: I think
12:32 pmoreau: And I just saw that it can include “local memory that may be needed by an implementation to execute the kernel”.
12:32 karolherbst: yeah
12:32 pmoreau: It is not, as far as I can read
12:32 karolherbst: it's good to know how much of local space you have left for that device
12:32 pmoreau: “Returns the minimum amount of private memory, in bytes, used by each work-item in the kernel.”
12:33 karolherbst: pmoreau: sure.. but it has no implications
12:33 karolherbst: with what would that help to know this number?
12:33 karolherbst: and how would a CTS can verify the value is correct
12:33 pmoreau: I guess some people might need it, otherwise it wouldn’t be in the spec?
12:34 karolherbst: for what do you need it?
12:34 karolherbst: the CL spec contains all sort of garbage nobody needs :p
12:35 karolherbst: CL_KERNEL_COMPILE_WORK_GROUP_SIZE is also super pointless...
12:35 karolherbst: the application already knows that
12:35 pmoreau: As a dev, I would be interested in knowing whether spilling occurred and private memory ends up being used. I do not care how much though, so just a boolean would be enough for me.
12:36 karolherbst: yeah.. just that it includes other things as well
12:36 pmoreau: True
12:36 karolherbst: spilling to VRAM, yes, that's useful information, but that's not what the spec can give you
12:36 karolherbst: I mean.. we also need to spill indirectly accessed arrays
12:37 karolherbst: but that's also spilling
12:38 pmoreau: “CL_KERNEL_COMPILE_WORK_GROUP_SIZE is also super pointless...” I could see it being useful if your app is using a library of OpenCL kernels and you do not necessarily know those values, but that’s about it
12:39 karolherbst: pmoreau: mhh.. actually if you compile from a binary that could be useful actually
18:10 Lyude: pmoreau: it very well might, although I have no idea where the MMIO read warningsare coming from
18:11 Lyude: pmoreau: can you get me another log with drm.debug=0x116 log_buf_len=50M nouveau.debug=disp=trace
18:11 pmoreau: Lyude: I am not the one with the issue, sjh is.
18:11 imirkin: mmio read warnings are emitted by the bus or whatever
18:12 imirkin: there's an intr which says what the issue was
18:12 Lyude: imirkin: yeah, but I don't think we see the interrupt without that last kernel parameter]
18:14 imirkin: mmio read/write failures are always reported as nvkm_error
18:14 Lyude: imirkin: no, I mean the interrupt from the evo channel that would have told us what the actual error was that triggered the timeout on the wndw channels
18:15 imirkin: ah, dunno
18:15 imirkin: i think errors still get logged pretty verbosely
18:15 imirkin: but perhaps not
18:15 imirkin: or not always
19:09 mixfix41: join #dri-devel
19:50 sjh: hi yes i can do that
22:41 haigedkinomehed: What you say means nothing to me, because i am an independent thinker, you tell me nothing what to do and how to do, especially even so when identified that most of you are really brain handycapped ones. So when i need something i always will completely ignore you, and to me it is important that communities like linux users have a good time on their platform despite the fact that you do not want the same thing.