levi

617 posts

levi banner
levi

levi

@levidiamode

365 days of GPU programming ▓▓▓▓▓▓▓▓▓▓▓▓▓▓▓▓▓▓▓░░░░░░░░░░░░░░░░ 192/365

out of memory Katılım Haziran 2025
739 Takip Edilen5.8K Takipçiler
Sabitlenmiş Tweet
levi
levi@levidiamode·
Day 83/365 of GPU Programming Looking at DeepSeek's Multi-Head Latent Attention today. The last part of the AMD challenge series is to optimize an MLA decode kernel for MI355X where the absorbed Q and compressed KV cache are given and your task is to do the attention computation. A resource that really helped internalize what MLA does was @rasbt's incredible visual guide to attention variants in LLMs (luckily he posted that last week!), which covers everything from MHA to GQA to MLA to SWA, et cetera. If there's one place to get a visual intuition for recent attention mechanisms, it's this blog post. @jbhuang0604's video on MQA, GQA,MLA and DSA was the best conceptual intro I found on the topic and progressively builds up the ideas from first principles. The Welch Labs analysis of MLA is a great watch as well. Beautiful visualization of the changes DeepSeek made for MLA. Tried out a few kernels once I had a basic understanding of MLA and I think I'm slowly getting more comfortable with at least analyzing kernels.
levi tweet medialevi tweet medialevi tweet medialevi tweet media
levi@levidiamode

Day 82/365 of GPU Programming Taking a closer look at Mixture of Experts today, so I can write better MoE kernels. Specifically, to optimize an MXFP4 MoE fused kernel for the GPU Mode challenge. I haven't had much prior exposure to MoEs, so lots of new concepts I learned today. Luckily I found the best intro to MoEs thanks to @MaartenGr visual overview of the topic. I then watched @tatsu_hashimoto's amazing Stanford CS336 lecture on MoEs, which added deeper context around why MoEs are gaining popularity, FLOPs, OLMoE, infra complexity, routing functions (mindblown this works so well...), expert sizes, training objectives, top k routing and DeepSeek variations. Once I had a basic understanding I started playing around with the some AITER kernels but progress there is tbd. Also had a nice chat with @juscallmevyom (who was kind enough to reach out!) about the AMD kernels and the challenge of materialization overhead.

English
22
145
1.4K
122.2K
Ian
Ian@dumbfook·
@levidiamode apart from the great content the avi is so funny bro
English
1
0
1
30
levi
levi@levidiamode·
192/365 of GPU Programming Learning more about PTX has been really fun. Don't think I'm at a point where I can effectively handwrite PTX yet but it's nice to be able to contextualize the use of PTX and its applications in the kernel stack. I'm especially excited about automated kernel approaches that tackle the PTX layer (e.g. Standard Kernel via condensed representations from various DSLs or Patrick with PyPTX).
levi tweet media
levi@levidiamode

191/365 of GPU Programming Stumbled upon this wonderful introduction to PTX. As someone who never learned assembly in school, PTX always seemed a bit daunting . So I've recently been trying to understand it as well as SASS at a level where I feel more comfortable reading through PTX code in kernels. It's been nice to be able to recognize certain instructions when looking at the profiler. If you've never given PTX a try, would recommend starting with this blog post!

English
3
2
48
1.4K
levi
levi@levidiamode·
@denniszdes no exact roadmap. mainly based on whatever i feel like learning on a particular day
English
0
0
1
12
levi
levi@levidiamode·
191/365 of GPU Programming Stumbled upon this wonderful introduction to PTX. As someone who never learned assembly in school, PTX always seemed a bit daunting . So I've recently been trying to understand it as well as SASS at a level where I feel more comfortable reading through PTX code in kernels. It's been nice to be able to recognize certain instructions when looking at the profiler. If you've never given PTX a try, would recommend starting with this blog post!
levi tweet media
levi@levidiamode

190/365 of GPU Programming Reviewing and continuing my Blackwell TMEM studies. There's an excellent "programming tensor cores" talk by Cris Cecka (creator of CuTe) I've watched a while back and had a hard time following along at that time. Revisited the video today and the sections on TMA and TMEM are much more comprehensible now. As I'm doing this challenge part-time whenever I have time outside of work, it's sometimes difficult for me to feel tangible progress, so it's nice to revisit previous materials that make that clearer.

English
9
32
231
10.3K
levi
levi@levidiamode·
@VipulS_1 amazing, thanks for sharing!
English
0
0
0
17
levi
levi@levidiamode·
@NatKokoromyti these are great, thanks for sharing! love the standard kernel idea
English
0
0
1
42
jun jie
jun jie@junnjiee_·
day 1 of learning databases a system design goal of a disk-oriented DBMS (think Postgres, SQLite) is to manage a DB where its size exceeds available memory, and there are certain aspects that make this possible: 1. memory/data management using buffer pools, this is where the magic happens, to make it seem like the entire DB is in-memory, but in reality there are operations under the hood to fetch and evict pages depending on demand 2. optimizations to maximise sequential I/O, as sequential access is generally faster than random access on-disk 3. optimizations to minimise I/O through compression techniques, choice of storage models and storage approaches. lesser I/O overall would result in lesser stalls, improving performance
English
1
0
1
82
levi
levi@levidiamode·
190/365 of GPU Programming Reviewing and continuing my Blackwell TMEM studies. There's an excellent "programming tensor cores" talk by Cris Cecka (creator of CuTe) I've watched a while back and had a hard time following along at that time. Revisited the video today and the sections on TMA and TMEM are much more comprehensible now. As I'm doing this challenge part-time whenever I have time outside of work, it's sometimes difficult for me to feel tangible progress, so it's nice to revisit previous materials that make that clearer.
levi tweet medialevi tweet medialevi tweet media
levi@levidiamode

189/365 of GPU Programming Taking some time to study Blackwell's Tensor Memory (TMEM) in more depth today (tcgen05, UMMA accumulation, layout/addressing, et cetera). As often is the case, Colfax has a great article on the topic. If you haven't given their blog a try, would really recommend it.

English
1
6
104
13.3K
levi
levi@levidiamode·
189/365 of GPU Programming Taking some time to study Blackwell's Tensor Memory (TMEM) in more depth today (tcgen05, UMMA accumulation, layout/addressing, et cetera). As often is the case, Colfax has a great article on the topic. If you haven't given their blog a try, would really recommend it.
levi tweet medialevi tweet media
levi@levidiamode

188/365 of GPU Programming One of my favorite ways to learn at the moment is to just have a conversation with GPT-live like you would when talking to an expert. I've been using it today to review some of the concepts around Tensor Memory Accelerators and mbarriers today and it was able to go into quite a lot of detail with essentially instantaneous responses. Will probably use this as my default way of reviewing topics going forward. Really feels like the future of education. A personal tutor in everyone's pocket.

English
3
4
72
7K
levi
levi@levidiamode·
188/365 of GPU Programming One of my favorite ways to learn at the moment is to just have a conversation with GPT-live like you would when talking to an expert. I've been using it today to review some of the concepts around Tensor Memory Accelerators and mbarriers today and it was able to go into quite a lot of detail with essentially instantaneous responses. Will probably use this as my default way of reviewing topics going forward. Really feels like the future of education. A personal tutor in everyone's pocket.
GIF
levi@levidiamode

187/365 of GPU Programming A couple of additional great resources on the topic of Tensor Memory Accelerators are Colfax's post on "Mastering the NVIDIA TMA" (which goes into TMA load, store, store reduce and load multicast) and @tqchenml's mini book "Modern GPU Programming For MLSys" (which includes a nice intro to TMAs and swizzling formats). Feel like I have much better understanding of TMA operations, tensor maps and mbarriers now. Learning more about TMEM feels like a natural next step here.

English
1
1
58
5.3K
levi
levi@levidiamode·
187/365 of GPU Programming A couple of additional great resources on the topic of Tensor Memory Accelerators are Colfax's post on "Mastering the NVIDIA TMA" (which goes into TMA load, store, store reduce and load multicast) and @tqchenml's mini book "Modern GPU Programming For MLSys" (which includes a nice intro to TMAs and swizzling formats). Feel like I have much better understanding of TMA operations, tensor maps and mbarriers now. Learning more about TMEM feels like a natural next step here.
levi tweet medialevi tweet media
levi@levidiamode

186/365 of GPU Programming It's time to learn more about Tensor Memory Accelerators (TMAs)! I've had a pretty shallow understanding of TMAs so far, hence want to invest more hours into comprehending their use/constraints/et cetera. Starting with this article "The Hitchhiker’s Guide to the Tensor Memory Accelerator" by MIT's 6.S894 course. The original version of 6.S894 was created by William Brandon who gave one of my favorite ever GPU Mode talks, so have high hopes for the resources from this course.

English
3
5
140
8.3K
levi
levi@levidiamode·
186/365 of GPU Programming It's time to learn more about Tensor Memory Accelerators (TMAs)! I've had a pretty shallow understanding of TMAs so far, hence want to invest more hours into comprehending their use/constraints/et cetera. Starting with this article "The Hitchhiker’s Guide to the Tensor Memory Accelerator" by MIT's 6.S894 course. The original version of 6.S894 was created by William Brandon who gave one of my favorite ever GPU Mode talks, so have high hopes for the resources from this course.
levi tweet media
levi@levidiamode

185/365 of GPU Programming One GPU kernel benchmark I keep coming across is SOLExecBench by Nvidia. Cursor, Recursive, Databricks, et cetera have all been testing their harnesses/agent systems to tackle the 235 blackwell kernels in this eval. Spending some time today to understand the paper/benchmark a bit better and since Nvidia is providing free compute to test your kernels, I'll take a closer look at some of the problems and see how close I can get to their SOL bounds.

English
2
10
192
16.5K

Keşfet

@dumbfook @denniszdes @VipulS_1 @NatKokoromyti @junnjiee_ @elonmusk @BarackObama @taylorswift13