Episode · ADSP: Algorithms + Data Structures = Programs
Episode 258: 🇳🇴 An Algorithm Taxonomy (Serial, Parallel, Cooperative)
31 Oct 2025 · 17 min
Episode notes
In this episode, Conor and Bryce record live from Norway! Bryce explains the taxonomy of algorithms: serial, parallel, and cooperative! Link to Episode 258 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterDate Recorded: 2025-09-23 Date Released: 2025-10-31 MPIIPCRow-wise Softmax in TritonRow-wise Softmax in ParrotCCCL - Parallel and Cooperative AlgorithmsI...
Transcript
Read the transcript · about 2,680 words, follows along as you listen
Bryce:There are three types of algorithms. There are serial algorithms, then there are parallel algorithms. The third category, the third thing, is a cooperative algorithm.
Conor:I guess I was in the nascency of my parallel thinking because before I joined NVIDIA, I was a serial thinker. I only knew about category one. And then I joined NVIDIA. And it probably took me a year or two really to convert to parallel algorithm thinking. But I still have not developed cooperative algorithm thinking.
Bryce:The beautiful thing about the parallel algorithms level is the seemingly sequential nature. The seemingly sequential nature is key.
Conor:When I do reach for cub, I reach for device algorithms. Because Thrust, I feel like our listeners can relate to, even if they've never touched Thrust, because a lot of our listeners are C people. And even if they're not C people, they're Rust people, they're Python people, they know the equivalence of a rotate, a copy, a whatever. But then as as soon as you drop down into block aware algorithms, so what are the what are the what's the taxonomy? And we gotta start there because you lost me and probably some of our listeners.
Conor:Bryce looks excited, folks.
Bryce:There are three types of algorithms. There are serial algorithms in which you a single thread calls the algorithm and um and a single thread performs the algorithm and the result is returned to a single thread. That's what we're all familiar with. Yes. Then there are parallel algorithms in which let's start with cooperative algorithms instead. No, we'll we'll we'll start with parallel algorithms, make more sense to people here. Um then there are parallel algorithms where um you have the appearance of serial execution in a parallel algorithm and in R nomenclature, the um in like thrust and in the C standard parallel algorithms, you have a single thread that a single caller that invokes the algorithm.
Bryce:And the algorithm then creates or acquires parallel execution resources and then runs the result runs the computations on those parallel resources, that collection of threads. And then that produces a result, and then that result is returned to the caller in thrust and in the C standard algorithms, it is returns asynchronously. It's just the return value of the invocation, and so it's completely seemingly sequential. But you could also imagine that this model also could apply to an asynchronous API, something that returns a future or a sender.
Bryce:Uh but the idea is that there's a single caller, a single thread initiates the operation. The third category, the third thing, is a cooperative algorithm. In a cooperative algorithm multiple threads decide to participate. So multiple threads call the primitive at the same time. And no no execution resources are created or acquired by a cooperative algorithm. Instead, the calling threads that have called into this primitive um uh they they donate their resources.
Bryce:So like examples of this is like uh if if people are familiar with with MPI programming, uh like MPI collectives, like an MPI reduce or um you know an MPI barrier, those are collective algorithms where all of the different ranks call it together. Or in in all of the different ranks? All of the different processes in a distributed computation, uh which are called ranks in the in the distributed world. Um or uh you know, if you're doing like like shared memory IPC, like you've got a like you know, you have a IPC stamp.
Bryce:Uh interprocess communication. Like you're using shared memory to have multiple processes communicate, and they're all invoking the same thing. They're all participating in some collective operation together. That's a cooperative algorithm. In the case of CUDA, we in CUDA you launch these kernels, which are parallel functions that are executed by all the threads at once. Within those kernels, you can invoke a cooperative algorithm like a reduction, for example, where a bunch of threads all decide to participate in the reduction.
Bryce:In each thre instead of having an input array, uh what each thread, the inputs are distributed among the threads, and each thread contributes a value into the reduction, or maybe maybe it's reduced some values locally, but it it each thread gives one piece of the input uh to the reduction. Um and all of the threads have to contribute their piece of the input and then contribute to the to the calculation. And these cooperative algorithms are the building blocks that we use to implement the parallel algorithms.
Bryce:And I was speaking in terms of the cooperative level, because uh that's how I often think about algorithms these days. I suppose that's always been how I've thought about algorithms. Um I think that's how I default to thinking about it. You default to thinking about it at the at the serial serial or the parallel level. I I've I've been I always think about it at this level, but then when we converse, I often raise it up to the level that you're used to thinking about it.
Bryce:And you're usually thinking about it at a higher level, a functional level, and then lowering it down for me.
Conor:Well, so what I'll say, folks, this is an inflection point in the podcast. When you I mean, actually, I don't know when you started listening, but back when we started this podcast in 2020? Was it 20? Was it 19? It was definitely well actually I think we've said this before and then we looked it up and it was like, oh no, it was actually midst of the pandemic. But I thought you we decided we were gonna start it right after we went to Belfast. And that was pre-pandemic.
Conor:That was 2019. But we're about to find out, folks. 252 episodes. 2020! I was wrong. I was I knew I was wrong because we've done this before. 2020 November, which was post-pandemic. I think we talked about it at Belfast, but then I didn't start it until a year later. Anyways, when we started recording this, I wasn't allowed. I guess I was in the nascency of my parallel thinking because before I joined NVIDIA, I was a serial thinker. I only knew about category one, and then I joined NVIDIA, and it probably took me a year or two really to convert to parallel algorithm thinking.
Conor:But I still have not developed cooperative algorithm thinking because I primarily work in thrust, and when I do work in Cub, I feel like it's the uh Captain Morgan, you know, when I do drink, I drink whatever.
Dos Equis (Clip):He is the life of parties he has never attended. If he were to punch you in the face, you would have to fight off the strong urge to thank him. Sharks have a week dedicated to him. He is the most interesting man in the world. I don't always drink beer, but when I do, I prefer no snackies.
Conor:Stay thirsty, my friend. When I do reach for cub, I reach for device algorithms, which is the equivalent of thrust. The only algorithms I reach for are the device segmented reduced and device fixed segment reduce, etc. Uh Bryce has something to add.
Bryce:When we started the podcast, I was at the place that you are currently at. I was a few years into NVIDIA, but I was still thinking at the parallel algorithm level, not at the cooperative algorithm level. Well, maybe to some degree I was starting to think at both, but having now five years later, I think about things first from the cooperative algorithm level. Um and and then from the parallel algorithm level. And I think you are very close to jumping off the deep end because of the project that you're working on.
Bryce:I think I think that a year or two from now you're gonna be where I am now. Because right now, you are where I was a couple of years ago.
Conor:I can almost guarantee you that's true because the Cub block primitives are what you use when you're writing kernels. Right. And I am about to start doing that, folks. We've only got segmented scans and reductions, and we need more. And also we don't need the segments. We need fixed segments. We need row-wise and column-wise, everything.
Bryce:With the way we want people to program is at the parallel algorithm level. Um because that is the level that's purest and closest to functional programming and array programming. At the cooperative level, you have to think about uh threads and synchronization. At the cooperative level, you're programming threads. And so it's always gonna be an implementation detail for the work that we do. It's never gonna be the programming model that we expose up to people.
Bryce:And so when you s what you just said, you think that I'm right, you that that you'll have to do that because you need X, you need Y, and Z. That's because you know that you need those things to be able to implement the back end of, like to be able to implement the algorithms that you want, not as a programming model that you'd expose to users. So I think that in terms of this podcast, we will continue to primarily talk about things in terms of the parallel algorithms level, because the beautiful thing about the parallel algorithms level is the seemingly sequential nature.
Bryce:The seemingly sequential nature is key.
Conor:I mean, and this is uh Bryce just admitted something that he may or may not have wanted to admit is that he said, you know, we want people to program at the functional or array level. And so he just said, of course, off off mic. And this is the point, you know, we are C, we are CUDA C programmers, developers. But the array programming model, the functional programming model, it's just easier to reason about. Everybody that's hating on all Haskell, OCam, well, guess what?
Conor:It's a better model, it's just not performant. But guess what? We here, a promise to the listener is we're gonna make the easy way be the efficient way. And that's the dream. That's the dream. And yes, there are trade-offs, but we're gonna make it happen. Bys got something to add.
Bryce:It's not always the case that it's less efficient to write the functional or or the array or just the higher level code. If you write the lower level code, you are baking in certain details and you're making certain assumptions about how the about the like the nature of what you're executing on. Like if you're programming to threads, as you do in a cooperative model, then you're you're making certain assumptions about quantities of threads and uh how those threads are going to communicate with with each other.
Bryce:Whereas if you express things at the arrayslash functional level, you're expressing things purely in terms of like the requirements. You're not saying anything about how it needs to be implemented, right? Like, like it's um it's it's just it's descriptive versus prescriptive, right? Array and functional is descriptive. You're describing what it does. Whereas cooperative is prescriptive, you're saying it must be done this way. And I think the history of parallel programming has actually shown that descriptive models tend to win over prescriptive models.
Conor:There it is, folks. Array, functional, it's the future. And we'll probably talk about cooperative and block algorithms, but I currently don't understand them, so I'm here to hold your hand.
Bryce:You gotta talk about your softmax.
Conor:Uh no, so I was Bryce just said off Mike, we gotta talk about your softmax. Guess what? We're at the 47 minute mark. And the next time, Bryce and I are one-on-one recording tomorrow night, today is Tuesday, September 23rd. Tomorrow will be September 24th. Actually, wait, I'm making a promise I can't keep. The next time we're gonna record is actually probably gonna be in person at uh Under the Sea in the Netherlands. So once we finish recording there, I mean maybe we might steal away and talk about Softmax, but it might be after C Under the Sea.
Conor:Maybe we'll steal away, find some time to talk about Softmax, but I've been working on Ashwin posted internally about a Row Wise Softmax in NumPy and what his aspirations for that were. I tried to do it in I think I can mention it now. It's you just heard a record scratch. Bryce made me cut everything I just said. Uh we'll be talking about softmax and a special project that uh we will be announcing on the pod. Bryce has something to say.
Bryce:Avid listeners will know that I have been hinting and suggesting at least two or three times that eventually on this podcast, where we are clearly obsessed with AI, we are gonna have to train some sort of machine learning model. We're gonna have to go and build something and get under the hoods. And Connor doesn't even realize that he started down this exploration because he's gotten nerd-sniped into implementing this uh softmax. Softmax, which is one of the key ML algorithms at the heart of mini attention mechanisms, which is the Ashra.
Bryce:I'm not saying I nerd sniped, I'm saying Ashra nerd sniped you, but softmax is at the key of mini attention algorithms, such as such as flash attention.
Conor:Four operations exponential, division, I understand.
Bryce:But and I told Connor, this is cool, this will get people very interested in your project if you show them an efficient softmax. But what'll get people even more interested is if you can show them how to implement flash attention in your thing. And then Connor's gonna go and do that, and then next after that, we're gonna have to go and implement some sort of open source model, inference for some sort of open source model in his thing, and then we're gonna end up talking more and more about this.
Bryce:And before you know it, folks, this is gonna be a machine learning podcast, and Connor's gonna have implemented some machine learning models in his thing. I tell you, I predict right now, one to two years from now, that's where we'll be. Training will come after the inference, though. Training will come after the inference.
Conor:I got nothing to say, folks. We got a conference to go to tomorrow.
Bryce:Oh man. Do you think you have Do you think you have the energy for it? Starting to rally. I'm starting to rally. I was I was feeling drained over the weekend, but now I'm starting to rally.
Conor:Today I slept until 7.30. The day before that I was up at 5.45. Luckily, my plan tomorrow is also to sleep into 7.45. Uh so I should be good. I do plan to attend at least the start of the keynote. It's it's not Hannah's. It's not Hana. You predicted it correctly. At one point, was it yesterday, two days ago, you saw that there was a second keynote that had been added. And I I kid you not, Bryce said, hmm, I wonder if Hana can't make it anymore. And sure enough, a day later, uh Hana's keynote was removed and it was replaced by another keynote.
Conor:Anyways, how are you feeling about the conference tomorrow?
Bryce:You got the energy? I'm rallying.
Conor:All right, folks, you will hear from us in a week on Wednesday, September 24th, live from the NDC social. Until then, be sure to check these show notes either in your podcast app or at adspthepodcast.com for links to anything we mentioned in today's episode, as well as a link to a GitHub discussion where you can leave thoughts, comments, and questions. Thanks for listening. We hope you enjoyed and have a great day.
Bryce:Low quality, high quality. That is the tagline of our podcast.
Conor:That's not the tagline. Our tagline is chaos with sprinkles of information.
Transcript supplied by the publisher with the episode.
Podcast
ADSP: Algorithms + Data Structures = Programs
Conor Hoekstra, Bryce Adelstein Lelbach & Ben Deane · Tech & Science · English
Deep dives into algorithms, data structures and C++ — exploring the STL, functional programming and computer science fundamentals with three experienced engineers.
More from ADSP: Algorithms + Data Structures = Programs All episodes →
-
5 Dec 2025 · 42 min
Episode 263: 🇳🇱 The Sean Parent Origin Story
In this episode, Conor and Bryce interview Sean Parent about the origin story of his career in software engineering! Link to Episode 263 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterAbout the Guest: Sean Parent is a senior principal scientist and software architect managing Adobe's Software Technology Lab. Sean first joined Adobe in 1993 working on Pho...
-
28 Nov 2025 · 37 min
Episode 262: 🇳🇱 Chaos with Sean Parent! Bikes, Buses, Waffles & More!
In this episode, Conor and Bryce interview Sean Parent about generic programming and there is much chaos! Link to Episode 262 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterAbout the Guest: Sean Parent is a senior principal scientist and software architect managing Adobe's Software Technology Lab. Sean first joined Adobe in 1993 working on Photoshop and ...
-
21 Nov 2025 · 26 min
Episode 261: 🇳🇱 C++ Under the Sea 🇳🇱 Bernhard, Koen & C++26 Reflection!
In this episode, Conor and Bryce record live from C++ Under the Sea! We interview Bernhard, Koen, talk about C++26 Reflection and more! Link to Episode 261 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterAbout the Guests: Bernhard is a senior system software engineer at NVIDIA, where he extends, optimizes and maintains the CUDA Core Compute Libraries (CCC...
-
14 Nov 2025 · 24 min
Episode 260: 🇳🇱 C++ Under the Sea 🇳🇱 Ray, Paul, Parrot & Scanman!
In this episode, Conor and Bryce record live from C++ Under the Sea! We interview Ray and Paul from NVIDIA, talk about Parrot, scans and more! Link to Episode 260 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterAbout the Guests: Ray is a Senior Systems Software Engineer at NVIDIA since 2022. Studied Software Engineering at the University of Amsterdam. Fou...
-
7 Nov 2025 · 47 min
Episode 259: 🇳🇴 NDC TechTown 🇳🇴 Vittorio Romeo & JF Bastien
In this episode, Conor and Bryce record live from NDC TechTown in Norway! We interview Vittorio Romeo and JF Bastien about C++, training, their talks and more! Link to Episode 259 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterAbout the Guests: Vittorio is a passionate C++ expert with over a decade of professional and personal experience. His expertise c...
-
24 Oct 2025 · 35 min
Episode 257: 🇳🇴 Live from Norway! Replicate, Scatter, Gather & RLD (Part 3)
In this episode, Conor and Bryce record live from Norway! They continue their chat about the replicate, scatter, gather and run length decode algorithms! Link to Episode 257 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterDate Recorded: 2025-09-23 Date Released: 2025-10-24 thrust::gatherthrust::scatterthrust::permutation_iteratorthrust::counting_iteratort...
-
17 Oct 2025 · 41 min
Episode 256: 🇩🇰 Algorithms: Replicate, Scatter, Gather & RLD (Part 2)
In this episode, Conor and Bryce record live from Denmark! They talk about the replicate, scatter, gather and run length decode algorithms! Link to Episode 256 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterDate Recorded: 2025-09-20 Date Released: 2025-10-17 thrust::gatherthrust::scatterthrust::permutation_iteratorthrust::counting_iteratorthrust::sequenc...
-
10 Oct 2025 · 28 min
Episode 255: 🇩🇰 C++ Copenhagen Meetup & Replicate
In this episode, Conor and Bryce record live from Denmark! They recap the C++ Copenhagen Meetup hosted by Symbion, the replicate algorithm and much more! Link to Episode 255 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterDate Recorded: 2025-09-20 Date Released: 2025-10-10 Roku Engineering SymposiumDenmark DGX SuperpodTweet of Beer PouringAPL replicateC++...
-
3 Oct 2025 · 38 min
Episode 254: 🇩🇰 Live from Denmark! Roku, EVs, Aarhus (Ø) and more!
In this episode, Conor and Bryce record live from Denmark! They recap the Roku Engineering Symposium as well as many random topics including Polestar vs Tesla vs Prius, the city of Aarhus and much more! Link to Episode 254 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterDate Recorded: 2025-09-19 Date Released: 2025-10-03 Roku Engineering SymposiumRoku ros...
-
26 Sep 2025 · 25 min
Episode 253: Sean Parent on AI and Cursor
In this episode, Conor and Bryce chat with Sean Parent about AI and Cursor! Link to Episode 253 on WebsiteDiscuss this episode, leave a comment, or ask a question (on GitHub)Socials ADSP: The Podcast: TwitterConor Hoekstra: Twitter | BlueSky | MastodonBryce Adelstein Lelbach: TwitterAbout the Guest: Sean Parent is a senior principal scientist and software architect managing Adobe's Software Technology Lab. Sean first joined Adobe in 1993 working on Photoshop and is one of the creators of Phot...
