Groups | Search | Server Info | Keyboard shortcuts | Login | Register [http] [https] [nntp] [nntps]

Groups > sci.physics.relativity > #671597

Why do you even need a mpmc queue? [Thunder Kittens] (Re: Deadlock Exorcism: Switch from Push to Pull)

Cross-posted to 2 groups.

Show all headers | View raw

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?

Mild Shock schrieb:
> Hi,
> 
> This is quite fun, how some TLA+ guy fears
> the full state of queue like the devil in
> itself. But I guess if a service rate is
> 
> low and the producer has not much to do to
> produce its work items, the arrival rate
> has nevertheless to adapt, and dealing
> 
> with "full states", which are wrongly
> called deadlock here, is the normal:
> 
> Tutorial-style talk - BlockingQueue
> https://github.com/lemmy/BlockingQueue/tree/main
> 
> Prolog is in good position. The bird box
> model has a redo port. So sometimes switching
> from push to pull, can help without doing
> 
> Deadlock Exorcism. You can also translate
> the bird box ports into pi-calculus:
> 
> A pi-calculus Specification of Prolog
> Benjamin Z. Li - University of Pennsylvania
> 11 Apr 1994, European Symposium on Programming,
> Prolog, Unification, Backtracking
> https://scispace.com/pdf/a-pi-calculus-specification-of-prolog-3qf2pf04ud.pdf 
> 
> 
> Have Fun!
> 
> Bye
> 
> Mild Shock schrieb:> Hi,
>  >
>  > CAS and XADD have no looping, they
>  > are atomic operations, that take some
>  > time but basically have some outcome
>  >
>  > with some ACID property and a result
>  > value. What loops is the ADT, the Abstract
>  > Data Type that you implement. Respectively
>  >
>  > the client that uses the Abstract Data Type.
>  > In your case you added the loop inside the
>  > Abstract Data Type or lower level aggregate
>  >
>  > code of a higher level operation:
>  >
>  > Chris M. Thomasson wrote:
>  > void producer(double state) {
>  >      uint32_t ver = XADD(&head, 1);
>  >      cell& c = cells[ver & (N - 1)];
>  >      while (LOAD(&c.ver) != ver) backoff(); /** Looping **/
>  >      c.state = state;
>  >      STORE(&c.ver, ver + 1);
>  > }
>  > https://groups.google.com/g/lock-free/c/acjQ3-89abE/m/a6-Di0GZsyEJ
>  >
>  > In my case I added the loop during the client
>  > usage of the ADT:
>  >
>  > From: Mild Shock <janburse@fastmail.fm>
>  > Subject: Source of the benchmark for DmitryVyukov
>  > Date: Tue, 21 Jul 2026 01:44:21 +0200
>  >
>  >      private static void producer(Queue q) {
>  >          for (int i = 0; i < WORK; i++) {
>  >              Integer val = Integer.valueOf(i);
>  >              while (!enqueue(q, val)) ; /** Looping **/
>  >          }
>  >      }
>  >
>  > Do you see the two loops, in your C code
>  > and in my Java code? They are marked with a
>  > comment /** Looping **/ .
>  >
>  > You see them, don't you? But I don't know
>  > exactly what backoff() does. Sometimes loops
>  > are spurious yield loops, required because
>  >
>  > an ADT cannot gurantee that every yield
>  > implies a certain condition. This is for
>  > example already found in the intrinsinc
>  >
>  > monitor of Java, the wait(). You might consult
>  > Doug Lea about the matter and how idiomatic
>  > Java code looks like dealing with
>  >
>  > spurious yields.
>  >
>  > Bye
> 
> Mild Shock schrieb:
>> Hi,
>>
>> Remember when first all local AI was Python
>> and PyTorch APIs. And then suddently people strated
>> using bare metal C/C++ Code. Here is the story:
>>
>> How it started:
>>
>> GPT-J or GPT-J-6B is an open-source large
>> language model (LLM) developed by EleutherAI
>> in 2021. As the name suggests, it is a
>> generative pre-trained transformer model
>> designed to produce human-like text that
>> continues from a prompt.
>> https://www.eleuther.ai/
>>
>> How it was going [Georgi Gerganov]:
>>
>> So a few days later comes out the LLaMA, I do
>> some calculations and I figure out “Okay, 65
>> billion parameters. You probably need about
>> 40 gigs of RAM, with 4-bit quantization. So
>> this can run on a MacBook. Why not do it?”
>>
>> Why I was able to do it so quickly - basically,
>> for all that I saw it’s pretty much GPT-J architecture
>> with some modifications, like some extra memorization
>> layers. It’s minor changes. Basically, again, the
>> existing code for the GPT-J, I just simply
>> modified it there, it happened pretty quickly.
>> https://changelog.com/podcast/532
>>
>> Georgi Gerganov, Bulgarian, now with Hugging
>> Face, ggml-cann also running on Chinese AI chips.
>> ggml Manifesto https://github.com/ggml-org/ggml
>>
>> Bye
>>
> 

Back to sci.physics.relativity | Previous | Next — Previous in thread | Next in thread | Find similar | Unroll thread

Thread

The Wuhan Virus that destroyed Python [ggml Manifesto] Mild Shock <janburse@fastmail.fm> - 2026-07-22 21:00 +0200
  Deadlock Exorcism: Switch from Push to Pull [A pi-calculus Specification of Prolog] (Re: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 00:23 +0200
    Why do you even need a mpmc queue? [Thunder Kittens] (Re: Deadlock Exorcism: Switch from Push to Pull) Mild Shock <janburse@fastmail.fm> - 2026-07-23 08:43 +0200
      Trivial balancing example for (int i=0; i<global_id; i++) (Re: Why do you even need a mpmc queue? [Thunder Kittens]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 08:57 +0200
        Enqueue/dequeue need not be fast and can spinn ["fairness" questions] (Was: Trivial balancing example for (int i=0; i<global_id; i++)) Mild Shock <janburse@fastmail.fm> - 2026-07-23 09:11 +0200
          The Pixel Phone AI Experiment Song (Re: Enqueue/dequeue need not be fast and can spinn ["fairness" questions] ) Mild Shock <janburse@fastmail.fm> - 2026-07-23 09:21 +0200
      Re: Why do you even need a mpmc queue? [Thunder Kittens] (Re: Deadlock Exorcism: Switch from Push to Pull) Ross Finlayson <ross.a.finlayson@gmail.com> - 2026-07-23 08:24 -0700
  Potential Python Recovery: Free Threading [3.13 release] (Was: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 10:19 +0200
    Re: Potential Python Recovery: Free Threading [3.13 release] (Was: The Wuhan Virus that destroyed Python [ggml Manifesto]) Ross Valikhanov <kavna@rl.ru> - 2026-07-23 16:01 +0000
  Re: The Wuhan Virus that destroyed Python [ggml Manifesto] Ramon Dubenkov <omd@nnk.ru> - 2026-07-23 13:38 +0000
  The things XILINX braught to the AMD table (Was: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 18:47 +0200
    NIVIDIA evacuated its Chinese market [Tau Scaling] (Was: The things XILINX braught to the AMD table) Mild Shock <janburse@fastmail.fm> - 2026-07-23 19:11 +0200
    NVIDIA evacuated its Chinese market [Tau Scaling] (Re: The things XILINX braught to the AMD table) Mild Shock <janburse@fastmail.fm> - 2026-07-23 19:12 +0200
      Re: NVIDIA evacuated its Chinese market [Tau Scaling] (Re: The things XILINX braught to the AMD table) Lane W <cactus_DAC@yahoo.com> - 2026-07-23 11:22 -0600

csiph-web