Гайд по GPU-программированию. Большой глубины нет, но для вката норм.

База есть: устройство SM и warp'ов, coalescing и bank conflicts, intra/inter warp примитивы, shared-memory и register tiling, occupancy, редукции, сканы и базовый профайлинг. Хотя скорее обзорно.

Инфы для уважаемых господ нет: cp.async/TMA, пайплайны через mbarrier, ldmatrix/mma, WGMMA и tcgen05, warp specialization с producer/consumer и прочими прелестями, thread-block clusters и DSM, persistent kernels, разбор PTX/SASS.

Всякое задротство даже вскользь не упомянуто: реверс latency/throughput и стадий execution pipeline для SASS, issue/dispatch и scoreboard-зависимостей, dependency barriers и operand reuse flags, укладка GMMA/TMEM/SMEM дескрипторов и swizzle'ов, register-bank conflicts, spills и occupancy cliffs.

Возможно, когда-нибудь я соберу все известные мне шаманские мантры в одном месте.

Если этот пост наберет 30 масюнь, сделаю мини-книжку по шаманству 😀: про то, как read-only loads оказываются на L1/TEX path, как микробенчами реверсить register banks и operand collectors, как вручную перекладывать dependency barriers и reuse flags в SASS, как все это автоматизируют компиляторы для гетерогенных ml железок и где ручные кернелы критичны будут всегда.

https://arpanpathak.github.io/gpu-parallel-book/

@learning_stones