ChanServ changed the topic of #dri-devel to: <ajax> nothing involved with X should ever be unable to find a bar
RAOF has quit [Ping timeout: 480 seconds]
RAOF_ is now known as RAOF
robmur01 has quit [Ping timeout: 480 seconds]
robmur01 has joined #dri-devel
a-865 has quit [Quit: ChatZilla 0.17 [SeaMonkey 2.53.17/20230727221859]]
q66 has quit [Quit: WeeChat 4.0.2]
Kayden has quit [Quit: params!!!!!!!]
Peuc has joined #dri-devel
<karolherbst> dcbaker: seems to be flaky..
<karolherbst> ehh wait
<karolherbst> there are also fails in the test I've written.. uhh
<karolherbst> huh.. something is weird
<karolherbst> heh.. that generated wrapper doesn't include the origina header file... maybe something changed with an updated bindgen
<karolherbst> I'll look into that tomorrow
columbarius has joined #dri-devel
co1umbarius has quit [Ping timeout: 480 seconds]
<karolherbst> and ubuntu "rolling" doesn't have a new enough bindgen...
<karolherbst> *sigh*
a-865 has joined #dri-devel
YuGiOhJCJ has joined #dri-devel
linkmauve has left #dri-devel [Disconnected: Replaced by new connection]
linkmauve has joined #dri-devel
<dcbaker> karolherbst: soon we’ll be able to build bindgen in the build process, lol
<HdkR> Oh wow, how long would that take?
<HdkR> That sounds like my idea of forcing people to rebuild a clang fork :D
<soreau> HdkR: upstream the patches, then you wont have to ;)
<HdkR> :)
q66 has joined #dri-devel
Danct12 is now known as Guest554
Danct12 has joined #dri-devel
neniagh_ has joined #dri-devel
yuq825 has joined #dri-devel
DPA has quit [Quit: ZNC 1.8.2+deb3.1 - https://znc.in]
DPA has joined #dri-devel
Kayden has joined #dri-devel
neniagh has quit [Ping timeout: 480 seconds]
Danct12 has quit [Quit: WeeChat 4.0.4]
Danct12 has joined #dri-devel
Danct12 has quit []
Danct12 has joined #dri-devel
Danct12 has quit [Quit: WeeChat 4.0.4]
Danct12 has joined #dri-devel
crabbedhaloablut has joined #dri-devel
<kurufu> Since anv exports yuv formats a single plane, is that essentially set in stone now? Asking as vulkan video makes that format more popular the single plane format seems to make lots of things harder (if only because the world expects 3 planes).
<Sachiel> what do you mean? The number of planes depends on the specific format, no?
aravind has joined #dri-devel
<kurufu> Despite the format being x planes, drivers dont neccessarily need to export a representation with multiple planes.
<Sachiel> An image’s memory planecount (as returned by drmFormatModifierPlaneCount) is distinct from its format planecount (in the sense of multi-planar Y′CBCR formats). In VkImageAspectFlags, each VK_IMAGE_ASPECT_MEMORY_PLANE_i_BIT_EXT represents a memory plane and each VK_IMAGE_ASPECT_PLANE_i_BIT a format plane.
cengiz_io has quit [Quit: Connection closed for inactivity]
<Company> fun fact: Disabling optimizations when generating spirv with glslang/glslc can result in significantly faster shader code
<Company> almost as fast as zink now
Duke`` has joined #dri-devel
neggles has quit [Quit: bye friends - ZNC - https://znc.in]
pekkari has joined #dri-devel
aravind has quit [Ping timeout: 480 seconds]
itoral has joined #dri-devel
tzimmermann has joined #dri-devel
Duke`` has quit [Ping timeout: 480 seconds]
kzd has quit [Ping timeout: 480 seconds]
neggles has joined #dri-devel
Company has quit [Quit: Leaving]
aravind has joined #dri-devel
kem has quit [Ping timeout: 480 seconds]
kem has joined #dri-devel
Danct12 has quit [Quit: WeeChat 4.0.4]
Danct12 has joined #dri-devel
mvlad has joined #dri-devel
Danct12 has quit []
Danct12 has joined #dri-devel
orbea1 has joined #dri-devel
orbea has quit [Ping timeout: 480 seconds]
pekkari has quit [Quit: Konversation terminated!]
<daniels> soreau: could you please file an MR?
pekkari has joined #dri-devel
pochu has joined #dri-devel
frieder has joined #dri-devel
pochu_ has joined #dri-devel
pochu has quit [Ping timeout: 480 seconds]
pochu_ has quit []
frieder has quit [Remote host closed the connection]
kts has joined #dri-devel
sghuge has quit [Remote host closed the connection]
sghuge has joined #dri-devel
qyliss has quit [Quit: bye]
kts has quit [Quit: Konversation terminated!]
qyliss has joined #dri-devel
qyliss has quit [Quit: bye]
qyliss has joined #dri-devel
frieder has joined #dri-devel
An0num0us has joined #dri-devel
kxkamil has quit []
lynxeye has joined #dri-devel
kxkamil has joined #dri-devel
pochu has joined #dri-devel
pekkari has quit [Quit: Konversation terminated!]
pekkari has joined #dri-devel
DodoGTA has quit [Quit: DodoGTA]
An0num0us has quit [Ping timeout: 480 seconds]
DodoGTA has joined #dri-devel
pcercuei has joined #dri-devel
vliaskov has joined #dri-devel
pekkari has quit [Quit: Konversation terminated!]
Danct12 has quit [Remote host closed the connection]
<soreau> daniels: Now that I thought about it, it needs zmike's MR to be merged first since it's against the egl branch of !24700
Danct12 has joined #dri-devel
sima has joined #dri-devel
_jannau__ is now known as _jannau_
Danct12 has quit [Quit: WeeChat 4.0.4]
consolers has joined #dri-devel
<soreau> daniels: also I'd like some feedback on fixing the crash http://pastie.org/p/6tqXQok1L7hJhIMv7mCQa9/raw with resizing weston-simple-egl regardless of vblank_mode/swap_interval http://ix.io/4GKs
heat_ has joined #dri-devel
An0num0us has joined #dri-devel
donaldrobson has joined #dri-devel
heat_ has quit [Remote host closed the connection]
heat_ has joined #dri-devel
idr_ has joined #dri-devel
idr has quit [Ping timeout: 480 seconds]
<soreau> in that trace, wsi_wl_surface is NULL because it's using the old swapchain after it has been replaced with a new swapchain but before it was pruned
<soreau> attempts to simply prune before it was used were unsuccessful
aravind has quit [Ping timeout: 480 seconds]
hansg has joined #dri-devel
<karolherbst> HdkR: building bindgen takes like a minute
Danct12 has joined #dri-devel
Danct12 has quit [Quit: WeeChat 4.0.4]
heat_ has quit [Remote host closed the connection]
dviola has quit [Quit: WeeChat 4.0.4]
heat_ has joined #dri-devel
Danct12 has joined #dri-devel
dviola has joined #dri-devel
<karolherbst> dcbaker: okay.. so the arch container simply had bindgen 0.64 and there the static inline stuff is mostly broken..
mripard has joined #dri-devel
heat_ has quit [Read error: Connection reset by peer]
heat_ has joined #dri-devel
kts has joined #dri-devel
pekkari has joined #dri-devel
kts has quit [Quit: Leaving]
qyliss has quit [Quit: bye]
wideopen has joined #dri-devel
qyliss has joined #dri-devel
<wideopen> The solver itself has all the documentation, it's incremental solving through the server, there is more room to extensions, so if this is a code point like variable in clause i.e constraint based thingy, it tests it, and takes on a assume path then continues and does it again (the test), there is very little that needs a change. It's all done cause bitwuzla has softfloat core too, so mmiotrace through the solver should be also very
<wideopen> easy, you just add all the integers to be polled, and you maximize them and annotate a name to them. What squad psykose talks about their anal tranny stalkers, as all the last ones alike have been reported missing teeth and likely if they have anything other to bully or violate more, they will report a missing bully altogether, so i care none about your squad, and idiots or brain dead he was in the moment he was born, that has nothing
<wideopen> to do with me , only thing that has to do with me, is that bully stalks my businesses along with many others, and since my hand bone was broken 2012 years ago, there is a set of people who will crucify such, i.e beat up , lock up and if more problems are met, then already wheelchair etc. At overseas legal powers such as Khmer army handled that, but i can offer some wrestlers to help them. In the meantime this commune is a filthy trash,
<wideopen> i am better off anyway not to deal with you. That i am bio-chiped is another thing i am aware about, but this battle i win, and i say upfront i can elect something to tattoo on your butts. And my choice is "Decade long Markov Chain abusive stalker anal house parrot"
<wideopen> good bye
wideopen has quit [Quit: Leaving]
camus has quit [Remote host closed the connection]
camus has joined #dri-devel
<daniels> soreau: the patch looks reasonable to me, but it would need zmike to comment on it since the flush_queue is zink-specific
<soreau> daniels: unfortunately, zmike said he can't reproduce the resize crash
<soreau> but no telling what sort of dragons he has installed ;)
hansg has quit [Quit: Leaving]
<soreau> or maybe it's only on radv? idk which chip he tried
<daniels> right, but he's the one who knows whether flushing the flush queue at that point is theoretically correct or not
<daniels> I didn't even know zink had its own internal queue of stuff
Danct12 has quit [Quit: WeeChat 4.0.4]
qyliss has quit [Quit: bye]
qyliss has joined #dri-devel
aravind has joined #dri-devel
idr_ has quit [Ping timeout: 480 seconds]
qyliss has quit [Quit: bye]
qyliss has joined #dri-devel
kts has joined #dri-devel
consolers has quit [Ping timeout: 480 seconds]
pekkari has quit [Quit: Konversation terminated!]
pochu has quit [Quit: leaving]
jkrzyszt has joined #dri-devel
itoral has quit [Quit: Leaving]
<Wallbraker> Does https://docs.mesa3d.org/envvars.html#envvar-MESA_VK_TRACE do anything? Is it a newish variable?
aravind has quit [Read error: Connection reset by peer]
<pixelcluster> Wallbraker: it does something on 23.2 and newer
<pixelcluster> the different trace modes have had different environment variables to enable before that, but docs are generated from latest main
<Wallbraker> Ah yes, I'm on 23.1.7, that clears that up. Really those things should have version tags so you know for which version they apply to.
jdavies has joined #dri-devel
jdavies is now known as Guest625
Guest625 has quit [Ping timeout: 480 seconds]
cmichael has joined #dri-devel
<zmike> soreau: I slept on it and I think https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25288 is what you need
donaldrobson has quit [Ping timeout: 480 seconds]
<soreau> zmike: Thanks, I'll try it. Were you ever able to reproduce the resize crash? (on radv?)
<zmike> eventually
donaldrobson has joined #dri-devel
<soreau> super good
qyliss has quit [Quit: bye]
kts has quit [Ping timeout: 480 seconds]
qyliss has joined #dri-devel
<DavidHeidelberg[m]> running GPT-2 on iris (TGL): ~ 112ms per query; llvmpipe (i7-1185G7 @ 3.00GHz): ~150ms (but CPU heats up quickly, so power consumption would be very different)
<DavidHeidelberg[m]> outcome: OpenCL may be not best, but still better than CPU on integrated GPU with shared memory for LLVM
<DavidHeidelberg[m]> *LLM damn
yyds has joined #dri-devel
<soreau> zmike: yes it works
<soreau> zmike: did you get a chance to glance at the swap interval patch?
idr has joined #dri-devel
<karolherbst> DavidHeidelberg[m]: wondering how good it runs on Intel's stack
<karolherbst> but yeah..
<karolherbst> on Intel's iGPUs it's mostly about power efficiency
<karolherbst> and.. there might be some overhead randomly :D
<karolherbst> always hard to tell with tests running quickly
kts has joined #dri-devel
<zmike> soreau: no, but I'm also probably not the ultimate reviewer for such a thing
<DavidHeidelberg[m]> karolherbst: RUSTICL_ENABLE=iris GPU=1 PYTHONPATH="." JIT=1 python examples/gpt2.py --model_size=gpt2 --prompt "Hello." --count 46 --temperature 0 --timing
<karolherbst> heh...
<karolherbst> DavidHeidelberg[m]: what repo do I need to clone?
<karolherbst> ahh, it's tinygrad
<soreau> zmike: ok
<karolherbst> DavidHeidelberg[m]: looks like intel is a bit faster :D but I'm also using a debug build
<karolherbst> What is Rusticl?: Rusticl is a library for building and manipulating Rusticl objects. It is a library for building and manipulating Rusticl objects. It is a library for building and manipulating Rusticl objects. It is a
<karolherbst> oh well..
<karolherbst> but that question breaks with rusticl
<karolherbst> "ValueError: probabilities do not sum to 1" :D
<karolherbst> guess I might want to fix that
<karolherbst> works with llvmpipe...
<DavidHeidelberg[m]> Interesting :D
<DavidHeidelberg[m]> Usually so far rusticl w/ iris worked best for me
<karolherbst> I break it way too easily
orbea1 has quit []
<karolherbst> but yeah...
<karolherbst> I'm sure it's a 0.9999 vs 1.0 thing
<karolherbst> and some precision is slightly off
orbea has joined #dri-devel
<karolherbst> let's see...
<karolherbst> yeah.. works with count 45 :D
<DavidHeidelberg[m]> Iris doing trolololo
qyliss has quit [Quit: bye]
<DavidHeidelberg[m]> Which gen?
<karolherbst> but yeah.. intels stack is significantly faster.. I guess there is something I should optimize then
<karolherbst> 9.5
<karolherbst> "Rusticl is an EULA agreement-based and open source project which aims to provide a framework for building web apps based on Rust. This means that the project is open source."
<daniels> well, those are certainly all words
<karolherbst> it's fascinating how sure it is and how much garbage it is saying
<karolherbst> *sure of itself"
heat_ has quit [Remote host closed the connection]
<tnt> mmm, I just checked pocl/rusticl iris/rusticl llvmpipe/intel compute and yeah, rusticl iris is quite a bit slower than the 3 others for some reason.
heat_ has joined #dri-devel
qyliss has joined #dri-devel
<karolherbst> probably something debug build
<DavidHeidelberg[m]> it's GPT-2, with llama (but it takes 13G of VRAM) it's "sometimes" reasonable
<karolherbst> clinfo is also 3x quicker with release builds
<DavidHeidelberg[m]> yeah, I'm using Debian nightly builds (so release)
<karolherbst> ahh
<tnt> I should be on a release build too
<karolherbst> anyway.. I'm sure there are optional CL extensions or something
<karolherbst> or the runtime overhead is just too high
<karolherbst> but glad to know that it also produces same nonsense with Intel
<karolherbst> so the math seems to be alright
qyliss has quit []
qyliss has joined #dri-devel
<karolherbst> nice.. seems to work with zink as well... kinda
<karolherbst> I kinda have to find a solution for the gpu getting reset too quickly
<karolherbst> and I get tons of "MESA: error: zink: couldn't allocate memory: heap=4 size=20102"
<karolherbst> oh well...
kzd has joined #dri-devel
<karolherbst> okay.. yeah it's mostly GPU side things ...
<karolherbst> and some memcpies, but whatever
<karolherbst> 83% is just the python runtime
<karolherbst> GPU is 90% busy, so there is that
yuq825 has left #dri-devel [#dri-devel]
<karolherbst> and it's always creating a context with profiling enabled.. oh boi.. anyway, I suspect we can be more optimized on the kernel side somewhere
linusw has quit [Read error: Network is unreachable]
olv has quit [Read error: Network is unreachable]
rg3igalia has quit [Read error: Network is unreachable]
steev has quit [Read error: Network is unreachable]
hashar has quit [Read error: Network is unreachable]
linusw has joined #dri-devel
daniels has quit [Read error: Network is unreachable]
olv has joined #dri-devel
rg3igalia has joined #dri-devel
robher has quit [Read error: Network is unreachable]
kerneltoast has quit [Write error: connection closed]
arnd has quit [Read error: Network is unreachable]
austriancoder has quit [Read error: Network is unreachable]
jimjams has quit [Read error: Network is unreachable]
hfink has quit [Read error: Network is unreachable]
i509vcb has quit [Read error: Network is unreachable]
mdnavare has quit [Read error: Network is unreachable]
markco has quit [Read error: Network is unreachable]
hashar has joined #dri-devel
zzag has quit [Read error: Network is unreachable]
hfink has joined #dri-devel
mdnavare has joined #dri-devel
i509vcb has joined #dri-devel
markco has joined #dri-devel
rib has quit [Read error: Network is unreachable]
zzag has joined #dri-devel
kode54 has quit [Read error: Network is unreachable]
rodrigovivi has quit [Read error: Network is unreachable]
austriancoder has joined #dri-devel
eric_engestrom has quit [Read error: Network is unreachable]
zx2c4 has quit [Write error: connection closed]
zmike has quit [Read error: Network is unreachable]
dianders has quit [Read error: Network is unreachable]
dschuermann has quit [Read error: Network is unreachable]
steev has joined #dri-devel
daniels has joined #dri-devel
rib has joined #dri-devel
steve--w has quit [Read error: Network is unreachable]
rodrigovivi has joined #dri-devel
kode54 has joined #dri-devel
zmike has joined #dri-devel
zx2c4 has joined #dri-devel
dianders has joined #dri-devel
eric_engestrom has joined #dri-devel
steve--w has joined #dri-devel
jimjams has joined #dri-devel
kerneltoast has joined #dri-devel
robher has joined #dri-devel
arnd has joined #dri-devel
dschuermann has joined #dri-devel
aswar002 has quit [Remote host closed the connection]
alyssa has quit [Quit: alyssa]
aswar002 has joined #dri-devel
a-865 has quit [Ping timeout: 480 seconds]
YuGiOhJCJ has quit [Quit: YuGiOhJCJ]
yyds has quit [Remote host closed the connection]
a-865 has joined #dri-devel
pekkari has joined #dri-devel
karolherbst has quit [Ping timeout: 480 seconds]
alyssa has joined #dri-devel
heat_ has quit [Remote host closed the connection]
heat_ has joined #dri-devel
soreau has quit [Ping timeout: 480 seconds]
heat_ has quit [Remote host closed the connection]
heat has joined #dri-devel
Duke`` has joined #dri-devel
soreau has joined #dri-devel
karolherbst has joined #dri-devel
frieder has quit [Remote host closed the connection]
Company has joined #dri-devel
kzd has quit [Quit: kzd]
<karolherbst> mhhh.. maybe I have some synchronization bug somewhere... zink is also hitting something all the time
pekkari has quit [Quit: Konversation terminated!]
tzimmermann has quit [Quit: Leaving]
hansg has joined #dri-devel
pekkari has joined #dri-devel
alyssa has quit [Quit: alyssa]
jkrzyszt has quit [Ping timeout: 480 seconds]
pannage has joined #dri-devel
cmichael has quit [Quit: Leaving]
<karolherbst> oof... ralloc_asprintf accounts for 33% of my launch kernel overhead...
lynxeye has quit [Quit: Leaving.]
jrpan has joined #dri-devel
<karolherbst> "SIMD16 skipped because workgroup size 1 already fits in SIMD8" getting tons of those
<karolherbst> or "SIMD32 skipped because workgroup size 16 already fits in SIMD16"
<karolherbst> Kayden: ^^ seems like this is a significant CPU overhead for launching compute jobs
<karolherbst> like.. the biggest part of it
<karolherbst> might make sense t skip those ralloc_asprintf calls
<karolherbst> inside brw_simd_should_compile that is
<jrpan> Hi, I'm new to graphics rendering and drivers and please forgive me for asking dumb questions. I've been playing with the vulkan intel driver. I see that when anv_cmd_buffer_bind_descriptor_set is called, the descriptor is saved to the cmd_buffer at cmd_buffer->state.gfx.base->descriptors[set_index].
<jrpan> If there are multiple draws in a command buffer, and for each draw, new descriptors are binded, how are the desciptors being distinguished between drawcalls that are within the same command buffer (when the buffer is submitted to GPU)?
<jrpan> Becuase it's the same command buffer and the later binded descriptor would just overwrite previous binded descriptor. Or my understanding is just wrong?
<jrpan> I just want to reference all descriptors that are being used in a command buffer at queuesubmit. But so far I can only get the last descriptor used.
<pendingchaos> mostly likely vkCmdDraw/etc read from state.gfx.base->descriptors to create commands that actually bind the descriptors
<jrpan> So at vkCmdDraw, the state.gfx.base->descriptors is being "saved" to the vkcmdDraw?
<Sachiel> anv_CmdBindPipeline in anv_cmd_buffer.c takes care of that
flom84 has joined #dri-devel
pekkari has quit [Quit: Konversation terminated!]
<mareko> tarceri_: I've noticed that gl_program::sh::UniformBlocks is populated before UBO linking for GLSL, but not SPIR-V. If I add UBOs before UBO linking, do I also need to update gl_program::sh::UniformBlocks?
pannage has quit [Ping timeout: 480 seconds]
<Kayden> karolherbst: that isn't handled by shader cache?
jkrzyszt has joined #dri-devel
gouchi has joined #dri-devel
gouchi has quit []
<karolherbst> Kayden: nope, that's all at the SIMD selection level
<karolherbst> Kayden: iris_launch_grid -> iris_upload_compute_state -> iris_upload_compute_walker -> brw_cs_get_dispatch_info -> brw_simd_select_for_workgroup_size -> brw_simd_should_compile -> ralloc_asprintf
sarnex has quit [Ping timeout: 480 seconds]
RSpliet has quit [Remote host closed the connection]
<karolherbst> guess that's only hit if the workgroup size is variable...
jkrzyszt has quit [Ping timeout: 480 seconds]
RSpliet has joined #dri-devel
sarnex has joined #dri-devel
zf` has joined #dri-devel
zf` has quit []
An0num0us has quit [Ping timeout: 480 seconds]
pochu has joined #dri-devel
flom84 has quit [Ping timeout: 480 seconds]
junaid has joined #dri-devel
DemiMarie has left #dri-devel [#dri-devel]
donaldrobson has quit [Ping timeout: 480 seconds]
jrpan has quit [Quit: Page closed]
ngcortes has joined #dri-devel
<karolherbst> it kinda feels all a bit suboptimal
<karolherbst> but I can't really put my finger on it yet
<karolherbst> ehh
<karolherbst> nvm, maybe I should check what kernels are actually used often
<karolherbst> but I'm wondering if there are any magic intel isntructions we might want to wire up to speed up matrix multiplications
<karolherbst> I kinda need GPU perf profiling tools...
jrpan has joined #dri-devel
pannage has joined #dri-devel
vliaskov has quit [Remote host closed the connection]
<anholt> hakzsam: I've got some wip ci stuff that I'd like to have the 1.3.6 cts for. are you still working on polishing that today?
melonai33 has joined #dri-devel
melonai3 has quit [Ping timeout: 480 seconds]
jrpan has quit [Remote host closed the connection]
nchery has joined #dri-devel
melonai339 has joined #dri-devel
TMM has quit [Ping timeout: 480 seconds]
TMM has joined #dri-devel
melonai33 has quit [Ping timeout: 480 seconds]
sima has quit [Ping timeout: 480 seconds]
<anholt> hakzsam: I guess we're waiting for !25284 for cts uprev?
kzd has joined #dri-devel
fab has quit [Quit: fab]
pcercuei has quit [Quit: brb]
wideopen has joined #dri-devel
<wideopen> https://github.com/JonathanSalwan/Triton/issues/284 https://github.com/jaredsofteng/gini/blob/master/doc/crisp/crisp.pdf those are the docs that show how to run those things, and no i did not test, but the crisp part is in fact encoding the models with very high compression, which is technically the maximum of 10bit elias fano too coincidentally , probably something like this they use, only DMA and gpu backends can be written, it's not
<wideopen> a lot of work. You need to read through a paragraph of 2.10 models, where it is described, not that varint encoding, they went so far, that they encode wire transfers too maximum ways, where they send the compressed model. I mean the code is round about already there, French people wrote the thing, triton concrete and bitwuzla cnf with gini backend. I am at my last year with computers, i do not care what you do, if you bother me again
<wideopen> on my territory people will shoot you. You are trash.
wideopen has quit [Quit: Leaving]
hansg has quit [Remote host closed the connection]
pcercuei has joined #dri-devel
heat has quit [Read error: No route to host]
heat has joined #dri-devel
pannage has quit [Remote host closed the connection]
junaid has quit [Remote host closed the connection]
heat_ has joined #dri-devel
alanc has quit [Remote host closed the connection]
alanc has joined #dri-devel
heat has quit [Read error: No route to host]
heat_ has quit [Read error: No route to host]
heat has joined #dri-devel
vyivel has quit [Read error: Connection reset by peer]
vyivel has joined #dri-devel
pcercuei has quit [Quit: brb]
pcercuei has joined #dri-devel
Duke`` has quit [Ping timeout: 480 seconds]
pcercuei has quit []
konstantin_ has joined #dri-devel
konstantin has quit [Ping timeout: 480 seconds]
heat has quit [Remote host closed the connection]
heat has joined #dri-devel
illwieckz_ has joined #dri-devel
illwieckz has quit [Read error: Connection reset by peer]
illwieckz_ has quit []
illwieckz has joined #dri-devel
pcercuei has joined #dri-devel
qyliss has quit [Quit: bye]
qyliss has joined #dri-devel
<karolherbst> okay
<karolherbst> I have a fun idea for an optimization of compute kernels
<karolherbst> so all threads execture the lower loop, which is doing the exact same thing in all threads
<karolherbst> and stores the same result at the same location in all threads
<karolherbst> I wonder if we could just vectorize or optimize that code in a way, that it's not doing something silly like this
An0num0us has joined #dri-devel
mvlad has quit [Remote host closed the connection]
<karolherbst> mhhh
<karolherbst> this entire loop could be an atomic... actually.. the entire shared memory array could be one atomic
<karolherbst> I wonder how terrible it would be to match this pattern...
illwieckz has quit [Remote host closed the connection]
<airlied> uggh yeah I wonder if tinygrad itself could do better in it's generator code there
<karolherbst> the pain point is that it is shared code with the CPU stuff...
<karolherbst> somehow
<karolherbst> but I do wonder if we actually just want to match that pattern, because it's actually not _that_ complicated
<karolherbst> the loop is uniform, which we should be able to proof
<karolherbst> and if every thread stores it at the same location, we should do something smarter
<karolherbst> yeah
<karolherbst> `con 32 %54 = fmul! %62, %53 (0.000781)`
<karolherbst> so yeah.. nir detects it as uniform
<karolherbst> the biggest question is simply to what should we optimize this
illwieckz has joined #dri-devel
<karolherbst> maybe subgroup ops and just cut that loop by the subgroup size?
<karolherbst> but yeah.. if LLVM is able to optimize this to something not stupid, no wonder Intel is faster
An0num0us has quit [Ping timeout: 480 seconds]
illwieckz has quit [Remote host closed the connection]
crabbedhaloablut has quit []
pcercuei has quit [Quit: dodo]
illwieckz has joined #dri-devel
shashanks_ has joined #dri-devel
shashanks__ has quit [Ping timeout: 480 seconds]
illwieckz has quit [Remote host closed the connection]
illwieckz has joined #dri-devel
alyssa has joined #dri-devel
JohnnyonFlame has joined #dri-devel
JohnnyonFlame has quit []
gio has quit [Remote host closed the connection]
gio has joined #dri-devel
Company has quit [Remote host closed the connection]
Company has joined #dri-devel
Company has quit [Remote host closed the connection]
Company has joined #dri-devel