Hello X ...
Introduction I recently built a small (mainly pedagogical) distributed training engine to understand how different forms of model parallelism work in practice. This post compares sequence parallelism as used in Megatron-LM and DeepSpeed. Although both systems keep the transformer residual stream partitioned as [B, S/P, H] (where P is group size), the mechanisms they use around attention are quite…
Introduction What is being asked in these experiments is if some future information, in the form of a rollout being successful or unsuccessful, is "leaked" (in the form of entropy over time) somewhere early enough in a completion that by intervening mid-generation or otherwise you can induce likely-to-fail rollouts into becoming successful. Background The experiment was done using TRL's AsyncGRPO…
Intro This post is about the bubbles in you see in your Nsys traces and a couple interesting methods (tvm-ffi and Nvidia's Programmatic Dependent Launch) that aim to close those gaps. TVM-FFI The Nvidia docs strongly suggest compiling your CuTe kernel with tvm-ffi because of its ability to reduce launch overhead. The couple images below show an nsys trace with (top) and without (bottom) tvm-ffi.…
Introduction The core function of an epilogue is to write data to gmem. The accumulator values sit in registers in the layout chosen by wgmma, but the final output needs to land in global memory as a normal (M, N) tensor, usually in a different dtype. In this post, we’ll look at the machinery that bridges these worlds, including retiled accumulator fragments, stmatrix stores into shared memory,…
Intro In a recent blog , in part we spoke about how a given tile scheduler might affect the speed of your kernel. But throughout that blog, we never spoke to the constraints that exist upon the A or B tiles themselves or their roles in the actual matrix-multiply portion of a gemm kernel. The aim of this blog is to speak more to various ways we add "depth" to a gemm (CuTeDSL) kernel. Hardware…
Motivating the issue A matrix multiply is embarrassingly parallel. Each output tile can be computed independently. There are no data dependencies, no synchronization, and no ordering constraints. This is a kernel that can run at about 1.05x the speed of CuBLAS at M=N=K=8192 . Yet, when we naively order our tiles, don't use a persistent kernel, and don't use TMA multicast, our performance…
Motivating the issue A lot of the blogs that I've read on swizzling don't motivate the problem in a way that I can understand. So my aim for this blog is to do just that. Let me know what you think. The crossbar switch is the physical bottleneck connecting processing cores to memory banks. The crossbar is the underlying technology, which handles data access routing, and explains why this problem…