Olivia
Online
Pavel Simó
@pavelsimo
if a tree falls in a forest and no one is around to hear it, does it make a sound?
Berlin, Germany
Joined August 2011
579 Following
67 Followers
103 Posts
https://t.co/Oz7KBj1Ppk
Day 18/100 of learning CUDA
> solved the LeetGPU FP16 Dot Product problem and tried to optimize it as much as possible using warp shuffles and vectorized float4.
> i'm still amazed by how much effort it takes to write a performant dot product. the end result is pretty much unreadable to someone with no CUDA experience. a few things make it feel "magical". for instance, __shfl_down_sync is really weird. it blows my mind that you can copy a value directly from another thread register in the same warp.
> it also takes some effort not to mess up the indices once you start vectorizing (float4, half2, ...). you have to read more data at a time, so your loop conditions change, and you also have to take care of boundary conditions because the data size is not always a multiple of the vector size, etc.
link to the code below
Day 17/100 of learning CUDA
> solved the 2D version of yesterday's problem: 2D Subarray Sum, again with three implementations and different levels of optimization (source code below).
> finished chapter 4 of Programming Massively Parallel Processors, link to my notes: https://t.co/OT3qrAN9Fc
https://t.co/Oz7KBj1Ppk
Who to follow
Day 17/100 of learning CUDA
> solved the 2D version of yesterday's problem: 2D Subarray Sum, again with three implementations and different levels of optimization (source code below).
> finished chapter 4 of Programming Massively Parallel Processors, link to my notes: https://t.co/OT3qrAN9Fc
Day 16/100 of learning CUDA
> solved another LeetGPU problem: Subarray Sum, the problem is interesting enough so i ended up implementing it in three different ways. the first approach calculates the prefix sum and uses the generated array to compute the answer. it is a bit of overkill, but it gives O(1) queries once the prefix sum array is pre-calculated with a kernel. the other two approaches are classical reductions, we just need to make sure any value we consider are within the range [S, E], otherwise we assign a 0 to the new input array.
> read a few more pages of Programming Massively Parallel Processors, almost done with chapter 4.
![pavelsimo's tweet photo. Day 16/100 of learning CUDA
> solved another LeetGPU problem: Subarray Sum, the problem is interesting enough so i ended up implementing it in three different ways. the first approach calculates the prefix sum and uses the generated array to compute the answer. it is a bit of overkill, but it gives O(1) queries once the prefix sum array is pre-calculated with a kernel. the other two approaches are classical reductions, we just need to make sure any value we consider are within the range [S, E], otherwise we assign a 0 to the new input array.
> read a few more pages of Programming Massively Parallel Processors, almost done with chapter 4.](https://pbs.twimg.com/media/HO03GcXWcAAYjKS.jpg)
Day 16/100 of learning CUDA
> solved another LeetGPU problem: Subarray Sum, the problem is interesting enough so i ended up implementing it in three different ways. the first approach calculates the prefix sum and uses the generated array to compute the answer. it is a bit of overkill, but it gives O(1) queries once the prefix sum array is pre-calculated with a kernel. the other two approaches are classical reductions, we just need to make sure any value we consider are within the range [S, E], otherwise we assign a 0 to the new input array.
> read a few more pages of Programming Massively Parallel Processors, almost done with chapter 4.
Day 15/100 of learning CUDA
> no new problems today, i revisited the LeetGPU GEMM problem. i wanted to improve the naive implementation from yesterday, and i found this blog post by Simon Boehm: How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog: https://t.co/5J4nVLsUoJ
> it's an excellent step-by-step guide, going from a naive GEMM kernel to something close to cuBLAS, explaining the decisions behind each optimization, highly recommended!
> i had the time to apply the first two optimizations from the article: Global Memory Coalescing and Shared Memory Caching, which took the naive solution from 26 ms to 6.85 ms. still a long way to go... the top solutions in LeetGPU are in the 0.15 - 0.60 ms range.
link to the source code below
Meet Qwen3.8-Max: A New Bar for Coding and Cowork.
@elliotarledge https://t.co/RjOCduzaAH
https://t.co/Oz7KBj1Ppk
Day 15/100 of learning CUDA
> no new problems today, i revisited the LeetGPU GEMM problem. i wanted to improve the naive implementation from yesterday, and i found this blog post by Simon Boehm: How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog: https://t.co/5J4nVLsUoJ
> it's an excellent step-by-step guide, going from a naive GEMM kernel to something close to cuBLAS, explaining the decisions behind each optimization, highly recommended!
> i had the time to apply the first two optimizations from the article: Global Memory Coalescing and Shared Memory Caching, which took the naive solution from 26 ms to 6.85 ms. still a long way to go... the top solutions in LeetGPU are in the 0.15 - 0.60 ms range.
link to the source code below
Day 14/100 of learning CUDA
> solved two more LeetGPU medium problems: Softmax Attention and GEMM (FP16).
> the GEMM (General matrix multiply) one is a naive implementation of matmul (for now), the problem requires to use half precision, but you need to make sure all calculations are done in float to avoid rounding errors.
> the softmax attention solution i came up with is 3 kernels: (1) the portion that multiplies the queries (Q) by the keys transpose (K^T), (2) the softmax part, and (3) finally a matmul with the values (V). technically i could have used the same matrix multiplication kernel for both 1 and 3, but then a memory allocation would be needed for K^T, so i decided to just keep those two separated (animation below).
> finished chapter 3 of Programming Massively Parallel Processors, link to my notes: https://t.co/PAsOFOyswF
> started chapter 4. in the next few days i'll revisit some of my LeetGPU solutions and optimize them with what i've learned from the book.
link to the source code below
@Rasmic nice, congrats!
https://t.co/Oz7KBj1Ppk
Day 14/100 of learning CUDA
> solved two more LeetGPU medium problems: Softmax Attention and GEMM (FP16).
> the GEMM (General matrix multiply) one is a naive implementation of matmul (for now), the problem requires to use half precision, but you need to make sure all calculations are done in float to avoid rounding errors.
> the softmax attention solution i came up with is 3 kernels: (1) the portion that multiplies the queries (Q) by the keys transpose (K^T), (2) the softmax part, and (3) finally a matmul with the values (V). technically i could have used the same matrix multiplication kernel for both 1 and 3, but then a memory allocation would be needed for K^T, so i decided to just keep those two separated (animation below).
> finished chapter 3 of Programming Massively Parallel Processors, link to my notes: https://t.co/PAsOFOyswF
> started chapter 4. in the next few days i'll revisit some of my LeetGPU solutions and optimize them with what i've learned from the book.
link to the source code below
Day 13/100 of learning CUDA
> solved another LeetGPU medium problem: Prefix Sum. the trickiest one so far... i was not able to solve it on my own, i watched the lectures Prefix Sum Scan Part 1 & Part 2 (links in the comments).
> til:
> the prefix sum is part of a problem family called "Scan". the name comes from the scan operator in Ken Iverson's APL (1962): https://t.co/Ga8Y9xgZG0
> many libraries implement scan operations: CUB -> BlockScan and DeviceScan, Thrust -> inclusive_scan, C++17 -> std::inclusive_scan, and MPI -> MPI_Scan.
https://t.co/Oz7KBj1Ppk
Day 13/100 of learning CUDA
> solved another LeetGPU medium problem: Prefix Sum. the trickiest one so far... i was not able to solve it on my own, i watched the lectures Prefix Sum Scan Part 1 & Part 2 (links in the comments).
> til:
> the prefix sum is part of a problem family called "Scan". the name comes from the scan operator in Ken Iverson's APL (1962): https://t.co/Ga8Y9xgZG0
> many libraries implement scan operations: CUB -> BlockScan and DeviceScan, Thrust -> inclusive_scan, C++17 -> std::inclusive_scan, and MPI -> MPI_Scan.
Day 12/100 of learning CUDA
> solved two more LeetGPU medium problems: Count 2D Array Element and Count 3D Array Element. same reduction pattern.
> til:
> warp-level primitives: threads in the same warp can read each others registers with __shfl_down_sync, no shared memory needed. used it in a new version of Reduce.
> memory coalescing, how the gpu merges a warps loads into one request and fetches as few 128-byte segments as possible.
> shared memory bank conflicts: shared memory __shared__ is split into 32 banks, and when threads in a warp hit different words in the same bank the access becomes sequential, one thread at a time.
link to the source code below
Prefix Sum Scan - Part 2
https://t.co/ivOJMEZEsM
Prefix Sum Scan - Part 1
https://t.co/4nfbXJlQGv
https://t.co/Oz7KBj1Ppk
Day 12/100 of learning CUDA
> solved two more LeetGPU medium problems: Count 2D Array Element and Count 3D Array Element. same reduction pattern.
> til:
> warp-level primitives: threads in the same warp can read each others registers with __shfl_down_sync, no shared memory needed. used it in a new version of Reduce.
> memory coalescing, how the gpu merges a warps loads into one request and fetches as few 128-byte segments as possible.
> shared memory bank conflicts: shared memory __shared__ is split into 32 banks, and when threads in a warp hit different words in the same bank the access becomes sequential, one thread at a time.
link to the source code below
Day 11/100 of learning CUDA
> solved three more LeetGPU medium problems: 2D Convolution (see animation below), Gaussian Blur and Jacobi Stencil. nothing super special about these, but they demand being extremely careful with the indices and boundary conditions. i'm starting to feel comfortable with row-major indexing. would like to try optimized versions of these later, currently my solutions are super naive.
> read the article A Gentle Introduction to CUDA PTX: https://t.co/Phw9KSnrj8, indeed a really "gentle" introduction to PTX: it covers the basic instructions, where PTX fits in the CUDA compilation pipeline, and the cli commands to inspect and work with it. highly recommended. while reading found a NVIDIA guide to Inline PTX Assembly: https://t.co/IV8pKzavGu, which teach you how to inline PTX in your kernels. i'm not at that level of ninja yet, but soon...
> started chapter 3 of Programming Massively Parallel Processors. i've already picked up some of the chapter concepts by doing LeetGPU problems, but it's good to take a step back and revisit them properly.
link to the source code below
Last Seen Users on Sotwe
Most Popular Users
Elon Musk
@elonmusk
241.2M followers
Barack Obama
@barackobama
119.1M followers
Cristiano Ronaldo
@cristiano
112.7M followers
Donald J. Trump
@realdonaldtrump
111.8M followers
Narendra Modi
@narendramodi
107.1M followers
Rihanna
@rihanna
98.2M followers
NASA
@nasa
92.2M followers
Justin Bieber
@justinbieber
91.4M followers
KATY PERRY
@katyperry
88.9M followers
Taylor Swift
@taylorswift13
82.8M followers
Lady Gaga
@ladygaga
74.3M followers
Virat Kohli
@imvkohli
71.8M followers
Kim Kardashian
@kimkardashian
70.4M followers
YouTube
@youtube
68.8M followers
Neymar Jr
@neymarjr
64.7M followers
Bill Gates
@billgates
64.6M followers
The Ellen Show
@theellenshow
62.4M followers
Selena Gomez
@selenagomez
62M followers
CNN
@cnn
61.8M followers
X
@x
60.8M followers