[jpgvm]

Yeah pretty much this.

I would separate the knowledge into maybe 3 distinct buckets.

The baseline: device/host boundary, SIMT programming etc.

The intermediate: kernel architecture, CUDA graph vs persistent kernels, warp specialisation/divergence avoidance techniques etc.

The advanced: architecture specifics so tcgen05, TMA, SMEM/HBM, memory throughput vs compute biases in various arch impls., GEMM, FHMA, all the tricks that make modern fused kernels very fast. Also would bucket most GPU Direct RDMA/GPU NetIO/friends here too.

The baseline hasn't changed much and probably won't, the intermediate knowledge has also remained pretty reliably stable for ~10 years with only things like graphs changing stuff. Tile might become more relevant than it is today but for now CUDA, cuBLAS, friends are where it's worth investing knowledge.

[fransje26]

> [..] all the tricks that make modern fused kernels very fast

This would require very different (re-written?) kernels than a few years back, wouldn't it?

Would you have any good resources on the topic?

[namibj]

There's actually little that changed in a way too fundamentally to matter other than _perhaps_ getting the async load-from-global-to-shared-memory DMA memcpy that avoided blocking register file space as target buffers for in-flight read-from-global operations. Shared after all is just a partition of L1d$ since iirc Volta (since they offered non-fixed/at-launch-requested expanded shared capacity support), so it made sense to provide this not-just-a-hint "prefetch into this user-managed slice of what is otherwise L1d$": it's AFAIK basically just some special load-like units that ask special L1d$-miss-fill units to deliver to a now-explicitly-specified target location in the non-automatic-cache partition of the local SRAM and signal completion in otherwise fairly normal local semaphore/barrier fashion.

The major difference is that this doesn't have a natural moment to transform/touch the values after read from global and before storage to shared.

Otherwise, tiled MMA (gemm) kernels where normal even in Maxwell days (after the classic K80, before the P100; Maxwell is when H.265 support landed).

[KeplerBoy]

I wish there were any good literature on GPU Direct RDMA and GPU NetIO. Got any tips?

[jpgvm]

So I would say the most important thing is that the APIs these are using as in mlx5 DevX (essentially direct fw access) or ibverbs are exactly the same regardless if it's CPU or GPU talking to it. So with that in mind the source of rdma-core, DPDK, ucx etc may be the most elucidating when it comes to low level details.

For higher level patterns again the APIs are the same so anything building on libibverbs or aforementioned ucx etc are pretty compatible from a high level ideas perspective. If you are new to RDMA in general definitely start with raw verbs instead of using abstractions like MPI if you really want to build a good intuition and then move to MPI once you understand what it is doing for you.

[KeplerBoy]

Thanks I appreciate the answer!