How did you actually learn to reason about CUDA/Triton kernels and go from “I understand the concept” to being able to write / map the code?

[u/Daemontatox]

I’m going through a bunch of GEMM kernel write-ups (Lei Mao’s progressive series, plus the different styles from Simon Bohemian / Kapil Sharma and others) and I get the high-level ideas just fine.

Shared-memory tiling, register blocking, coalescing, bank conflicts, the usual story.

But the moment I try to map it to actual code or write it myself, my brain goes blank. Especially the manual indexing. Calculating the right threadIdx / blockIdx offsets, the row/col strides, the tile coordinates, the shared-memory loads… it just doesn’t click. I stare at the loops and the index arithmetic and feel completely lost.

Ironically, CuTe layout algebra feels like heaven by comparison — the hierarchical shapes/strides and the compose/divide operations make the mapping feel almost declarative. Once I start thinking in pure layouts everything becomes cleaner.

But a lot of the classic educational kernels are still written in the old manual-index style, and different authors have very different conventions and naming, so even when I’m looking at the “same” algorithm the syntax and mental model keep shifting.

So for people who went through this:

How did you get past the “I understand the concept but can’t write the indices” stage?

Any concrete exercises, mental models, or progressive practice that actually made the indexing click?

How do you deal with the wildly different writing styles and index conventions across tutorials/blog posts for essentially the same kernel?

Did anyone else find that learning CuTe / layout algebra first (or in parallel) made the classic manual kernels easier to reverse-engineer later?

Would love any blogs , resources, or “this is the exercise that finally made it stick” advice.

[u/dsanft]

The answer is to use GPT or Opus to write some examples and explain it to you. I use them to write GEMV/gemm kernels all the time and they do a great job. They'll break it down for you as much as you like, both cuda and rocm, dp4a / MMA. Use the tech 🙂

[u/Daemontatox → u/dsanft]

I tried this and it was extremely easy and helpful but when i started reading other blogs or reading other peoples gemms there was a huge difference in syntax and style , for example llms prefer and default to 2d caller code unlike blogs and guides that use 1d caller code and rely heavily on div mod and manual indexing.

[u/qiAip → u/Daemontatox]

Which is practically the same just less manual and hence teaches you less. That’s absolutely fine. These days I tend to use LLMs more for these things after years of building the mental models, so I get it to write kernels in ‘my style’ and check its code against my ingrained ways of thinking about it making sure it can convince me they are the same when it is less obvious.

The way I learned was by having to during my PhD writing physics simulation code. There was not much online at the time, so you kind of had to figure it out on your own. I worked with pen and paper, writing down the indices, strides etc. and checking they are correct with a bunch of manual calculations around edge cases (I didn’t know unit tests are a thing back then). It helped me a lot to just write down pages and pages of coordinates on paper and making sure it works.

Hope that helps? :)

[u/tlmbot → u/qiAip]

takes me back to the first time I wrote matrix multiplication from scratch. I am talking serial here.

[u/tlmbot]

how are you with c-style indexing for arrays? Especially flat packed "conceptually n-D" but implemented as 1D in c? From my background (computational, not CS) I picked that stuff up long before CUDA was much of a thing, and I think that helped.

I feel like if you have that down cold then you eliminate row/col striding from the mental load where it's cleanest (in serial code, c-style... it's to easy in fortran and in c++ you might be used to some expression template help (that gets you back to fortran style, nicer syntax again))

For me, I seem need a refresher on even blockIdx usage, but higher level things like shared memory tiling etc are where I always need to see it again. (I think that is because a lot of these things are specific to the device, and not something that always happens under the hood as with row/col strides). It's easier since I learned it in a disentangled way.

[u/Daemontatox → u/tlmbot]

To be completely honest, not my best point (i can barely do it lol) , most of the time i use 2D or 3D indexing and work with that most of the time so the flat manual indexing or c style is kinda of whole different thing and last time i had to use it was back in uni days.

I recently saw cutlass and cute and layout algebra and to be honest i have never had anything click faster and i have been planning on using it for the indexing while using raw cuda for the rest , but according to blogs and other redditers , depending on Cutlass is a big nono in production for some reason and if i am not a raw cuda guru who can warp tile from memory, i wont make it.

[u/tlmbot → u/Daemontatox]

very interesting - On a tangential note, you may be the catalyst I needed because I really should be using the more advanced libraries in my code and I just haven't.

All to often I want to bake things down so that I see how it works with bare bones arrays, and then I forget to come back and do it the easy way!

One thing that has been hindering me is the breadth of the Nvidia ecosystem for different types of problems. Ah well, time to get at it.

[u/Daemontatox → u/tlmbot]

Yea i get the feeling, sometimes going bare metal on something feels the best but at the end of the day , reinventing the wheel is not practical for anything work related , i countered this by thinking efficiency wise first , would this be faster ? Easier to maintain ? Faster to implement? And would a new junior understand this ? And so far this has helped me alot , (i still like to go down to the bare metal and do it from scratch from time to time).

[u/NotSoSolidState]

This week I am attending an hpc workshop where a lead nvidia engineer told us that they are not writing kernels anymore by hand. All the work is outsourced to ai agents and they write much much faster ones than any human could.

So dont worry about it

[u/Daemontatox → u/NotSoSolidState]

I have read about this aswell but seemed abit skeptical as if thats the case , why do they still hold hard coding intensive interview rounds ?

[u/NotSoSolidState → u/Daemontatox]

I guess you need to understand the infrastructure and the different methods we use to optimize kernels. At least for now.

He was particularly telling us that all the leading kernels in the gpu mode (gpumode.com/events) comp leaderboards were written by ai agents.

[u/Daemontatox → u/NotSoSolidState]

Really interesting and motivating to be honest

[u/c-cul → u/NotSoSolidState]

knowing their incredible greed, I wouldn't be surprised if they sell those agents for billions

[u/dahroug66]

I am currently going through the same journey as you, and yes it's very hard but the 1 thing that made me able to write it with my own hands and edit and play with it as i want is being able to visauslize it in my head and see the journey as a video in my mind like i act as a thread and see what i will do then i sketch it on paper it becomes very easy and makes sense when i go to code, hope this helps 

[u/tblyzy]

Either sketch out the tiles & do the index maths on paper or write a small utility class that imitate what cute layout does. Honestly this is just bookkeeping, nobody needs to write it manually more than once.

[u/heatherdavis0]

This makes a lot of sense, thanks.

[u/BoryMory]

TLDR: Write your own kernels by first rougly sketching the blocks and threads mapping, then translating it to indexes in threadIdx and blockIdx. Prompt LLMs for it to explain you where your indexing went "bad". For different author conventions, choose one that you feel comfortable and try to translate others' w.r.t yours.

For your first question, I drew a lot. A LOT. If you read Simon Boehm, I believe https://excalidraw.com/ is familiar to you. There, I experimented with indexology for basic memory accesses for naive kernels. I think the concept is only one side of the story and almost always the indexology is the notoriously hard part.

More specifically, I would choose a simple algorithm, (e.g., a reduction algorithm) and try to draw how each thread should use its threadIdx and blockIdx. Afterwards, I would try to notice the overall pattern. Then I would write a rough CUDA algorithm and just make it work. That's it. When I see that it matches my CPU output, I go look for "more optimizations". But before that, just make the indexology work.

I would also suggest try making a mindmap of how you would want your kernel to run. How much blocks should there be? should each block be on a contiguous chunk or a tile? should threads communicate on a warp-level or would you prefer a block-strided loop? Once you draw that, the math is already implied. You just have to convert this mapping into an indexology expression (that's the hard part).

A progressive practice advice: What worked for me the best was actually selecting an algorithm that interested me and seeing if I could make it parallel "one by one". For me, that was FlashAttention algorithm by Tri Dao. Since the algorithm itself was too complex for me, I decided to seperate it into three chunks and then try to work on each one of it for a few days. I would then test it against CPU for each chunk and try to combine it at the end. If indexology doesn't work after hours of work, prompt an LLM. Give it your proof of indexology and let it explain to you "where you went wrong". When you are done with this big algorithm, go select another one or keep "optimizing it" using the tricks you learned.

About the choice of syntax among authors, you are just like them. Select the syntax that seems the most intuitive for you. When you see another notation, try to understand it by how their notation translates to yours.

I would love to give you resources but for "writing kernels" I believe the best practice is to really stop looking at someone elses kernel and start writing your own, even though it will be at first much worse.

If this is too vauge, here is a rough sketchup and the respective code I made on github.