Why do you even need a mpmc queue? [Thunder Kittens] (Was: Do you see the loops, in C code and in Java code? / Looping /)

From	Mild Shock <janburse@fastmail.fm>
Newsgroups	sci.logic, comp.lang.prolog, sci.physics
Subject	Why do you even need a mpmc queue? [Thunder Kittens] (Was: Do you see the loops, in C code and in Java code? / Looping /)
Date	2026-07-24 14:44 +0200
Message-ID	<113vmnr$bed5$4@solani.org> (permalink)
References	(5 earlier) <113mco7$1j9n2$1@dont-email.me> <113n4m3$5oih$1@solani.org> <113n5j3$5ouq$3@solani.org> <113r9gn$396f7$4@dont-email.me> <113rf37$846i$1@solani.org>

Cross-posted to 3 groups.

Show all headers | View raw

Hi,

You don't pay attention, right! I am little
bit disappointed that your attention span is
near zero. I already posted:

From: Mild Shock <janburse@fastmail.fm>
Subject: Why do you even need a mpmc queue? [Thunder Kittens]
Date: Thu, 23 Jul 2026 08:43:03 +0200

 > Hi,
 >
 > Because I use WebGPU and not WebGL. And
 > because WebGPU can adresss modern GPU
 > developed with the NVIDIA Volta evolution,
 >
 > which happened in 2017. Namley that compute
 > shaders are not any more subject to the
 > realization restriction of lock step
 >
 > execution, but have independent thread state.
 > And because there is independent thread state
 > there is also independent time spent for a
 >
 > a work item by each logical thread, if the
 > submitted logical thread uses a lot of branching
 > logic or even loops. But the use of branching
 >
 > and loops is encouraged in independent thread
 > state programming of compute shaders. The variables
 > that can drive such logic are the scalar variables:
 >
 > Tour of WGSL - Control Flow
 > https://google.github.io/tour-of-wgsl/control-flow/
 >
 > Then not to waste GPU compute time, by logical
 > threads doing nothing. You will need to
 > introduce some load balancing among multiple
 >
 > logical threads. And MPMC queues are one way to
 > readize load balancing. Compute shaders with
 > producer and consumer entry points are proposed
 >
 > as fundamental architecture by Thunder Kittens:
 >
 > ThunderKittens: Simple, Fast, and Adorable AI Kernels
 > https://arxiv.org/abs/2410.20399
 >
 > They are used by this SpaceX acquisition:
 >
 > Composer 2 Technical Report
 > https://arxiv.org/abs/2603.24477
 >
 > Thunder Kittens uses Hardware support, i.e. tma_expect().
 >
 > Bye
 >
 > Chris M. Thomasson schrieb:
 >>> never meant to be used in a GPU.
 >>> Dmitry CAS version can be used, but
 >>>
 >>> Why do you even need a mpmc queue
 >>> in your compute shader anyway?

Chris M. Thomasson schrieb:
 > I don't think he knows exactly what he is doing...
 > Why does he need a lock/wait-free queue in a
 > compute shader? What is he trying to do?

Thread

I'm a spinner, I'm a sinner [Dmitry Vyukov for pi-WAM] (Re: Paul Tarau versus Mr. Taskmanager, who would win? [A PDP-11 Humunkulus from 1979] Mild Shock <janburse@fastmail.fm> - 2026-07-19 11:55 +0200
  Gemini, DeepSeek, OpenAI all know Dmitriy V'jukov (Re: I'm a spinner, I'm a sinner [Dmitry Vyukov for pi-WAM]) Mild Shock <janburse@fastmail.fm> - 2026-07-20 08:33 +0200
    The Cache Identity Crisis by Micro Penis (Re: Gemini, DeepSeek, OpenAI all know Dmitriy V'jukov) Mild Shock <janburse@fastmail.fm> - 2026-07-20 13:54 +0200
      Just RTFM the RDNA 3.5 specs! [GPU Cache Lines] (Re: The Cache Identity Crisis by Micro Penis) Mild Shock <janburse@fastmail.fm> - 2026-07-20 14:10 +0200
        The large memory tax: ECC RAM (Re: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines]) Mild Shock <janburse@fastmail.fm> - 2026-07-20 14:23 +0200
        Re: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines] (Re: The Cache Identity Crisis by Micro Penis) Ross Finlayson <ross.a.finlayson@gmail.com> - 2026-07-20 12:26 -0700
          Re: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines] (Re: The Cache Identity Crisis by Micro Penis) R Kym Horsell <kymhorsell@gmail.com> - 2026-07-20 22:20 +0000
            Friendly Reminder: GPU 10x more performant than CPU (Was: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines]) Mild Shock <janburse@fastmail.fm> - 2026-07-21 00:39 +0200
              Breaking the CUDA edge in AI by WebGPU (Was: Friendly Reminder: GPU 10x more performant than CPU) Mild Shock <janburse@fastmail.fm> - 2026-07-21 00:56 +0200
                Like WebAssembly before it, WebGPU has "escaped" the browser. (Was: Breaking the CUDA edge in AI by WebGPU) Mild Shock <janburse@fastmail.fm> - 2026-07-21 01:05 +0200
                If you are paranoid you can use Falco [Agentic AI] (Re: Like WebAssembly before it, WebGPU has "escaped" the browser.) Mild Shock <janburse@fastmail.fm> - 2026-07-21 22:57 +0200
                What would an EMACS guru say [Windows Recall] (Re: If you are paranoid you can use Falco [Agentic AI]) Mild Shock <janburse@fastmail.fm> - 2026-07-21 23:35 +0200
                Decide what you critique tiny winy penis (Re: What would an EMACS guru say [Windows Recall]) Mild Shock <janburse@fastmail.fm> - 2026-07-22 08:17 +0200
                How confused is tiny winy penis? (Was: Decide what you critique tiny winy penis) Mild Shock <janburse@fastmail.fm> - 2026-07-22 08:28 +0200
                Maybe change your hobby, become a dog owner? (Was: How confused is tiny winy penis?) Mild Shock <janburse@fastmail.fm> - 2026-07-22 09:37 +0200
                Even dogs know Switzerland != Germany [Syphilis Brain Micro Penis] (Re: Decide what you critique tiny winy penis ) Mild Shock <janburse@fastmail.fm> - 2026-07-22 11:19 +0200
                node.js has also Worker isolation , headless (Was: Even dogs know Switzerland != Germany [Syphilis Brain Micro Penis]) Mild Shock <janburse@fastmail.fm> - 2026-07-22 11:29 +0200
                Abyss hobby interpretation vs. professional understanding [Keep up with the Kardashians] (Re: node.js has also Worker isolation , headless) Mild Shock <janburse@fastmail.fm> - 2026-07-22 11:38 +0200
                Play stupid games, win classic Usenet prizes [CCCP Troll] (Re: Abyss hobby interpretation vs. professional understanding [Keep up with the Kardashians]) Mild Shock <janburse@fastmail.fm> - 2026-07-22 12:16 +0200
                Glue Sniffing 5-Year Old Moron [CCCP Troll] (Was: Play stupid games, win classic Usenet prizes [CCCP Troll]) Mild Shock <janburse@fastmail.fm> - 2026-07-22 14:29 +0200
                My Swift Go 16 AI has no IMEI, are you nuts? (Re: Even dogs know Switzerland != Germany [Syphilis Brain Micro Penis]) Mild Shock <janburse@fastmail.fm> - 2026-07-22 14:15 +0200
                Same nickname and email, could post faster [5 year old moron] (Re: My Swift Go 16 AI has no IMEI, are you nuts? ) Mild Shock <janburse@fastmail.fm> - 2026-07-22 14:25 +0200
                I have nothing to hide, you can find me in search.ch (Re: My Swift Go 16 AI has no IMEI, are you nuts?) Mild Shock <janburse@fastmail.fm> - 2026-07-22 14:41 +0200
                Where did I confirm German via .ch, you are more than nuts! (Re: I have nothing to hide, you can find me in search.ch) Mild Shock <janburse@fastmail.fm> - 2026-07-22 14:48 +0200
                Ask a Ukrainian Neighbour to do Detective [CCCP Troll] (Re: Where did I confirm German via .ch, you are more than nuts! ) Mild Shock <janburse@fastmail.fm> - 2026-07-22 14:51 +0200
          Re: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines] (Re: The Cache Identity Crisis by Micro Penis) Ross Finlayson <ross.a.finlayson@gmail.com> - 2026-07-20 21:36 -0700
            rossy boy is going paranoid (Was: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines]) Mild Shock <janburse@fastmail.fm> - 2026-07-21 09:01 +0200
              Re: rossy boy is going paranoid (Was: Just RTFM the RDNA 3.5 specs! [GPU Cache Lines]) Ross Finlayson <ross.a.finlayson@gmail.com> - 2026-07-21 00:56 -0700
      L1,..,Ln caches are located on the CPU AND on the GPU (Was: The Cache Identity Crisis by Micro Penis( Mild Shock <janburse@fastmail.fm> - 2026-07-20 19:35 +0200
        GPU Cache Hierarchy: Understanding L1, L2, and VRAM (Was: L1,..,Ln caches are located on the CPU AND on the GPU) Mild Shock <janburse@fastmail.fm> - 2026-07-20 19:40 +0200
          Well thats good, co-location, onto the same processor die (Re: GPU Cache Hierarchy: Understanding L1, L2, and VRAM (Was: L1,..,Ln caches are located on the CPU AND on the GPU) Mild Shock <janburse@fastmail.fm> - 2026-07-20 22:02 +0200
            Where is micro penis mental error? (Re: Well thats good, co-location, onto the same processor die (Re: GPU Cache Hierarchy: Understanding L1, L2, and VRAM (Was: L1,..,Ln caches are located on the CPU AND on the GPU) Mild Shock <janburse@fastmail.fm> - 2026-07-20 22:07 +0200
    I didn't find Futex in WebGPU / WGSL (Re: Gemini, DeepSeek, OpenAI all know Dmitriy V'jukov) Mild Shock <janburse@fastmail.fm> - 2026-07-20 23:26 +0200
      There is no imageAtomicAdd in WGSL (Re: I didn't find Futex in WebGPU / WGSL) Mild Shock <janburse@fastmail.fm> - 2026-07-21 00:15 +0200
        OpenGL is dead. Apple said bye bye / Wayland Compositor (Re: There is no imageAtomicAdd in WGSL) Mild Shock <janburse@fastmail.fm> - 2026-07-21 00:26 +0200
        Flogging a Dead Horse, OpenGL is EOL (Re: There is no imageAtomicAdd in WGSL) Mild Shock <janburse@fastmail.fm> - 2026-07-21 01:39 +0200
          imageAtomicAdd trivial, Dmitry Vyukov requires capacity (Re: Flogging a Dead Horse, OpenGL is EOL) Mild Shock <janburse@fastmail.fm> - 2026-07-21 01:40 +0200
            capacity = 2^n for some n / systolic system (Re: imageAtomicAdd trivial, Dmitry Vyukov requires capacity) Mild Shock <janburse@fastmail.fm> - 2026-07-21 01:41 +0200
              Source of the benchmark for DmitryVyukov (Was: capacity = 2^n for some n / systolic system) Mild Shock <janburse@fastmail.fm> - 2026-07-21 01:44 +0200
        Re: There is no imageAtomicAdd in WGSL (Re: I didn't find Futex in WebGPU / WGSL) "Chris M. Thomasson" <chris.m.thomasson.1@gmail.com> - 2026-07-20 16:59 -0700
          I never used OpenGL Version 4.2 and later (Was: There is no imageAtomicAdd in WGSL) Mild Shock <janburse@fastmail.fm> - 2026-07-21 08:47 +0200
            Because of MIMD you have to reassess algorithms (Re: I never used OpenGL Version 4.2 and later) Mild Shock <janburse@fastmail.fm> - 2026-07-21 09:02 +0200
              Why MIMD is interesting for pi-WAM? Mild Shock <janburse@fastmail.fm> - 2026-07-21 09:18 +0200
              Re: Because of MIMD you have to reassess algorithms (Re: I never used OpenGL Version 4.2 and later) "Chris M. Thomasson" <chris.m.thomasson.1@gmail.com> - 2026-07-22 13:34 -0700
                Do you see the loops, in C code and in Java code? /** Looping **/ (Was: Because of MIMD you have to reassess algorithms) Mild Shock <janburse@fastmail.fm> - 2026-07-23 00:09 +0200
                Re: Do you see the loops, in C code and in Java code? /** Looping **/ (Was: Because of MIMD you have to reassess algorithms) Mild Shock <janburse@fastmail.fm> - 2026-07-23 00:53 +0200
                Re: Do you see the loops, in C code and in Java code? /** Looping **/ (Was: Because of MIMD you have to reassess algorithms) "Chris M. Thomasson" <chris.m.thomasson.1@gmail.com> - 2026-07-22 18:00 -0700
                Re: Do you see the loops, in C code and in Java code? /** Looping **/ (Was: Because of MIMD you have to reassess algorithms) "Chris M. Thomasson" <chris.m.thomasson.1@gmail.com> - 2026-07-22 18:20 -0700
                Why do you even need a mpmc queue? [Thunder Kittens] (Was: Do you see the loops, in C code and in Java code? /** Looping **/) Mild Shock <janburse@fastmail.fm> - 2026-07-24 14:44 +0200
                What does pi in pi-WAM mean? (Re: Why do you even need a mpmc queue? [Thunder Kittens]) Mild Shock <janburse@fastmail.fm> - 2026-07-24 14:53 +0200
                OR-parallelism or AND-parallelism? [MapReduce] (Was: What does pi in pi-WAM mean?) Mild Shock <janburse@fastmail.fm> - 2026-07-24 15:00 +0200

csiph-web