Meanwhile, vision folks

Looking at RoPE. Why does everyone do RoPE in BF16 for LLMs 👁️

Generate videos in just a few seconds. Try Grok Imagine, free for a limited time.

Remind me of the Importance Sampling stuff when I did Peter Shirley's raytracing series.

I was trying to add Wan2.2 to my repo just ytd.

Kinda hate that autocast is sprinkled generously across the codebase. Reasons to NOT use autocast.- Dtype casting rules can be unexpected.- Different casting behavior from FSDP's.- Non-trivial overhead.Would be better to explicitly do .float() / .bfloat16() instead

just found out Wan2.2 was released in FP32 👁️. Was wondering why I got OOM 🤣

Let's see. My kernel is only off by 1e36. Doesn't look too shabby

Prototyped an attention with MXFP8 for QK.T.Good news: it's faster than BF16 attention.Bad news: MXFP8 matmul becomes slower in e2e 🤡.Have a feeling it's related to power limit / throttling

SageAttention folks are ingenious. Per-thread quantization scaling based on MMA accumulator thread layout to utilize INT4 Tensor Cores🤯.Figure taken from SageAttention2 paper -

How to beat Cutlass (on 5090 FP8)? Use a faster instruction 🤣. Lolz aside, >100% SOL is ofc not possible. With the faster instruction, SOL is doubled, so % SOL is actually only ~60%

This is funny since the more complicated instruction is apparently faster. Follow up q. Why can't PTX->SASS compiler emit the faster SASS automatically when the slow PTX is used?.

I can confirm that mma.sync.aligned.m16n8k32.row.col.kind::mxf8f6f4.block_scale.f32.e4m3.e4m3.f32.ue8m0 is faster than mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 on 5090. mma.sync.aligned.m16n8k32.row.col.kind::mxf8f6f4.f32.e4m3.e4m3.f32 has the same speed as old one.

I should be reading SageAttention papers. .

In attention, the 1st matmul Q @ K.T does reduction over head_dim, which is typically small e.g. 128. This seems ideal of low-bit matmul? I think even int8 with 1 scale per head_dim would work fine?.The 2nd matmul P @ V is a different beast, due to reduction along sequence dim. .

Managed to get mma.block_scale working. The good thing is, the scale layout is more sane than I thought initially. The bad thing is, it means the diagram is wrong. Or it's just very confusing. Not sure where is the lower/higher address/index. NVIDIA folks perhaps can clarify.

Was quite surprised that Qwen3 impl in HF transformers looks very clean. And some layers use kernels from HF kernels hub 👀

Have been staring at this for the last few days 😵‍💫

CuDNN seems to be only using sm80 features though. So maybe with TMA, it can be even faster.

Faster than FA2 on 5090 👀.Hopefully I can get to CuDNN speed

SageAttention3 has NVFP4 attention for 5090 👀.Doesn't look like code has been released yet though.