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

Groups > sci.logic > #348214

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

Cross-posted to 3 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,
>>
>> Remember when first all local AI was Python
>> and PyTorch APIs. And then suddently people started
>> 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.logic | 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:01 +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:25 +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:45 +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:55 +0200
        The Pixel Phone AI Experiment Song (Enqueue/dequeue need not be fast and can spinn ["fairness" questions]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 09:19 +0200
          Enqueue/dequeue need not be fast and can spinn ["fairness" questions] (Re: The Pixel Phone AI Experiment Song (Enqueue/dequeue need not be fast and can spinn ["fairness" questions]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 09:23 +0200
      And, where did I talk about rockets? [Hint its about xAI's Grok] (Re: Why do you even need a mpmc queue? [Thunder Kittens]) Mild Shock <janburse@fastmail.fm> - 2026-07-25 01:25 +0200
        Why forget something, that was never on my mind (Re: And, where did I talk about rockets? [Hint its about xAI's Grok]) Mild Shock <janburse@fastmail.fm> - 2026-07-25 09:49 +0200
      Example Mandel Brot rendering [Faster with MIMD] (Was: Why do you even need a mpmc queue? [Thunder Kittens]) Mild Shock <janburse@fastmail.fm> - 2026-07-25 09:56 +0200
  Potential Python Recovery: Free Threading [3.13 release] (Re: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 10:21 +0200
  The things XILINX braught to the AMD table (Re: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-23 18:48 +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:13 +0200
      Micro penis mother sung arias (Re: NVIDIA evacuated its Chinese market [Tau Scaling]) Mild Shock <janburse@fastmail.fm> - 2026-07-24 14:40 +0200
        Micro penis brain is in constant hiatus (Re: Micro penis mother sung arias) Mild Shock <janburse@fastmail.fm> - 2026-07-24 15:27 +0200
          Ignoramus or Ignorabimus: I don't care [(Re: Micro penis brain is in constant hiatus (Re: Micro penis mother sung arias) Mild Shock <janburse@fastmail.fm> - 2026-07-24 15:35 +0200
            You are a moron, brainless putin payed (Re: Ignoramus or Ignorabimus: I don't care) Mild Shock <janburse@fastmail.fm> - 2026-07-24 18:01 +0200
              Yeah keep reading my posts, uninspired fool (Re: You are a moron, brainless putin payed) Mild Shock <janburse@fastmail.fm> - 2026-07-24 19:47 +0200
            Out of the blue accusation span 15 days [Empirical USENET study] (Re: Ignoramus or Ignorabimus: I don't care) Mild Shock <janburse@fastmail.fm> - 2026-07-24 20:27 +0200
  Little Data Center on Your Palm [AI Laptops for 500 USD] (Re: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-24 17:59 +0200
    2008: 4 Blades + Tesla S1070 versus 2026: 1 AI Laptop (Re: Little Data Center on Your Palm [AI Laptops for 500 USD]) Mild Shock <janburse@fastmail.fm> - 2026-07-24 18:15 +0200
  Hurry the blue bus doesnt stop indefinitely (Re: The Wuhan Virus that destroyed Python [ggml Manifesto]) Mild Shock <janburse@fastmail.fm> - 2026-07-24 20:37 +0200
    Not SIMD, a MIMD design for NVIDIA Volta (Re: Hurry the blue bus doesnt stop indefinitely) Mild Shock <janburse@fastmail.fm> - 2026-07-24 20:58 +0200
      Could take 3-4 months find machine / browser (Re: Not SIMD, a MIMD design for NVIDIA Volta) Mild Shock <janburse@fastmail.fm> - 2026-07-24 21:16 +0200

csiph-web