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

Groups > sci.physics.relativity > #671599

Enqueue/dequeue need not be fast and can spinn ["fairness" questions] (Was: Trivial balancing example for (int i=0; i<global_id; i++))

Cross-posted to 2 groups.

Show all headers | View raw

Hi,

Since the main goal is to change the
geometry of how work items are assigned
and processed, and since work item processing

costs much more time then queue API calls.
It is irrelevant whether the queue API are
ultra fast or not, or whether they spin or

not. Usually one can affort both things,
since balancing outweights these small extra
costs or silly occupation of a platform thread.

It could be also feasible, in the total
balance sheet of execution and energy consumption
of your work items along workers to use

Atomic.pause() equivalents, special instructions,
for spinning. This will be added in 2027 to JavaScript:

For example, in Intel x86, it may be a
pause instruction as per Intel's optimization manual.
https://developer.mozilla.org/en-US/docs/Web/JavaScript/Reference/Global_Objects/Atomics/pause

But I have to see what works and what doesn't
work. The above are only theoretical
considerations about the friction that queues

induce. Usually friction hurts much less
than balant scheduling mistakes, i.e. unfavorable
schedulings. golang is very much bugged by

unfavorable schedulings. While on paper golang
looks like brilliant idea, examples that create
similar patterns of runtimes as in the below

for (int i=0; i<global_id; i++) can create
horrible scheduling and very bad an erratic
behaviour, especially if the synchronization

mechanism and backpressure signals have to travel
long chains of channel waits. So you will get not
into the terrain of "dead lock" questions, but

basically enter the terrain of "fairness" questions.

Bye

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