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

Groups > sci.physics.relativity > #671598

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 2 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,
>>  >
>>  > 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