Groups | Search | Server Info | Keyboard shortcuts | Login | Register [http] [https] [nntp] [nntps]
| From | Mild Shock <janburse@fastmail.fm> |
|---|---|
| Newsgroups | sci.math, sci.physics.relativity |
| Subject | The Pixel Phone AI Experiment Song (Re: Enqueue/dequeue need not be fast and can spinn ["fairness" questions] ) |
| Date | 2026-07-23 09:21 +0200 |
| Message-ID | <113sfdt$8o56$2@solani.org> (permalink) |
| References | <113r40o$8g50$1@solani.org> <113rfso$84i6$1@solani.org> <113sd5n$8mfl$1@solani.org> <113se1d$8mvd$2@solani.org> <113ser4$8npg$1@solani.org> |
Cross-posted to 2 groups.
Hi,
For graphics, rendering, in the worst case your
FPS can go down. Then you have a 2x as performant
graphic card, and sundently the FPS is ok again!
Or you render into a smaller screen, with less
number of pixels, and things turn good again.
So for my pixel phone AI experiment, or what I
will toy around on a AI laptop, I am singing:
I'm a spinner, I'm a sinner
I spin on CAS loops for my dinner
Some call it busy-wait, I call it fate
When the queue is empty, I just rotate
Bye
Whats better Steve Miller or Muddy Waters?
The Joker
https://www.youtube.com/watch?v=dV3AziKTBUo
Hoochie Coochie Man
https://www.youtube.com/watch?v=e_l6A7krjrQ
Mild Shock schrieb:
> 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.math | Previous | Next — Previous in thread | Next in thread | Find similar | Unroll 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