a blog about things i am interested in. you can contact me on linkedin or x....
Radix Top-K is an algorithm for finding the top-k elements in an array without sorting the full array. For simplicity, assume the values are unsigned integers. The same idea can be extended to other representations. Initial setup: Choose TOP_K and BITS_PER_ITER , and assume every element can be represented with NUM_BITS bits. Then iteratively apply the following procedure: Extract the next…
In CuTeDSL, it is often convenient to use StaticPersistentTileScheduler when writing persistent kernels. In this blogpost I give a brief intro to how it works. Usage in Kernel Typical usage within a CuTeDSL program is that we first set up the grid and the scheduler params in a dedicated function: tile_sched_params = utils . PersistentTileSchedulerParams ( problem_shape_ntile_mnl , self .…
SBO and LBO are essential to write correct Kernels on B200 GPUs that leverage Tensor Cores. That is because they need to be provided to the SMEM descriptor that the tcgen05 MMA op expects the programmer to provide as an operand. The corresponding section in the PTX docs commonly is perceived as confusing and in this blogpost I try to explain the mechanism for both K-Major and M-Major matrix…
In this short note we will briefly derive a helpful identity to speedup GDN prefill algorithm. In my simple Torch implementation of the algorithm that gave already a good speedup of about 18%. I assume for custom CUDA C++ kernel the gains can even be more pronounced. Please read my previous post on GDN for background. Reminder The state transition for GDN is as follows: S t = S t − 1 ( α t ( I − β…
Chunkwise Gated Delta Rule is important when performing operations such as prefilling or training with the recently popular Gated Delta Attention formulation. In essence, in the chunkwise algorithm, we want to transform the naive recurrent equations for the transition from one timestep to the next into a form that is more GPU-friendly. This works by splitting the sequence length into chunks and…
In this blogpost we will briefly explain Gated Delta Net decoding. This workload is one of the workloads to be optimised in the Flashinfer Competition and understanding its mathematical properties can help in optimising the corresponding kernel. I hope this helps to onboard others quickly with the necessary understanding of Gated Delta Net decoding. Reference Code and its interpretation import…
In this blogpost we will discuss the kernel of grouped blockscaled GEMM on B200. This rounds up our recent series on grouped blockscaled GEMM on B200 GPUs. In this blogpost we focus on analysing on the difference between a persistent grouped GEMM in CuTeDSL vs the ordinary persistent GEMM. kernel @cute . kernel def kernel ( self , tiled_mma : cute . TiledMma , tiled_mma_sfb : cute . TiledMma ,…
This is the next blogpost where I aim to explain grouped blockscaled GEMM for B200 GPUs in a top down manner. Before studying the internals of Grouped Blockscaled GEMM Kernel we should understand the setup of MMA and TMA, tile scheduler and other parts which are important during kernel execution. This blogs aims to explain the parts that are different from the usual persistent blockscaled dense…
This is the first blogpost on a series which will be about Grouped Blockscaled GEMM on Blackwell of which the code can be found here . In ordinary english that means we can calculate a group of GEMMS, i.e. multiple blockscaled GEMMS with potentially distinct problem sizes. Before we study the kernel itself it is useful exercise to understand the setup within the run function because we can learn…
Warp specialisation is an optimisation that splits the GEMM mainloop into two parts: We have one warp that does the TMA (i.e. copy tiles to SMEM) and one warp that does the MMA (i.e. multiply together these tiles). CuTeDSLs pipelining abstraction makes it particularly convenient for the user to implement this optimisation. In this short note I briefly show how to turn ordinary non persistent…