wafer✓

Vercel Developers @vercel_dev

DeepSeek V4.1 Flash 1017 is now included in Vercel Token Pass. • New weights from DeepSeek • 1M context window • No config change Get $80 of tokens for $30/mo ↓

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.

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.

wafer @wafer_ai

we launched the most comprehensive ai performance engineering repo in the world we're posting every single resource follow and save to keep up with the series. links in thread 🧵 part 14: Blackwell Tuning Guide NVIDIA covers the execution limits, memory hierarchy, and communication mechanisms behind tuning AI workloads on Blackwell: - register budgets, resident warps, and occupancy in fused AI kernels. - shared-memory limits for attention and matmul tiles across Blackwell compute capabilities. - thread-block clusters and distributed shared memory for exchanging tensor data across blocks. - 8-block portable clusters, B200's optional 16-block clusters, and the parallelism tradeoff of larger clusters. - L1/shared-memory carveouts and the balance between caching and explicit tensor staging. - L2 cache persistence for inference kernels that reuse the same tensors. - NVLink peer access and direct tensor transfers between GPUs in model-parallel inference. the guide connects Blackwell's hardware limits to kernel design decisions like tile size, block cooperation, and tensor reuse, helping engineers choose what to test when tuning attention, matmul, and multi-GPU inference.

wafer @wafer_ai

we launched the most comprehensive ai performance engineering repo in the world we're posting every single resource follow and save to keep up with the series. links in thread 🧵 part 14: Blackwell Tuning Guide NVIDIA covers the execution limits, memory hierarchy, and communication mechanisms behind tuning AI workloads on Blackwell: - register budgets, resident warps, and occupancy in fused AI kernels. - shared-memory limits for attention and matmul tiles across Blackwell compute capabilities. - thread-block clusters and distributed shared memory for exchanging tensor data across blocks. - 8-block portable clusters, B200's optional 16-block clusters, and the parallelism tradeoff of larger clusters. - L1/shared-memory carveouts and the balance between caching and explicit tensor staging. - L2 cache persistence for inference kernels that reuse the same tensors. - NVLink peer access and direct tensor transfers between GPUs in model-parallel inference. the guide connects Blackwell's hardware limits to kernel design decisions like tile size, block cooperation, and tensor reuse, helping engineers choose what to test when tuning attention, matmul, and multi-GPU inference.

wafer @wafer_ai

we launched the most comprehensive ai performance engineering repo in the world follow and save to keep up with the series. links in thread 🧵 part 12: CUDA C++ Best Practices Guide NVIDIA's guide helps AI performance engineers connect a slow training step or inference request to a change they can test. it's a reference for deciding whether to change a tensor layout, reduce data movement, or tune a kernel's resource use. NVIDIA covers the measurement and optimization decisions behind CUDA applications: - profiling representative workloads and using Amdahl's law to estimate the benefit of accelerating a hotspot. for inference, use relevant batch sizes and sequence lengths to decide which kernels deserve engineering time. - timing GPU work with CUDA events and comparing requested memory traffic with actual transfers. for tensor copies and layout conversions, these measurements help check how much bandwidth the access pattern wastes. - keeping intermediate tensors on the GPU between model operations to avoid CPU round trips. pinned host memory and separate non-default streams can overlap incoming transfers with independent computation when the hardware supports it. - coalescing global-memory accesses by matching thread assignment to tensor layout. for strided tensors and transposes, inspect which addresses neighboring lanes touch before choosing a new layout or writing a custom kernel. - reusing matrix tiles in shared memory to reduce repeated global loads. the guide's transposed-tile example shows how padding removes its bank conflicts, giving engineers a concrete layout change to evaluate inside tensor kernels. - balancing block size, registers, and shared memory when tuning custom AI kernels. more registers can reduce spilling while lowering occupancy, so measure runtime across configurations to decide whether that tradeoff helps. - checking numerical results after changing reduction order or using faster math functions. for custom normalization or softmax kernels, compare outputs against a reference with explicit error tolerances alongside performance measurements. the worked matrix multiplications connect memory access patterns to measured performance. use them to plan a kernel optimization, then check whether the change improves the full training or inference workload while preserving numerical accuracy.