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.
max fu retweeted
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…
max fu retweeted
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.
max fu retweeted
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.
Replying to @elliotarledge
yep, here are the most helpful resources i’ve used so far
1. siboehm.com/articles/22/CUDA… (@Si_Boehm)
2. abhik.ai/articles/cuda-matri… (@abhiksark)
3. robertzhang.me/blog/cuda-mmm (@robdobflob)
and PMPP v5
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.
max fu retweeted
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…
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.
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
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…
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~


















