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

Groups > sci.logic > #348215

Trivial balancing example for (int i=0; i<global_id; i++) (Re: Why do you even need a mpmc queue? [Thunder Kittens])

Cross-posted to 3 groups.

Show all headers | View raw

Hi,

Is a trivial control construct for(),
when used in a compute shader with
NVIDIA Volta evolution, i.e. MIMD,

can lead to different time spend by
individual compute shaders:

fn main(global_id : i32) {
    i : i32 = 0;
    while (i < globa_id) {
       i++;
    }
}

You can visiualize as the time spent
by each logical thread as follows:

global id, logical thread life line
1     [    ]
2     [        ]
3     [            ]
4     [                ]
5     [                    ]
Etc..

With work items and load balancing you
could run the above with a lower number
of logical threads, I am writing the

work item number now inside the sub life
line inside the overall life line of
the logical thread:

worker , worker work items
A     [3           ]
B     [4               ][2       ]
C     [5                   ][1   ]

The overall time slightly increased by 1,
i.e. the case global_id = k combined
with the case global_id = n-k+1 . Also

one worker didn't have two work items,
only one work item. But the number of
logical threads needed was halfed.

Ok, a mpmc queue will be not that
intelligent, concerning the work sheduling.
But one could experiment with mpmc queue

priority queues etc.. etc..

Have Fun!

Bye

Mild Shock schrieb:
> 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