John Hahn✓

steve @gpusteve

if i had to learn cuda compilation, ptx, register spills, and gpu disassembly from scratch because my "optimized" kernel got slower, i'd start here links in thread 🧵 the compilation and machine code section of our ai performance engineering repo contains four nvidia resources for understanding how cuda source becomes gpu instructions and inspecting the code behind ai kernels: understanding ptx explains how nvcc separates cpu and gpu code, why ptx is a virtual instruction set, and how ptxas turns it into a cubin. the article covers fatbins, driver jit compilation, and gpu compatibility, giving context for deployment targets and inference startup costs. nvcc compiler driver covers gpu architecture targets, the files each compilation stage produces, separate compilation, and device linking. the nvcc guide explains how to retain intermediate files and report register usage, shared memory, and spill loads and stores. those compiler reports help compare kernel builds after changes to tile size, unrolling, or fusion. ptx isa defines register, local, shared, and global state spaces; memory ordering and synchronization; and arithmetic and matrix instructions. ptx isa details operand layouts, data types, accumulation precision, and gpu requirements, plus completion rules for asynchronous copies. useful for checking data movement and numerical behavior in attention and matmul kernels. cuda binary utilities covers extracting embedded ptx and cubins, disassembling native gpu instructions (sass), and reporting register counts and static memory usage with cuobjdump. nvdisasm provides control-flow graphs when the cubin contains complete relocation information, plus register-liveness views. with debug line information, nvdisasm can annotate instructions with source lines. useful for inspecting generated loads, stores, branches, and arithmetic after a kernel rewrite. together, these resources connect the compilation process to the code being measured. they help identify what changed between kernel builds and what to investigate before the next correctness check or benchmark. follow and save to keep up with the series.

steve @gpusteve

if i had to learn cuda compilation, ptx, register spills, and gpu disassembly from scratch because my "optimized" kernel got slower, i'd start here links in thread 🧵 the compilation and machine code section of our ai performance engineering repo contains four nvidia resources for understanding how cuda source becomes gpu instructions and inspecting the code behind ai kernels: understanding ptx explains how nvcc separates cpu and gpu code, why ptx is a virtual instruction set, and how ptxas turns it into a cubin. the article covers fatbins, driver jit compilation, and gpu compatibility, giving context for deployment targets and inference startup costs. nvcc compiler driver covers gpu architecture targets, the files each compilation stage produces, separate compilation, and device linking. the nvcc guide explains how to retain intermediate files and report register usage, shared memory, and spill loads and stores. those compiler reports help compare kernel builds after changes to tile size, unrolling, or fusion. ptx isa defines register, local, shared, and global state spaces; memory ordering and synchronization; and arithmetic and matrix instructions. ptx isa details operand layouts, data types, accumulation precision, and gpu requirements, plus completion rules for asynchronous copies. useful for checking data movement and numerical behavior in attention and matmul kernels. cuda binary utilities covers extracting embedded ptx and cubins, disassembling native gpu instructions (sass), and reporting register counts and static memory usage with cuobjdump. nvdisasm provides control-flow graphs when the cubin contains complete relocation information, plus register-liveness views. with debug line information, nvdisasm can annotate instructions with source lines. useful for inspecting generated loads, stores, branches, and arithmetic after a kernel rewrite. together, these resources connect the compilation process to the code being measured. they help identify what changed between kernel builds and what to investigate before the next correctness check or benchmark. follow and save to keep up with the series.

wafer @wafer_ai

1 trillion tokens a day on @OpenRouter!!

wafer @wafer_ai

1 trillion tokens a day on @OpenRouter!!

You.com @youdotcom

Introducing Inference Router: enterprise-grade inference on demand. Kimi K3, GLM 5.3, DeepSeek V4.1 Flash, Qwen3.8 Flash, Muse Spark 1.3, Mercury 3 and more, available with an uptime SLA of at least 99.9%. US hosting and web search built in.

You.com @youdotcom

Introducing Inference Router: enterprise-grade inference on demand. Kimi K3, GLM 5.3, DeepSeek V4.1 Flash, Qwen3.8 Flash, Muse Spark 1.3, Mercury 3 and more, available with an uptime SLA of at least 99.9%. US hosting and web search built in.

You.com @youdotcom

Introducing Inference Router: enterprise-grade inference on demand. Kimi K3, GLM 5.3, DeepSeek V4.1 Flash, Qwen3.8 Flash, Muse Spark 1.3, Mercury 3 and more, available with an uptime SLA of at least 99.9%. US hosting and web search built in.

wafer @wafer_ai

we're doing a massive update on Wafer's popular open source guide to GPU performance engineering! a few cool resources going into this update: ⠀ 1. Stream-K: why a matmul can leave part of the GPU idle while the last blocks finish, and how dividing the multiply-accumulate work more evenly helps. ⠀ 2. tcgen05 for dummies: gau-nernst builds a Blackwell matmul step by step. Tensor Memory, asynchronous copies and pipelining, with code you can follow as the kernel develops. ⠀ 3. Thinking Machines' Defeating Nondeterminism: how changing the batch can change the order of floating-point additions, and why that can change inference results. then, how to make the kernels batch-invariant. ⠀ 4. Triton's Linear Layouts: how tensor elements are arranged across threads and registers, and the data movement needed to change that arrangement. ⠀ the update also expands the quantization, speculative decoding and hardware references. ⠀ the guide starts with one inference request, then moves through GPU fundamentals, kernels, serving engines and distributed systems. start at the beginning if you're new; use the individual sections when you hit a specific problem. ⠀ if you know a primary source that explains one of these topics better, open a PR.

wafer @wafer_ai

we're doing a massive update on Wafer's popular open source guide to GPU performance engineering! a few cool resources going into this update: ⠀ 1. Stream-K: why a matmul can leave part of the GPU idle while the last blocks finish, and how dividing the multiply-accumulate work more evenly helps. ⠀ 2. tcgen05 for dummies: gau-nernst builds a Blackwell matmul step by step. Tensor Memory, asynchronous copies and pipelining, with code you can follow as the kernel develops. ⠀ 3. Thinking Machines' Defeating Nondeterminism: how changing the batch can change the order of floating-point additions, and why that can change inference results. then, how to make the kernels batch-invariant. ⠀ 4. Triton's Linear Layouts: how tensor elements are arranged across threads and registers, and the data movement needed to change that arrangement. ⠀ the update also expands the quantization, speculative decoding and hardware references. ⠀ the guide starts with one inference request, then moves through GPU fundamentals, kernels, serving engines and distributed systems. start at the beginning if you're new; use the individual sections when you hit a specific problem. ⠀ if you know a primary source that explains one of these topics better, open a PR.