Pranjal

I implemented H100 cuda matmul kernel from scratch, taking inspiration from @Si_Boehm's blog. Our final kernel outperforms cuBLAS by 7% for N=4096. It fits in a single C++ file without any dependencies. Full-blown blog post with all details: cudaforfun.substack.com/p/outperformin…

@PatrickToulme Does it ever read flashattention4 implementation or write from scratch

PyPTX is now the fastest FlashAttention kernel I’ve benchmarked. Using an improved PyPTX harness, GPT 5.6 Sol Max fully generated the kernel over a multi day loop that beats both FlashAttention-4 and cuDNN. Excited to share more results in the coming weeks. The future of kernel authoring is 100% agentic: models writing directly in the ISA.

@VictorTaelin You have done it once, just do it again

it is ok, I'll do better

... honestly, I'm too naive to be a founder. People are raising serious money (money I never raised) to create startups based on algorithms *I invented*, and open-sourced. That's going to be an expensive lesson. Taking a morning to appreciate people are just evil and it is all a game of interests, after all

@ellev3n11 composer chungus wen

le cursor humongous

@_arohan_ Agreed, but gaps will always remain. Let people tokenmaxx and you can build a training data pipeline from this and in the limit tokenmaxxing -> 0

Tokenmaxxing is what you should be doing during training. Inference should just one shot solutions

@tenderizzation GPUs can run dynamic shapes pretty well

one can dream

Glad to see they keep my skillset artificially relevant for a while longer.

@elliotarledge Idea is to fix more bugs than it creates

Who's working on LLMs making puzzles by themselves?

@real_deep_ml Claude can already do this

We just launched a new project that teaches you how to build Flash Attention with CUDA, step by step. By the end, you’ll have a working Flash Attention kernel built from the ground up. The project covers: -CUDA primitives warm-up -Matrix operations -Naive attention baseline -Online softmax math -Tiled attention building blocks -Fused Flash Attention kernel -Causal Flash Attention It will be open to everyone for the first 2 weeks, then it will become part of our premium projects.

Hard times create strong engs Strong engs create compilers Compilers create weak engs Weak engs create hard times

@ellev3n11 Use subagents aggressively and do long running tasks

what do people want to see in composer-next? any new particular capabilities?

@maharshii Once you go ptx, there's no way back.

1. The only place where CuTe layouts help is for TMA copying row-major quantized A and B from GMEM to SMEM with 128B swizzling 2. For scale factors, mbarriers and tensor memory, you can just allocate them as per-stage buffers using the Smem/Tmem allocator util that CuTeDSL provides. 3. There is warp specialization such that 1 warp handles TMA from gmem to smem, 1 warp just issues MMAs, and 4 warps do the epilogue. 4. For loading tensor memory to registers, tcgen05.ld requires 4 warps (a warpgroup) and thus epilogue needs to have 4 warps atleast. 5. CuTeDSL's nvvm mlir dialect provides the functions for tcgen05 ops that can be easily integrated as dsl user op. 6. For ops not in nvvm dialect like cp async bulk 1D gmem to smem (used for scale factors), you can just write the llvm inline asm (PTX) and expose it as a dsl user op. 7. Tcgen05 ops take descriptors i.e. shared memory or instruction descriptors which are just int64 with the information about MMA/Cp packed into them. 8. We move both scale factors to tensor memory first before staring with MMA, that's what cutlass does too. 9. Tcgen05 MMAs have an implicit unit of 8x2 (32B) tiles so in the K dimension, 1 MMA can take 64 FP4s, 16 BF16s, and so on... which happens to be the MMA_K for these dtypes! 10. GlobalScale is only loaded by the epilogue warps and multiplied with the accum tensor before casting it back to output dtype. Here's the code: github.com/Maharshi-Pandy…

I wrote a custom NVFP4 GEMM kernel in CuTeDSL stripping away almost all the fancy CuTe layouts "headache" in the official examples and doing the PTX, TMA, and Tcgen05 manually. It's crazy how low-level you can go with this and still be performant! My notes and code are below:

@HeinrichKuttler @pmddomingos next blog?

@pmddomingos this appears to be a misunderstanding of the relevant physics

All our measurements of space and time are quantized, and yet we still believe they're continuous.

@HeinrichKuttler Daddy Heiner teaching us how this works

How this actually works: Hazard functions. heiner.ai/blog/2021/03/2…

@HeinrichKuttler Ok I'll bite

@pranjalssh yeah. now ask an llm

Finally doing my CA driver license. Guess: When is it permitted to drive across double solid yellow lines? A When making a U-turn or left turn. B When the road is clear to pass the vehicle in front of you. C When instructed by construction or other signs due to a road closure.

This is a very common misconception. Biggest wins are overlapping. If you have 2 memory bound kernels back to back reading and writing 1GB, you pay 4GB of HBM cost. If you overlap them while streaming intermediate state of 128MB(fits in L2 cache) - you only pay 2GB HBM cost. For communication overlapping, benefits are even higher.