san francisco, ca
Joined May 2024
Update on my progress with Inference Engineering: Since the last update, I have been moving across states, meeting family members and what not, while trying to do the minimum amount of work to still make meaningful progress. But to be completely honest, the biggest bottleneck right now is fighting pangram since my writing is constantly being flagged as AI generated slop. I've been going back and forth with pangram for the sake of credibility and providing value to others without being perceived as someone who produces AI slop. I think I've completely accepted the fact that my writing will always be flagged as partially AI. I’ve concluded that the only metric that matters is if the writing actually makes sense. Anyways, I’ll try to finish the whole article soon. Thank you to all my friends for reaching out. 💪 I’ll be back.
7
12
203
11,360
New post: what happens when a GPU reads memory We follow an LDG instruction through the memory hierarchy on an RTX 4090, from the warps down to the DRAM and back, reverse-engineering the undocumented parts along the way. blog.doubleword.ai/what-happ…
1
50
1
328
20,766
this is probably my favorite article ive written. in this blog, i try to profile llm inference served using sglang and reason about the patterns, kernels and bottlenecks you would usually find in production.
22
100
6
920
104,305
If you are looking for a guide for learning cuda matmuls, skim through these and pick which one you think will engage your attention and curiousity the best. Learning this WILL pay off.
6
8
164
21,559
Day 28/90 of Inference Engineering I learned how to optimize an SGEMM with 1D block tiling in CUDA. Progressing from shared memory block tiling, 1D block tiling increases arithmetic intensity by increasing the amount of computation performed per byte. In my previous post, we implemented shared memory tiling, which loads data into shared memory so we can avoid redundant loads from global memory. Now, 1D tiling operates one level lower in the memory hierarchy by addressing redundant shared memory traffic. With 1D tiling, each thread loads a value from the shared memory tile into a register once, then reuses that register across multiple output computations instead of repeatedly reading the same value from shared memory. Practically speaking, every cell along row 0 of the output matrix tile needs the same row 0 values from A. Without 1D tiling, multiple threads repeatedly read those same A values from shared memory. With 1D tiling, a thread loads an A value into a register once, then reuses it across multiple output computations while loading different B values from shared memory. Implementing 1D tiling will decrease occupancy because each thread now uses more registers. Since the register file is finite, fewer threads, warps, and blocks can be resident on an SM simultaneously. But a decrease in occupancy is okay, it's a tradeoff when implementing this optimization. We're sacrificing some occupancy to increase computation per thread, increase register reuse, and reducing shared memory traffic.
8
26
250
10,356
Introducing SparklingTree: a mashup of DSpark (first-order markov corrector) and DDTree (tree diffusion drafting) that achieves 30-40% speedups over DSpark! I wrote every experiment, everything I learned, and the results in a blog here! jwlabs.vercel.app/post/spark…
13
3
66
3,461
Day 27/90 of Inference Engineering Optimizing a SGEMM CUDA kernel with SMEM tiling. Here’s my biggest takeaway on the difference between the two optimizations and how tiling is implemented. Global memory coalescing optimizes how threads within a warp access global memory so their requests can be combined into as few memory transactions as possible. Shared memory tiling reduces global memory traffic by loading tiles of matrices A and B into shared memory, allowing each loaded element to be reused by multiple threads. This reduces redundant global memory loads by having each tile loaded once and reused many times from shared memory. By dividing the output matrix into sub-matrices of TxT tiles, each threadblock loads the required A and B tiles into shared memory. Multiple threads then reuse the shared memory tiles to compute partial sums without repeatedly fetching the same data from global memory. I hope to fully implement this and move on to 1D block tiling within a few days.
6
12
161
6,865
Day 26/90 of Inference Engineering I didn't understand global memory coalescing until I took a step back and looked at coalescing at the warp level in parallel execution. It's not hard to understand at the surface level; it's a technique that ensures the 32 threads of a warp access adjacent addresses, so all 32 requests merge into a single transaction. Coalescing a matrix multiplication immediately helps with the output matrix C, since a warp no longer writes to a new row for every entry. Instead it writes 32 adjacent cells along one row. But coalescing also helps with matrices A and B: A's address is now identical across the warp, so it broadcasts in a single transaction, and B is coalesced from the warp's perspective too. Memory access is more efficient because the 32 threads of a warp now sit next to each other in memory at every step, instead of being scattered a full row/stride apart. Below is an animation showing how memory is accessed when it's coalesced versus strided. --- A little life update: After visiting SF, I visited some family. My mini PC parts also came so it took me 2 days to build everything. I didn’t have much time to really tinker with kernels and read up on inference related papers, but the momentum is slowly picking back up now that I’ve fully settled in. While I was in SF, I did finish reading PMPP up to chapters 6, I just have to work through all of the problem sets now. As for what I’ve accomplished today: 1. Fix visualkernels.com with an actual tree-sitter 2. Fixed the UI so matmul is properly laid out 3. Figured out global memory coalescing 4. Learned how to connect a custom CUDA kernel to pytorch 5. Tinkered with my vimrc
3
9
201
7,285
Day 24/90 of Inference Engineering Building a mini PC…
4
44
1,918
so tell me why i need to an extension cable for my GPU
2
551
Day 21/90 of Inference Engineering Over the last two days I’ve met a lot of truly exceptional and inspiring people in SF. Thank you for having me. - Updated on SGEMM optimization: hashed out global memory coalescing, now piecing together tiling.
1
1
73
3,218
Day 20/90 of Inference Engineering met a lot of cool people!
3
2
119
4,059
Day 19/90 of Inference Engineering Working on an optimized SGEMM with a cuBLAS like performance. Not much to say! Life is good. Also watched this banger video: youtube.com/watch?v=QQceTDjA…
2
2
101
6,167
Day 18/90 of Inference Engineering Bashing my head against the concept of global memory coalescing and tiling and its relationship with memory.
9
8
172
5,377
Day 17/90 of Inference Engineering Happy to say that I finally finished writing MNIST in CUDA. Here's an article I wrote on the little detail behind why cublasSgemm() works the way it does and why it's important. Thank you for following my journey thus far! I'm excited to write more CUDA and start profiling with Nsight~
3
5
54
2,786