Advanced GPU Assembly Programming Third Edition

Advanced GPU Assembly Programming

Most GPU performance problems are not source-code problems. They are machine-code problems.

A kernel can look clean in CUDA or HIP and still lose the war at the hardware level.

The compiler may choose an instruction sequence you did not expect. A branch may split a warp or wavefront into masked paths. A load pattern may explode into extra memory transactions. A tensor pipeline may sit underfed while the code looks “mathematically right.” Occupancy may look healthy while register pressure, wait states, barriers, cache behavior, or issue slots quietly cap throughput.

That is where this book begins.

Advanced GPU Assembly Programming is for advanced CUDA, HIP, AI-systems, HPC, and compiler engineers who need to read GPU machine code, understand NVIDIA and AMD execution behavior, and push kernels closer to the hardware performance ceiling.

This is not an introductory CUDA book.

It is not a beginner HIP guide.

It is not another surface-level explanation of “parallel programming on GPUs.”

This is a low-level technical reference for engineers who already understand kernels and now need to understand what those kernels become after compilation.

If you are optimizing AI inference, LLM kernels, GEMM, attention, scientific workloads, compiler output, CUDA-to-HIP portability, or architecture-specific performance, the question is no longer:

“Does the kernel run?”

The question is:

What is the machine actually doing, and how close is it to the real limit?

Inside, you will learn how to reason about:

• SIMT execution: warps, wavefronts, active masks, divergence, reconvergence, predication, and independent thread scheduling
• Machine code: PTX, SASS, AMD ISA, instruction encoding, disassembly, source correlation, and compiler idioms
• Execution resources: SMs, CUs, schedulers, issue slots, scoreboards, barriers, wait states, register files, and occupancy limits
• Memory behavior: coalescing, global memory, shared memory, LDS, cache policy, HBM bandwidth, alignment, sectors, and bank conflicts
• Tensor and matrix pipelines: tensor cores, MMA, MFMA, TMEM, FP8, FP6, FP4, block scaling, operand staging, and accumulator flow
• Asynchronous execution: cp.async, tensor-memory movement, producer/consumer roles, barriers, staged pipelines, and latency hiding
• Performance evidence: Nsight Compute, Nsight Systems, rocprof, Radeon GPU Profiler, Omniperf, roofline analysis, and microbenchmarking
• Real workloads: high-performance GEMM, attention, reductions, scans, sparse computation, atomics, scatter/gather, and irregular kernels

The value of this book is not that it tells you GPUs are fast.

You already know that.

The value is that it gives you the machinery to diagnose why a kernel is not fast enough.

Why did the compiler emit that instruction sequence?

Why did this memory access pattern create extra traffic?

Why are tensor units idle?

Why did a theoretically good tiling strategy lose throughput?

Why did NVIDIA and AMD behave differently?

Why did a change that looked harmless at source level move the bottleneck somewhere else?

This book helps you connect source code, compiler decisions, disassembly, profiler counters, memory transactions, lane masks, and architectural constraints into one coherent performance model.

Advanced GPU Assembly Programming was written for the engineer who wants the layer beneath CUDA, HIP, Triton, compiler output, and vendor libraries.