2058420153348391346

2026-08-29 10:55:28 -04:00
1 parent 412b288083
commit 4f949c20e4
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim Lots and lots of atomics."
+type: archive
+source: twitter
+source_url: "https://x.com/2Sexy4MyGPU/status/2058502561897030002"
+author: "Matthew Rusch"
+handle: 2Sexy4MyGPU
+post_id: "2058502561897030002"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Lots and lots of atomics."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/2Sexy4MyGPU/status/2058502561897030002
+- Author: Matthew Rusch (@2Sexy4MyGPU)
+- Posted: 2026-05-24 10:57:10
+
+## Branch
+
+**1/** **@2Sexy4MyGPU** ^2058502561897030002
+
+**@AgileJebrim**
+
+Lots and lots of atomics.
+
+**2/** **@AgileJebrim** ^2058541287825924369
+
+**@2Sexy4MyGPU**
+
+-
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim Shrink the array over and over until it's down to a single value."
+type: archive
+source: twitter
+source_url: "https://x.com/Apolloknius/status/2058534260387283164"
+author: "ApolloniusofTyana"
+handle: Apolloknius
+post_id: "2058534260387283164"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Shrink the array over and over until it's down to a single value."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Apolloknius/status/2058534260387283164
+- Author: ApolloniusofTyana (@Apolloknius)
+- Posted: 2026-05-24 13:03:08
+
+## Branch
+
+**1/** **@Apolloknius** ^2058534260387283164
+
+**@AgileJebrim**
+
+Shrink the array over and over until it's down to a single value.  We used to do this with textures to sample the overall luminance for tone mapping, as an example.
+
+**2/** **@AgileJebrim** ^2058539821576937494
+
+**@Apolloknius**
+
+Need details in relation to the hardware. Assume compute shaders are used, not multiple draw calls of fragment shaders. Those aren’t actually that efficient.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,45 @@
+---
+title: "@AgileJebrim Add up all the columns separately."
+type: archive
+source: twitter
+source_url: "https://x.com/ArcticDodger/status/2058657119247925435"
+author: "Lodger 🚙🏡🌤️🔋"
+handle: ArcticDodger
+post_id: "2058657119247925435"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Add up all the columns separately."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/ArcticDodger/status/2058657119247925435
+- Author: Lodger 🚙🏡🌤️🔋 (@ArcticDodger)
+- Posted: 2026-05-24 21:11:19
+
+## Branch
+
+**1/** **@ArcticDodger** ^2058657119247925435
+
+**@AgileJebrim**
+
+Add up all the columns separately. Give each additional task to a different GPU. Sum the sums. Is this supposed to be some kind of trick question?
+
+**2/** **@AgileJebrim** ^2058658618166100440
+
+**@ArcticDodger**
+
+What columns? It’s a 1D array.
+
+Different GPUs? That’s overkill. You simply won’t have enough data or bitwidth to sum at that point. A single GPU could even handle summing 4b 1-bit values.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,41 @@
+---
+title: "@AgileJebrim Large array of values is stored in float/double."
+type: archive
+source: twitter
+source_url: "https://x.com/CodeAndTears/status/2058560744694693968"
+author: "Jules"
+handle: CodeAndTears
+post_id: "2058560744694693968"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Large array of values is stored in float/double."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/CodeAndTears/status/2058560744694693968
+- Author: Jules (@CodeAndTears)
+- Posted: 2026-05-24 14:48:22
+
+## Branch
+
+**1/** **@CodeAndTears** ^2058560744694693968
+
+**@AgileJebrim**
+
+Large array of values is stored in float/double.
+
+Result must be bit exact reproductible between runs/machines/gpu.
+
+Have fun ..
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,55 @@
+---
+title: "@AgileJebrim They asked the same question to me at nvidia, it's a good one!"
+type: archive
+source: twitter
+source_url: "https://x.com/FUZxxl/status/2058471646474760280"
+author: "Robert Clausecker"
+handle: FUZxxl
+post_id: "2058471646474760280"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim They asked the same question to me at nvidia, it's a good one!"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/FUZxxl/status/2058471646474760280
+- Author: Robert Clausecker (@FUZxxl)
+- Posted: 2026-05-24 08:54:19
+
+## Branch
+
+**1/** **@FUZxxl** ^2058471646474760280
+
+**@AgileJebrim**
+
+They asked the same question to me at nvidia, it's a good one!
+
+**2/** **@AgileJebrim** ^2058530131191161252
+
+**@FUZxxl**
+
+I wish they’d done that. Unfortunately I got all the bad engineering teams doing interviews lol. I wasn’t particularly impressed.
+
+**3/** **@FUZxxl** ^2058531664003776993
+
+**@AgileJebrim**
+
+I kind of stumbled through it as I was interviewing for a CPU position and don't have much GPU experience, plus most work I've done so far was integer work and the interviewer asked lots of questions on catastrophic cancellation.
+
+**4/** **@AgileJebrim** ^2058540542229647393
+
+**@FUZxxl**
+
+We go all in on integers. Floats have too many problems and are best served as just a sit in for low bitwidth integers.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,39 @@
+---
+title: "@AgileJebrim I think this is a simple example of \"look how much I know about a niche topic, let's find something I know but others don't, so I feel superior\"."
+type: archive
+source: twitter
+source_url: "https://x.com/Gombally_/status/2058450878114271314"
+author: "Gombally"
+handle: Gombally_
+post_id: "2058450878114271314"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I think this is a simple example of \"look how much I know about a niche topic, let's find something I know but others don't, so I feel superior\"."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Gombally_/status/2058450878114271314
+- Author: Gombally (@Gombally_)
+- Posted: 2026-05-24 07:31:48
+
+## Branch
+
+**1/** **@Gombally_** ^2058450878114271314
+
+**@AgileJebrim**
+
+I think this is a simple example of "look how much I know about a niche topic, let's find something I know but others don't, so I feel superior".
+
+It's a retarded ass question, unless that job solely relies around utilising GPU's.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim M x N multipass reduction"
+type: archive
+source: twitter
+source_url: "https://x.com/Jason_Ewton/status/2058487515728785437"
+author: "Jason Ewton"
+handle: Jason_Ewton
+post_id: "2058487515728785437"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim M x N multipass reduction"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Jason_Ewton/status/2058487515728785437
+- Author: Jason Ewton (@Jason_Ewton)
+- Posted: 2026-05-24 09:57:23
+
+## Branch
+
+**1/** **@Jason_Ewton** ^2058487515728785437
+
+**@AgileJebrim**
+
+M x N multipass reduction
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,47 @@
+---
+title: "@AgileJebrim Then make them build it in Minecraft using pre-2016 redstone."
+type: archive
+source: twitter
+source_url: "https://x.com/Lord_KX_/status/2058578623360274916"
+author: "LordKX"
+handle: Lord_KX_
+post_id: "2058578623360274916"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Then make them build it in Minecraft using pre-2016 redstone."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Lord_KX_/status/2058578623360274916
+- Author: LordKX (@Lord_KX_)
+- Posted: 2026-05-24 15:59:25
+
+## Branch
+
+**1/** **@Lord_KX_** ^2058578623360274916
+
+**@AgileJebrim**
+
+Then make them build it in Minecraft using pre-2016 redstone.
+
+**2/** **@AgileJebrim** ^2058580585023484020
+
+**@Lord_KX_**
+
+👏
+
+**3/** **@AgileJebrim** ^2058581682303668551
+
+Unfortunately Minecraft doesn’t do parallelism very well. We’re building a way to program GPUs in a visual way that reflects the hardware though such that one can program it not too dissimilarly from Minecraft redstone. Text doesn’t offer the same insight into performance of different operations and data layouts as a visual sim would.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,83 @@
+---
+title: "@AgileJebrim many people don't write code that touches the gpu"
+type: archive
+source: twitter
+source_url: "https://x.com/Lotsoffoodinme/status/2058454499828904170"
+author: "el jack"
+handle: Lotsoffoodinme
+post_id: "2058454499828904170"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim many people don't write code that touches the gpu"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Lotsoffoodinme/status/2058454499828904170
+- Author: el jack (@Lotsoffoodinme)
+- Posted: 2026-05-24 07:46:11
+
+## Branch
+
+**1/** **@Lotsoffoodinme** ^2058454499828904170
+
+**@AgileJebrim**
+
+many people don't write code that touches the gpu
+
+**2/** **@AgileJebrim** ^2058531033809580044
+
+**@Lotsoffoodinme**
+
+We’re changing that.
+
+**3/** **@Lotsoffoodinme** ^2058534577136967750
+
+**@AgileJebrim**
+
+Out of interest, is it something drastically different to dividing it up into many slices and summing from there? I’m envisioning something like a map reduce
+
+**4/** **@AgileJebrim** ^2058539551895834686
+
+**@Lotsoffoodinme**
+
+It’s in that direction but I need greater details.
+
+**5/** **@Lotsoffoodinme** ^2058554249488064571
+
+**@AgileJebrim**
+
+As I don’t write code that targets the gpu I can’t. Is this a particularly complicated question? It seems to a complete lay person that it shouldn’t be too complicated.
+
+**6/** **@AgileJebrim** ^2058554931980705979
+
+**@Lotsoffoodinme**
+
+How would you answer this for a multicore CPU?
+
+**7/** **@Lotsoffoodinme** ^2058559658827382904
+
+Im assuming it’s a big enough set of numbers that chunking it makes sense. Maybe it’s hand wavy, but I’m also going to assume it fits in a 64bit number.
+
+Gameplan would be to chunk by core count, then have each core iterate over its slice. I rarely explicitly write vectorised code, but I’m reasonably confident I’d want more than one accumulator. By that I mean each iteration of the loop might do 4 additions at once and then increment by 32 or whatever the cpu supports.
+
+At the end reduce the accumulators and then reduce the threads.
+
+I hope that’s not a terrible answer.
+
+**8/** **@AgileJebrim** ^2058560525810692312
+
+**@Lotsoffoodinme**
+
+You can scale a lot more than just 4-wide SIMD.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim If the array is in RAM, do it sequentially."
+type: archive
+source: twitter
+source_url: "https://x.com/MichalPiotrK/status/2058669746799673769"
+author: "Michał"
+handle: MichalPiotrK
+post_id: "2058669746799673769"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim If the array is in RAM, do it sequentially."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/MichalPiotrK/status/2058669746799673769
+- Author: Michał (@MichalPiotrK)
+- Posted: 2026-05-24 22:01:30
+
+## Branch
+
+**1/** **@MichalPiotrK** ^2058669746799673769
+
+**@AgileJebrim**
+
+If the array is in RAM, do it sequentially. If it's in GPU memory, use partial sums sqrt(n) elements, then in a second run you get the result. Unless there are constraints, e.g. floating point with error tracking or the array is really very large, or do I optimize for power or...
+
+**2/** **@AgileJebrim** ^2058676643112136824
+
+**@MichalPiotrK**
+
+Nope.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim This is an idiotic interview question"
+type: archive
+source: twitter
+source_url: "https://x.com/MikeTreit/status/2058487016317399442"
+author: "Mike Treit"
+handle: MikeTreit
+post_id: "2058487016317399442"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim This is an idiotic interview question"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/MikeTreit/status/2058487016317399442
+- Author: Mike Treit (@MikeTreit)
+- Posted: 2026-05-24 09:55:24
+
+## Branch
+
+**1/** **@MikeTreit** ^2058487016317399442
+
+**@AgileJebrim**
+
+This is an idiotic interview question
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,51 @@
+---
+title: "@AgileJebrim I have no meaningful GPU coding experience."
+type: archive
+source: twitter
+source_url: "https://x.com/SpikyCaterpillr/status/2058630363509907942"
+author: "Spikier Caterpillar"
+handle: SpikyCaterpillr
+post_id: "2058630363509907942"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I have no meaningful GPU coding experience."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/SpikyCaterpillr/status/2058630363509907942
+- Author: Spikier Caterpillar (@SpikyCaterpillr)
+- Posted: 2026-05-24 19:25:00
+
+## Branch
+
+**1/** **@SpikyCaterpillr** ^2058630363509907942
+
+**@AgileJebrim**
+
+I have no meaningful GPU coding experience. That said, if it's a LARGE array in GPU memory, the obvious guess would be "tell several compute units to sum chunks of it in parallel, then add the outputs together". The optimum chunk size will vary by GPU and possibly other factors.
+
+**2/** **@AgileJebrim** ^2058630786803229053
+
+**@SpikyCaterpillr**
+
+"tell several compute units to sum chunks of it in parallel, then add the outputs together"
+
+Any idea how you’d structure it? How would you do this on multicore CPUs using SIMD?
+
+**3/** **@SpikyCaterpillr** ^2058640615412224473
+
+**@AgileJebrim**
+
+First determine the desired chunk size (I assume compute units have limited addressable space, for one), then iterate over the large array calling starting a summer on each chunk. When I run out of cores or GPU memory, wait and sum what I've got before continuing down the array.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim I would embarrass myself by trying to answer because from the application level a GPU is just a way of using money to make pytorch faster."
+type: archive
+source: twitter
+source_url: "https://x.com/VMises76153/status/2058464896971382831"
+author: "Nathaniel Nifong"
+handle: VMises76153
+post_id: "2058464896971382831"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I would embarrass myself by trying to answer because from the application level a GPU is just a way of using money to make pytorch faster."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/VMises76153/status/2058464896971382831
+- Author: Nathaniel Nifong (@VMises76153)
+- Posted: 2026-05-24 08:27:30
+
+## Branch
+
+**1/** **@VMises76153** ^2058464896971382831
+
+**@AgileJebrim**
+
+I would embarrass myself by trying to answer because from the application level a GPU is just a way of using money to make pytorch faster.
+
+**2/** **@AgileJebrim** ^2058530860370903075
+
+**@VMises76153**
+
+We are building a platform for creating general purpose hard real-time applications that natively run directly on the GPU. No PyTorch.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,67 @@
+---
+title: "@AgileJebrim Hierarchical map reduce?"
+type: archive
+source: twitter
+source_url: "https://x.com/atorstling/status/2058449690572501039"
+author: "Alexander Torstling"
+handle: atorstling
+post_id: "2058449690572501039"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Hierarchical map reduce?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/atorstling/status/2058449690572501039
+- Author: Alexander Torstling (@atorstling)
+- Posted: 2026-05-24 07:27:05
+
+## Branch
+
+**1/** **@atorstling** ^2058449690572501039
+
+**@AgileJebrim**
+
+Hierarchical map reduce?
+
+**2/** **@AgileJebrim** ^2058531652595253265
+
+**@atorstling**
+
+Okay. Now explain how that works.
+
+**3/** **@atorstling** ^2058614306871067052
+
+**@AgileJebrim**
+
+Use shared memory for in-batch coordination, and global memory for intra-batch coordination. Each batch takes a set of numbers and produce a sum. Next round of batches operates on the output of the previous round. If fancy, you could encode the problem in kernels/registers.
+
+**4/** **@atorstling** ^2058614911580701084
+
+**@AgileJebrim**
+
+Since you can use output of batches from different rounds as input, you can keep the pipeline full by starting new rounds eagerly
+
+**5/** **@AgileJebrim** ^2058616473262641450
+
+**@atorstling**
+
+I’m not really satisfied with this answer.
+
+**6/** **@atorstling** ^2058638219365818830
+
+**@AgileJebrim**
+
+You don't need to be, all good :)
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,88 @@
+---
+title: "I'm bored so I'm gonna give this question a good shot (I learned a bunch of GPU architecture stuff back in uni but I had to look a couple things up to make sure I wasn't mixing terms up)."
+type: archive
+source: twitter
+source_url: "https://x.com/brainage19/status/2058448551290839084"
+author: "thomas 🍙"
+handle: brainage19
+post_id: "2058448551290839084"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "I'm bored so I'm gonna give this question a good shot (I learned a bunch of GPU architecture stuff back in uni but I had to look a couple things up to make sure I wasn't mixing terms up)."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/brainage19/status/2058448551290839084
+- Author: thomas 🍙 (@brainage19)
+- Posted: 2026-05-24 07:22:33
+
+## Branch
+
+**1/** **@brainage19** ^2058448551290839084
+
+I'm bored so I'm gonna give this question a good shot (I learned a bunch of GPU architecture stuff back in uni but I had to look a couple things up to make sure I wasn't mixing terms up). I would assume the most efficient way is to mass add them in pairs across all streaming multiprocessors to halve the numbers you need to add, and keep doing that until you have 1 number left (the final result).
+
+You take M chunks of N numbers from global memory and load then into L2 cache, L1 caches in streaming multiprocessors, shared memory, then finally into registers, instructions to add them in pairs get scheduled to run on threads via the warp scheduler, and the results get written into registers. I'm not exactly sure how the instructions get loaded into instruction cache but I imagine it would be a similar process.
+
+Load the results from the registers all the way back up into global memory, now you have N/2 result numbers, N/2 original numbers, and the rest of the numbers. You repeat the process for the next chunk of numbers, making sure the results are read back into VRAM where the original numbers from the previous iteration start (this is to make sure after all chunks are processed, the results are contiguous in memory - working with contiguous memory is more efficient due to how memory is read/written in cache lines of several bytes at a time to minimise the cycles wasted waiting for data to be loaded into memory)
+
+You keep repeating this process until eventually you will have processed all of the chunks and you will have half of the numbers in memory that you started with in global memory contiguously. Before repeating this process of processing M chunks of N numbers, you would have to wait until all streaming multiprocessors have finished processing all chunks of data. This probably sounds really stupid given how much other stuff I know and the level of detail that I have explained it but I'm not sure how that really works at an architecture/instruction level.
+
+You repeat this process of processing M chunks of N numbers, repeatedly halving the numbers in VRAM until eventually you have 1 number left.
+
+Is this the right idea? Is there a more efficient way? Have I explained all the concepts correctly?
+
+**2/** **@AgileJebrim** ^2058532143576334428
+
+**@brainage19**
+
+Not really sure why you’re halving anything.
+
+**3/** **@shrydar** ^2058583721079722406
+
+**@AgileJebrim** **@brainage19**
+
+please tell me you know you can’t in general just use an accumulator (or small set of accumulators) to collect a partial sum from streaming a significant subset of the data..
+
+**4/** **@AgileJebrim** ^2058584367790788727
+
+**@shrydar** **@brainage19**
+
+You can use it for 99% of the time and only do the tree reduction at the very end.
+
+**5/** **@shrydar** ^2058705598947201535
+
+**@AgileJebrim** **@brainage19**
+
+Even chunking into groups of 1000 you’re throwing away nine or ten bits of accuracy. Depending on the application that could be fine or it could be fatal.
+
+**6/** **@AgileJebrim** ^2058707817314271710
+
+**@shrydar** **@brainage19**
+
+Why are you assuming floats?
+https://x.com/agilejebrim/status/2058562643158352101?s=46&t=ZbNecpPBsBMrvteeaA6IcA
+
+**7/** **@shrydar** ^2058709232938361308
+
+**@AgileJebrim** **@brainage19**
+
+I wasn’t assuming either way, but it’d be why Thomas was repeatedly halving the length of the list, so I was weirded out that you weren’t sure why he was even considering that. Clearly not needed for a 32 bit exact sum :)
+
+**8/** **@AgileJebrim** ^2058711742843166954
+
+**@shrydar** **@brainage19**
+
+The only thing I can infer from him adding pairs and repeatedly halving is if he was attempting to perform a reduction every single iteration. This is needlessly expensive.  Cross-communication between lanes, warps, and SMs needs to be kept to a minimum and only occur at the end.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim If you can fit Qwen 3.6 into your gpu, just load that with ollama and it will write the code for you."
+type: archive
+source: twitter
+source_url: "https://x.com/braincell__/status/2058530959029604567"
+author: "John Mell 🇺🇦🇦🇺"
+handle: braincell__
+post_id: "2058530959029604567"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim If you can fit Qwen 3.6 into your gpu, just load that with ollama and it will write the code for you."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/braincell__/status/2058530959029604567
+- Author: John Mell 🇺🇦🇦🇺 (@braincell__)
+- Posted: 2026-05-24 12:50:00
+
+## Branch
+
+**1/** **@braincell__** ^2058530959029604567
+
+**@AgileJebrim**
+
+If you can fit Qwen 3.6 into your gpu, just load that with ollama and it will write the code for you.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim It's for a frontend JavaScript position I see."
+type: archive
+source: twitter
+source_url: "https://x.com/dadbod2100/status/2058471050909020439"
+author: "dadbod2100"
+handle: dadbod2100
+post_id: "2058471050909020439"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim It's for a frontend JavaScript position I see."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/dadbod2100/status/2058471050909020439
+- Author: dadbod2100 (@dadbod2100)
+- Posted: 2026-05-24 08:51:57
+
+## Branch
+
+**1/** **@dadbod2100** ^2058471050909020439
+
+**@AgileJebrim**
+
+It's for a frontend JavaScript position I see.
+
+**2/** **@AgileJebrim** ^2058530234958238074
+
+**@dadbod2100**
+
+Definitely not.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim What kind of GPU?"
+type: archive
+source: twitter
+source_url: "https://x.com/deleted_1m/status/2058504770453754227"
+author: "Deleted"
+handle: deleted_1m
+post_id: "2058504770453754227"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim What kind of GPU?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/deleted_1m/status/2058504770453754227
+- Author: Deleted (@deleted_1m)
+- Posted: 2026-05-24 11:05:57
+
+## Branch
+
+**1/** **@deleted_1m** ^2058504770453754227
+
+**@AgileJebrim**
+
+What kind of GPU?
+
+**2/** **@AgileJebrim** ^2058541240056979920
+
+**@deleted_1m**
+
+Any. This must be portable compute shader code.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim No idea about GPU programming, cpu hints towards divide by half, but wouldn’t matrix multiplication works better here? Something like ask gpu to multiply source array by column full of 1 of the same length? Would love if somebody can explain if that would work"
+type: archive
+source: twitter
+source_url: "https://x.com/fatroom/status/2058498000666698206"
+author: "Roman Romanchuk"
+handle: fatroom
+post_id: "2058498000666698206"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim No idea about GPU programming, cpu hints towards divide by half, but wouldn’t matrix multiplication works better here? Something like ask gpu to multiply source array by column full of 1 of the same length? Would love if somebody can explain if that would work"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/fatroom/status/2058498000666698206
+- Author: Roman Romanchuk (@fatroom)
+- Posted: 2026-05-24 10:39:03
+
+## Branch
+
+**1/** **@fatroom** ^2058498000666698206
+
+**@AgileJebrim**
+
+No idea about GPU programming, cpu hints towards divide by half, but wouldn’t matrix multiplication works better here? Something like ask gpu to multiply source array by column full of 1 of the same length? Would love if somebody can explain if that would work
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim divide into small arrays, conquer each array on a core, then sum the sums?"
+type: archive
+source: twitter
+source_url: "https://x.com/feikcel/status/2058484875477303535"
+author: "гейоргий 🍏✌️| набирает массу | 170 см | фэтфоб"
+handle: feikcel
+post_id: "2058484875477303535"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim divide into small arrays, conquer each array on a core, then sum the sums?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/feikcel/status/2058484875477303535
+- Author: гейоргий 🍏✌️| набирает массу | 170 см | фэтфоб (@feikcel)
+- Posted: 2026-05-24 09:46:53
+
+## Branch
+
+**1/** **@feikcel** ^2058484875477303535
+
+**@AgileJebrim**
+
+divide into small arrays, conquer each array on a core, then sum the sums?
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim I am not knowlegable enough with GPUs, but does numerical stability of the result play a role in questions like these?"
+type: archive
+source: twitter
+source_url: "https://x.com/gulbaru/status/2058477347834060854"
+author: "Gulbaru"
+handle: gulbaru
+post_id: "2058477347834060854"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I am not knowlegable enough with GPUs, but does numerical stability of the result play a role in questions like these?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/gulbaru/status/2058477347834060854
+- Author: Gulbaru (@gulbaru)
+- Posted: 2026-05-24 09:16:59
+
+## Branch
+
+**1/** **@gulbaru** ^2058477347834060854
+
+**@AgileJebrim**
+
+I am not knowlegable enough with GPUs, but does numerical stability of the result play a role in questions like these?
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim Do floor(warpsize/arraysize) sums, keep using the same accumulated warpsize array until the whole array is done."
+type: archive
+source: twitter
+source_url: "https://x.com/heartoneto2012/status/2058466240033677747"
+author: "heartoneto"
+handle: heartoneto2012
+post_id: "2058466240033677747"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Do floor(warpsize/arraysize) sums, keep using the same accumulated warpsize array until the whole array is done."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/heartoneto2012/status/2058466240033677747
+- Author: heartoneto (@heartoneto2012)
+- Posted: 2026-05-24 08:32:50
+
+## Branch
+
+**1/** **@heartoneto2012** ^2058466240033677747
+
+**@AgileJebrim**
+
+Do floor(warpsize/arraysize) sums, keep using the same accumulated warpsize array until the whole array is done. Finally accumulate the partial sums
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,82 @@
+---
+title: "@AgileJebrim you shouldn’t do that on the GPU."
+type: archive
+source: twitter
+source_url: "https://x.com/jes5199/status/2058590784740540610"
+author: "Jes Wolfe!"
+handle: jes5199
+post_id: "2058590784740540610"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim you shouldn’t do that on the GPU."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/jes5199/status/2058590784740540610
+- Author: Jes Wolfe! (@jes5199)
+- Posted: 2026-05-24 16:47:44
+
+## Branch
+
+**1/** **@jes5199** ^2058590784740540610
+
+**@AgileJebrim**
+
+you shouldn’t do that on the GPU. do it on the network card
+
+**2/** **@AgileJebrim** ^2058591291059892482
+
+**@jes5199**
+
+Network cards can’t compete with the processing power of GPUs and a parallel summation is only one small aspect of a much larger amount of work that would need to be done. It’s worth it to become GPU resident.
+
+**3/** **@jes5199** ^2058598927071662429
+
+**@AgileJebrim**
+
+where’s your data come from
+did you do a massively
+parallel
+download
+
+**4/** **@AgileJebrim** ^2058599326553670066
+
+**@jes5199**
+
+An SSBO in VRAM
+
+**5/** **@AgileJebrim** ^2058599922706899325
+
+**@jes5199**
+
+A rgba32ui texel buffer if you want to do better. :P
+
+**6/** **@AgileJebrim** ^2058600334197395905
+
+**@jes5199**
+
+The original data was generated in compute shaders per real-time tick.
+
+**7/** **@jes5199** ^2058607245194383821
+
+**@AgileJebrim**
+
+oh imaginary data
+
+**8/** **@AgileJebrim** ^2058609452081615079
+
+**@jes5199**
+
+No, just freshly generated data that changes every tick in response to small user inputs.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,97 @@
+---
+title: "@AgileJebrim Easy."
+type: archive
+source: twitter
+source_url: "https://x.com/kechogarcia/status/2058522912731123861"
+author: "Kecho"
+handle: kechogarcia
+post_id: "2058522912731123861"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Easy."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/kechogarcia/status/2058522912731123861
+- Author: Kecho (@kechogarcia)
+- Posted: 2026-05-24 12:18:02
+
+## Branch
+
+**1/** **@kechogarcia** ^2058522912731123861
+
+**@AgileJebrim**
+
+Easy. WaveActiveSum, then LDS atomic sum, then global memory sum between groups (if GPU has intrinsica) then finally dram atomics. 
+Always go in cache hierarchy -> wave, LDS, global mem, vram
+
+Never do an atomic per thread.
+
+Also you can do it all in one persistent dispatch .
+
+**2/** **@AgileJebrim** ^2058524431547306220
+
+**@kechogarcia**
+
+Say you have many more data elements than you have SIMD lanes across the hardware. How are you handling that?
+
+**3/** **@kechogarcia** ^2058535071947145503
+
+**@AgileJebrim**
+
+Do it all in a persistent dispatch. Dispatch and fill the machine, then consume input.
+
+**4/** **@AgileJebrim** ^2058539350262989272
+
+**@kechogarcia**
+
+What I’m looking for an answer to is what the for loop inside the persistent dispatch is doing at the smallest level. Are you looping over WaveActiveSums?
+
+**5/** **@kechogarcia** ^2058546836697149904
+
+**@AgileJebrim**
+
+No. If you pad the buffer to wave size, just accumulate on 4 registers, 4 loads at a time. Or whatever the hw can do best. When you run out of work, then the problem is a cache hierarchical reduction. You can do it all in a single dispatch, no fancy  multipass bin reduction.
+
+**6/** **@kechogarcia** ^2058548990413504538
+
+**@AgileJebrim**
+
+ldsV = 0;
+GroupBarrier();
+float4 a = 0;
+while(has work())
+     a += load4(popnext());
+
+float4 perGroup = WaveActiveSum(a);
+
+if (is first lane(() atomicAdd(ldsV, perGroup);
+
+if (threadID == 0)
+    atomicAdd(ldsV, dramBuffer[0];
+
+Group mem shenanigans I have to read manual.
+
+**7/** **@AgileJebrim** ^2058550742453989755
+
+Okay good. I was looking for the partial sums per lane answer. You had left that out entirely in your initial answer. Reductions in a loop would be expensive.
+
+If I told you there were 16m unsigned integer elements to sum up, what other requirement would necessarily need to exist (or otherwise be addressed) and what other optimization opportunities would arise as a result of that knowledge?
+
+**8/** **@kechogarcia** ^2058558465270755676
+
+**@AgileJebrim**
+
+Hmmm. Means you could easily overflow, so now we gotta use 64 bits for accumulation? But this means we could double the number of sums on each thread by loading 64 bit primitives? Not sure I gotta think harder if you are looking for something else
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,53 @@
+---
+title: "@AgileJebrim Divide up the available memory between cores, each with out output slot and as many input slots as possible."
+type: archive
+source: twitter
+source_url: "https://x.com/matt_timmermans/status/2058551271389356330"
+author: "Matt Timmermans"
+handle: matt_timmermans
+post_id: "2058551271389356330"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Divide up the available memory between cores, each with out output slot and as many input slots as possible."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/matt_timmermans/status/2058551271389356330
+- Author: Matt Timmermans (@matt_timmermans)
+- Posted: 2026-05-24 14:10:43
+
+## Branch
+
+**1/** **@matt_timmermans** ^2058551271389356330
+
+**@AgileJebrim**
+
+Divide up the available memory between cores, each with out output slot and as many input slots as possible.   Load memory, have each core accumulate its part, repeat.
+
+Might be able to go faster with double-buffering so you can load and sum at the same time.
+
+**2/** **@AgileJebrim** ^2058551683420913730
+
+**@matt_timmermans**
+
+“have each core accumulate its part”
+
+Detail this.
+
+**3/** **@matt_timmermans** ^2058554402332409962
+
+**@AgileJebrim**
+
+I'm not actually familiar enough to know why you asking for detail (I should NOT get your job :-), but each pass is essentially an image resize.  Separate input + output memory, organize input to reduce contention.  Only a few thousand output pixels, so that part not important.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim Does the array initally exist in GPU memory or CPU?"
+type: archive
+source: twitter
+source_url: "https://x.com/minusYCore/status/2058488491399356688"
+author: "Mongo"
+handle: minusYCore
+post_id: "2058488491399356688"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Does the array initally exist in GPU memory or CPU?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/minusYCore/status/2058488491399356688
+- Author: Mongo (@minusYCore)
+- Posted: 2026-05-24 10:01:15
+
+## Branch
+
+**1/** **@minusYCore** ^2058488491399356688
+
+**@AgileJebrim**
+
+Does the array initally exist in GPU memory or CPU?
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,40 @@
+---
+title: "@AgileJebrim efficient reductions are a very deep rabbit hole even on single gpu https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf"
+type: archive
+source: twitter
+source_url: "https://x.com/miru_why/status/2058526448550608941"
+author: "miru"
+handle: miru_why
+post_id: "2058526448550608941"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim efficient reductions are a very deep rabbit hole even on single gpu https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/miru_why/status/2058526448550608941
+- Author: miru (@miru_why)
+- Posted: 2026-05-24 12:32:05
+
+## Branch
+
+**1/** **@miru_why** ^2058526448550608941
+
+**@AgileJebrim**
+
+efficient reductions are a very deep rabbit hole even on single gpu https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
+
+![](https://pbs.twimg.com/media/HJFa5VWaEAAGu1x?format=jpg&name=orig)
+![](https://pbs.twimg.com/media/HJFa5VYbQAIn9zW?format=jpg&name=orig)
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,49 @@
+---
+title: "@AgileJebrim i see it as a convolution task."
+type: archive
+source: twitter
+source_url: "https://x.com/noinodev/status/2058439885619622196"
+author: "noino🔸"
+handle: noinodev
+post_id: "2058439885619622196"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim i see it as a convolution task."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/noinodev/status/2058439885619622196
+- Author: noino🔸 (@noinodev)
+- Posted: 2026-05-24 06:48:07
+
+## Branch
+
+**1/** **@noinodev** ^2058439885619622196
+
+**@AgileJebrim**
+
+i see it as a convolution task. gather per thread in registers, gather in local thread group shmem, and atomic add assuming atomics are available. repeat for all cells, then keep scheduler happy by not going over register or thread limits. is there more to it than that?
+
+**2/** **@AgileJebrim** ^2058534147061109052
+
+**@noinodev**
+
+Give more details on each step. Assume more data elements than parallel hardware.
+
+**3/** **@AgileJebrim** ^2058534304808902987
+
+**@noinodev**
+
+Don’t assume you can use atomics.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "Use a parallel reduction, not atomics from every thread."
+type: archive
+source: twitter
+source_url: "https://x.com/ourtown2/status/2058459472557990030"
+author: "Brian Crabtree"
+handle: ourtown2
+post_id: "2058459472557990030"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "Use a parallel reduction, not atomics from every thread."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/ourtown2/status/2058459472557990030
+- Author: Brian Crabtree (@ourtown2)
+- Posted: 2026-05-24 08:05:57
+
+## Branch
+
+**1/** **@ourtown2** ^2058459472557990030
+
+Use a parallel reduction, not atomics from every thread.
+
+Split the array into blocks. Each thread loads one or more elements, accumulates a local partial sum in registers, then the block reduces those partials using warp-level shuffles and shared memory. Each block writes one partial sum to an output array. Then launch a second reduction over that smaller partial array until only one value remains.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim tree reduction?"
+type: archive
+source: twitter
+source_url: "https://x.com/saladpalad/status/2058442502773068209"
+author: "piztachios"
+handle: saladpalad
+post_id: "2058442502773068209"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim tree reduction?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/saladpalad/status/2058442502773068209
+- Author: piztachios (@saladpalad)
+- Posted: 2026-05-24 06:58:31
+
+## Branch
+
+**1/** **@saladpalad** ^2058442502773068209
+
+**@AgileJebrim**
+
+tree reduction?
+
+**2/** **@AgileJebrim** ^2058533901421715628
+
+**@saladpalad**
+
+Need details.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim Answering blind: hierarchically."
+type: archive
+source: twitter
+source_url: "https://x.com/steverockan/status/2058552472893837552"
+author: "mastodon.gamedev.place/@Steverock"
+handle: steverockan
+post_id: "2058552472893837552"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Answering blind: hierarchically."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/steverockan/status/2058552472893837552
+- Author: mastodon.gamedev.place/@Steverock (@steverockan)
+- Posted: 2026-05-24 14:15:30
+
+## Branch
+
+**1/** **@steverockan** ^2058552472893837552
+
+**@AgileJebrim**
+
+Answering blind: hierarchically. Split into chunks as large as possible, sum up chunks in a single pass, write results to a new texture, repeat until 1 pixel texture. Maybe have to becareful about precision.
+
+**2/** **@AgileJebrim** ^2058554730767335810
+
+**@steverockan**
+
+Textures would be slow.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,101 @@
+---
+title: "@AgileJebrim They need to know how GPU shaders are organized to even answer this."
+type: archive
+source: twitter
+source_url: "https://x.com/stikves/status/2058434571608490047"
+author: "sukru tikves"
+handle: stikves
+post_id: "2058434571608490047"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim They need to know how GPU shaders are organized to even answer this."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/stikves/status/2058434571608490047
+- Author: sukru tikves (@stikves)
+- Posted: 2026-05-24 06:27:00
+
+## Branch
+
+**1/** **@stikves** ^2058434571608490047
+
+**@AgileJebrim**
+
+They need to know how GPU shaders are organized to even answer this.
+
+But I hope nobody came with "I would write a simple for loop..."
+
+**2/** **@AgileJebrim** ^2058534593200894022
+
+**@stikves**
+
+Actually the correct answer does involve one or more for loops in the answer.
+
+**3/** **@stikves** ^2058680013483049048
+
+**@AgileJebrim**
+
+(Going down the rabbit hole)
+
+Yes, it would of course. 
+Some context goes out in short replies here.
+
+**4/** **@AgileJebrim** ^2058680437787041994
+
+**@stikves**
+
+You can comment longer. :P
+
+**5/** **@stikves** ^2058682097506931028
+
+**@AgileJebrim**
+
+I could... 
+
+Yet, at doing those at midnight usually backfires.
+"What was a thread block again? Or was it a..."
+
+**6/** **@AgileJebrim** ^2058684292662735191
+
+**@stikves**
+
+Here’s your hardware target. You have an arbitrary number of these SMs and a L2 cache and global memory available for cross communication between them. Describe how you would do it in relation to the hardware described here.
+
+![](https://pbs.twimg.com/media/HJHqc7dWkAEra7E?format=jpg&name=orig)
+
+**7/** **@stikves** ^2058687794470592548
+
+**@AgileJebrim**
+
+Since they work independently, you start from the nearest (register file) and move up (shared mem) to intra SM (L2)
+
+Though to be fair, for standard operations I would look this up
+
+Why?
+
+There is usually one corner case the prior art discovered but not obvious at first sight
+
+**8/** **@stikves** ^2058688045059301875
+
+**@AgileJebrim**
+
+(The standard solution to binary search for example had a off by one bug in the books for over a decade.
+
+Can I write binary search from the tip of my head?
+Of course, did 100s of times
+
+Will I actually do it for production?
+No)
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,67 @@
+---
+title: "@AgileJebrim mapreduce."
+type: archive
+source: twitter
+source_url: "https://x.com/tomcr2100/status/2058461731639161009"
+author: "Tom Keresztes"
+handle: tomcr2100
+post_id: "2058461731639161009"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim mapreduce."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/tomcr2100/status/2058461731639161009
+- Author: Tom Keresztes (@tomcr2100)
+- Posted: 2026-05-24 08:14:55
+
+## Branch
+
+**1/** **@tomcr2100** ^2058461731639161009
+
+**@AgileJebrim**
+
+mapreduce. simple enough answer ? 😉
+
+**2/** **@AgileJebrim** ^2058530982907490419
+
+**@tomcr2100**
+
+Design its implementation.
+
+**3/** **@tomcr2100** ^2059204310630568267
+
+**@AgileJebrim**
+
+start with processing in N chunks, every chunk is a wave, loading wide to feed ALU, and avoid trashing cache by different wave reads.  Waves write to a fixed location in a buffer of size N. Works for sets that don't fit into memory (transfers). Overflow TBD. Out of characters👾
+
+**4/** **@AgileJebrim** ^2059273072255844710
+
+**@tomcr2100**
+
+I don’t see any additions or reductions discussed, just memory loads and stores.
+
+**5/** **@tomcr2100** ^2059311468978262438
+
+**@AgileJebrim**
+
+Sorry, typical twitter superficial moment. I meant each block reduces a portion of the array, multiple kernel invocations to reduce sync, and try to rely on sequential addressing d[i] + d[i + offset] to avoid bank conflicts when writing the result to global memory.
+
+**6/** **@AgileJebrim** ^2059311961096020211
+
+**@tomcr2100**
+
+Can you show me what the first for loop looks like?
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim I had this problem as an example in a uni class covering cuda, curious how their solution and explanation holds up"
+type: archive
+source: twitter
+source_url: "https://x.com/uvadin/status/2058652215879794914"
+author: "UvadinVeliki"
+handle: uvadin
+post_id: "2058652215879794914"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I had this problem as an example in a uni class covering cuda, curious how their solution and explanation holds up"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/uvadin/status/2058652215879794914
+- Author: UvadinVeliki (@uvadin)
+- Posted: 2026-05-24 20:51:50
+
+## Branch
+
+**1/** **@uvadin** ^2058652215879794914
+
+**@AgileJebrim**
+
+I had this problem as an example in a uni class covering cuda, curious how their solution and explanation holds up
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,49 @@
+---
+title: "What precision are the numbers? Integer or floating point? Do we care about numerical stability? Can we assume we're on an NVIDIA GPU with threads, warps, CTAs, and grids? How much memory bandwidth do we have? Is that array of numbers big or really big?"
+type: archive
+source: twitter
+source_url: "https://x.com/xlrndo/status/2058691325302894659"
+author: "xlrndo"
+handle: xlrndo
+post_id: "2058691325302894659"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "What precision are the numbers? Integer or floating point? Do we care about numerical stability? Can we assume we're on an NVIDIA GPU with threads, warps, CTAs, and grids? How much memory bandwidth do we have? Is that array of numbers big or really big?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/xlrndo/status/2058691325302894659
+- Author: xlrndo (@xlrndo)
+- Posted: 2026-05-24 23:27:15
+
+## Branch
+
+**1/** **@xlrndo** ^2058691325302894659
+
+What precision are the numbers? Integer or floating point? Do we care about numerical stability? Can we assume we're on an NVIDIA GPU with threads, warps, CTAs, and grids? How much memory bandwidth do we have? Is that array of numbers big or really big?
+
+Too many questions to have any kind of single answer. It's not a bad interview question if it's used as a jumping off point for the candidate to demonstrate their knowledge.
+
+**2/** **@AgileJebrim** ^2058691753851384096
+
+**@xlrndo**
+
+https://x.com/agilejebrim/status/2058562643158352101?s=46&t=ZbNecpPBsBMrvteeaA6IcA
+
+**3/** **@AgileJebrim** ^2058692267641077855
+
+**@xlrndo**
+
+Portable across a wide range of GPU hardware. You can assume 32-wide waves, 256-wide workgroups, and 64KB shared memory. The memory bandwidth is unlikely to affect the answer.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,47 @@
+---
+title: "@AgileJebrim As a curious bystander: can you give a rough idea how suboptimal it is to do \"the naive thing\"?"
+type: archive
+source: twitter
+source_url: "https://x.com/DemetriSpanos/status/2058774747068891207"
+author: "Demetri Spanos"
+handle: DemetriSpanos
+post_id: "2058774747068891207"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim As a curious bystander: can you give a rough idea how suboptimal it is to do \"the naive thing\"?"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/DemetriSpanos/status/2058774747068891207
+- Author: Demetri Spanos (@DemetriSpanos)
+- Posted: 2026-05-25 04:58:44
+
+## Branch
+
+**1/** **@DemetriSpanos** ^2058774747068891207
+
+**@AgileJebrim**
+
+As a curious bystander: can you give a rough idea how suboptimal it is to do "the naive thing"?
+
+I mean one launch independently accumulating a vec4 per lane (Nx4M matrix yields 1xM array of vec4) and then a second launch to accumulate the 1xM in one lane?
+
+2x suboptimal? worse?
+
+**2/** **@AgileJebrim** ^2058776867981648367
+
+Accumulating a partial sum in a vec4 per lane is a great approach. I think you’re the first one here to suggest it that didn’t attempt to prematurely reduce the vec4 every iteration.
+
+There’s a significantly faster, albeit more involved way to reduce all of those parallel results than running a single lane in a second pass though. If you have 2000+ lanes across the hardware and 4 per lane, you’re looking at a serial accumulation in a single lane taking 8000+ steps on a half clock rate processor. You could do that in a small handful (~less than a dozen off the top of my head) instead.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,81 @@
+---
+title: "@AgileJebrim Using a hierarchical/structured layers on-device parallel reduction of cascading (multi-pass) kernels."
+type: archive
+source: twitter
+source_url: "https://x.com/JesseStojan/status/2058768097939472658"
+author: "Jesse S"
+handle: JesseStojan
+post_id: "2058768097939472658"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Using a hierarchical/structured layers on-device parallel reduction of cascading (multi-pass) kernels."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/JesseStojan/status/2058768097939472658
+- Author: Jesse S (@JesseStojan)
+- Posted: 2026-05-25 04:32:19
+
+## Branch
+
+**1/** **@JesseStojan** ^2058768097939472658
+
+**@AgileJebrim**
+
+Using a hierarchical/structured layers on-device parallel reduction of cascading (multi-pass) kernels.
+E.g. `hip/cub::DeviceReduce::Sum(...)`/`sycl::Reduction(...)` with more steps.
+(I saw you say 16M × uint8 in the replies, but it was one of the first questions in my head lol)
+
+**2/** **@AgileJebrim** ^2058769061480206761
+
+**@JesseStojan**
+
+Cool. Provide more details please.
+
+**3/** **@AgileJebrim** ^2058769188873732484
+
+**@JesseStojan**
+
+You’re implementing the low level details, no reusing an existing function.
+
+**4/** **@JesseStojan** ^2058785099756265546
+
+**@AgileJebrim**
+
+Well in that case, 1st kernel pass, each workgroup strides through the global array, sums the values, tree reduction in shared memory, writes the partial sum to a (uint32/64) partials[blockN] buffer. 2nd kernel pass reduces the partials[] down to the final sum.
+
+**5/** **@AgileJebrim** ^2058788993848766932
+
+**@JesseStojan**
+
+Tell me more about this step labeled “sums the values”
+
+**6/** **@JesseStojan** ^2058803240313979237
+
+**@AgileJebrim**
+
+Each thread sums many values using a stride (num_workgroups × threads_per_workgroup). Within each iteration, neighboring threads in a (subgroup/warp/wavefront) access consecutive addresses, enabling coalesced wide bursts (one transaction per subgroup).
+
+**7/** **@AgileJebrim** ^2058922041076592962
+
+**@JesseStojan**
+
+Do you think you can write a short for loop snippet for this in a shader language?
+
+**8/** **@JesseStojan** ^2058930323946152157
+
+**@AgileJebrim**
+
+Yeah, though just fair heads up, I'm not seeking employment lol, I'm just a nerd for rendering/CGI and massively parallel computing.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,49 @@
+---
+title: "@AgileJebrim Nvidia has good resources on this https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf"
+type: archive
+source: twitter
+source_url: "https://x.com/Karyuutensei/status/2058799615265132885"
+author: "Nick Tasios"
+handle: Karyuutensei
+post_id: "2058799615265132885"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim Nvidia has good resources on this https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Karyuutensei/status/2058799615265132885
+- Author: Nick Tasios (@Karyuutensei)
+- Posted: 2026-05-25 06:37:33
+
+## Branch
+
+**1/** **@Karyuutensei** ^2058799615265132885
+
+**@AgileJebrim**
+
+Nvidia has good resources on this https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
+
+**2/** **@AgileJebrim** ^2058922734327926914
+
+**@Karyuutensei**
+
+Yeah I’ve seen this before. They just waste time describing 6 shitty complex approaches before bothering to discuss the simplest and fastest approach lol.
+
+**3/** **@Karyuutensei** ^2058943120532750836
+
+**@AgileJebrim**
+
+I actually find it quite educative.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,43 @@
+---
+title: "@AgileJebrim I would allocate a variable in VRAM, then partition or map the array into a texture (or two)."
+type: archive
+source: twitter
+source_url: "https://x.com/ParkerRedford/status/2058843627053363467"
+author: "Parker Redford"
+handle: ParkerRedford
+post_id: "2058843627053363467"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I would allocate a variable in VRAM, then partition or map the array into a texture (or two)."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/ParkerRedford/status/2058843627053363467
+- Author: Parker Redford (@ParkerRedford)
+- Posted: 2026-05-25 09:32:26
+
+## Branch
+
+**1/** **@ParkerRedford** ^2058843627053363467
+
+**@AgileJebrim**
+
+I would allocate a variable in VRAM, then partition or map the array into a texture (or two). Have the GPU add the pixel values to the allocated variable in VRAM.
+
+**2/** **@AgileJebrim** ^2058917284773593591
+
+**@ParkerRedford**
+
+-
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,73 @@
+---
+title: "Use OpenMP :^)"
+type: archive
+source: twitter
+source_url: "https://x.com/Waffl3x/status/2058821469434523920"
+author: "Waffl3x ❤️‍🩹 🩹 👁‍🗨"
+handle: Waffl3x
+post_id: "2058821469434523920"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "Use OpenMP :^)"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/Waffl3x/status/2058821469434523920
+- Author: Waffl3x ❤️‍🩹 🩹 👁‍🗨 (@Waffl3x)
+- Posted: 2026-05-25 08:04:24
+
+## Branch
+
+**1/** **@Waffl3x** ^2058821469434523920
+
+Use OpenMP :^)
+Target region, map the memory, use parallel for with a reduction clause, yell at my compiler vendor if it fails to vectorize.
+Or just use Fortran (also with OpenMP).
+
+Answer changes if you need non-overflowing semantics. I am honestly not sure how I would do it efficiently in a scenario that effectively needs a bigint.
+
+I imagine you chunk it in such a way that lets you use native width ints for the largest proportion of the calculation possible
+
+**2/** **@AgileJebrim** ^2058921496240627917
+
+**@Waffl3x**
+
+OpenMP is for CPUs and isn’t really designed for the realtime workflows we require. ISPC would be better and actually has hard guarantees about vectorization.
+
+Our earlier tech was built on top of ISPC. Our next gen tech is in a custom shader language built for GPU-native apps.
+
+**3/** **@AgileJebrim** ^2058921653338361857
+
+**@Waffl3x**
+
+The goal should be to avoid BIGINT and have the input values be defined as smaller bitwidths instead.
+
+**4/** **@Waffl3x** ^2058961654243827959
+
+**@AgileJebrim**
+
+Mmm I can see it, also I suppose bigint isn't necessary anyway, 128bit is almost certainly goof enough for a sum... I hope.
+
+**5/** **@AgileJebrim** ^2058964499680420089
+
+**@Waffl3x**
+
+https://x.com/agilejebrim/status/2058562643158352101?s=46&t=ZbNecpPBsBMrvteeaA6IcA
+
+**6/** **@Waffl3x** ^2058969433499111811
+
+**@AgileJebrim**
+
+My mental model for this was off, it really does fit in a 32bit unsigned huh? OpenMP does this efficiently trivially I'm pretty sure, but that's just because it's such a simple case. So my answer really isn't useful for the interviewer.
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim when we were first evaluating \"compute shaders\" at sony it was my job to find the ideal answer to this very question, and i got it right."
+type: archive
+source: twitter
+source_url: "https://x.com/bmcnett/status/2058724377135558802"
+author: "bmcnett"
+handle: bmcnett
+post_id: "2058724377135558802"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim when we were first evaluating \"compute shaders\" at sony it was my job to find the ideal answer to this very question, and i got it right."
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/bmcnett/status/2058724377135558802
+- Author: bmcnett (@bmcnett)
+- Posted: 2026-05-25 01:38:35
+
+## Branch
+
+**1/** **@bmcnett** ^2058724377135558802
+
+**@AgileJebrim**
+
+when we were first evaluating "compute shaders" at sony it was my job to find the ideal answer to this very question, and i got it right. not everyone came around right away..
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,37 @@
+---
+title: "@AgileJebrim I have no idea about GPU hardware and how it works, can’t answer"
+type: archive
+source: twitter
+source_url: "https://x.com/sandovin34721/status/2058729746322620775"
+author: "kurku"
+handle: sandovin34721
+post_id: "2058729746322620775"
+date: 2026-05-25
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "@AgileJebrim I have no idea about GPU hardware and how it works, can’t answer"
+in_reply_to: ""
+parent_post_id: "2058420153348391346"
+---
+
+## Source
+
+- URL: https://x.com/sandovin34721/status/2058729746322620775
+- Author: kurku (@sandovin34721)
+- Posted: 2026-05-25 01:59:55
+
+## Branch
+
+**1/** **@sandovin34721** ^2058729746322620775
+
+**@AgileJebrim**
+
+I have no idea about GPU hardware and how it works, can’t answer
+
+## Related
+
+- Spine: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
@@ -0,0 +1,36 @@
+---
+title: "I’ve taken to using a simple interview question with folks."
+type: archive
+source: twitter
+source_url: "https://x.com/AgileJebrim/status/2058420153348391346"
+author: "Jebrim"
+handle: AgileJebrim
+post_id: "2058420153348391346"
+date: 2026-05-24
+archived: 2026-08-29
+draft: false
+tags:
+  - archive
+  - twitter
+  - AgileJebrim
+description: "I’ve taken to using a simple interview question with folks."
+in_reply_to: ""
+---
+
+## Source
+
+- URL: https://x.com/AgileJebrim/status/2058420153348391346
+- Author: Jebrim (@AgileJebrim)
+- Posted: 2026-05-24 05:29:42
+
+## Thread
+
+**1/** **@AgileJebrim** ^2058420153348391346
+
+I’ve taken to using a simple interview question with folks. It’s surprising how much people seem to struggle with answering it.
+
+Tell me how you would efficiently sum a large array of numbers on a GPU into a single accumulated value.
+
+Don’t need to see code, just explain it conceptually in relation to the hardware.
+
+Branches: [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-stikves-they-need-to-know-how-gpu-shaders-are-organized]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-noinodev-i-see-it-as-a-convolution-task-gather-per-thread]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-saladpalad-tree-reduction]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-brainage19-im-bored-so-im-gonna-give-this-question-a-good]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-atorstling-hierarchical-map-reduce]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-Gombally_-i-think-this-is-a-simple-example-of-look-how-much]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-Lotsoffoodinme-many-people-dont-write-code-that-touches-the-gpu]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-ourtown2-use-a-parallel-reduction-not-atomics-from-every]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-tomcr2100-mapreduce-simple-enough-answer]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-VMises76153-i-would-embarrass-myself-by-trying-to-answer]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-heartoneto2012-do-floor-warpsize-arraysize-sums-keep-using-the]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-dadbod2100-its-for-a-frontend-javascript-position-i-see]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-FUZxxl-they-asked-the-same-question-to-me-at-nvidia-its]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-gulbaru-i-am-not-knowlegable-enough-with-gpus-but-does]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-feikcel-divide-into-small-arrays-conquer-each-array-on-a]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-MikeTreit-this-is-an-idiotic-interview-question]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-Jason_Ewton-m-x-n-multipass-reduction]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-minusYCore-does-the-array-initally-exist-in-gpu-memory-or-cpu]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-fatroom-no-idea-about-gpu-programming-cpu-hints-towards]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-2Sexy4MyGPU-lots-and-lots-of-atomics]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-deleted_1m-what-kind-of-gpu]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-kechogarcia-easy-waveactivesum-then-lds-atomic-sum-then]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-miru_why-efficient-reductions-are-a-very-deep-rabbit-hole]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-braincell__-if-you-can-fit-qwen-3-6-into-your-gpu-just-load]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-Apolloknius-shrink-the-array-over-and-over-until-its-down-to]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-matt_timmermans-divide-up-the-available-memory-between-cores-each]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-steverockan-answering-blind-hierarchically-split-into-chunks]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-CodeAndTears-large-array-of-values-is-stored-in-float-double]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-Lord_KX_-then-make-them-build-it-in-minecraft-using-pre]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-jes5199-you-shouldnt-do-that-on-the-gpu-do-it-on-the]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-SpikyCaterpillr-i-have-no-meaningful-gpu-coding-experience-that]], [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question/2026-05-24-uvadin-i-had-this-problem-as-an-example-in-a-uni-class]], [[archLine truncated
@@ -12,5 +12,6 @@ tags:
 - [[archive/threads/AgileJebrim/2024-12-18-you-left-out-the-bit-where-i-also-said-no]]
 - [[archive/threads/AgileJebrim/2025-02-09-check-your-compiler]]
 - [[archive/threads/AgileJebrim/2026-05-12-we-have-code-reviews-at-isochron-and-i-require-my]]
+- [[archive/threads/AgileJebrim/2026-05-24-ive-taken-to-using-a-simple-interview-question]]
 - [[archive/threads/AgileJebrim/2026-05-30-i-dont-think-hardware-vendors-are-going-to]]
 - [[archive/threads/AgileJebrim/2026-08-22-start-with-spir-v-or-ptx-instead]]