cutlass

I'm telling your it works now. It's just not called `tcgen05`. Put this in nsight compute: https://github.com/NVIDIA/cutlass/blob/main/examples/79_blac... (I said 83, it's 79). If you want to know what NVIDIA really thinks, watch this repo: https://github.com/nVIDIA/fuser. The Polyhedral Wizards at play. All the big not-quite-Fields players are splashing around there. I'm doing lean4 proofs of a bunch of their stuff. https://v0-straylight-papers-touchups.vercel.app It works now. It's just not the PTX mnemonic that you want to see.

Seems this is likely due to ongoing work on FP8 support on nvidia/cutlass. From my reading, the alternative code path was likely added recently for testing by external contributors to the cutlass project, and other involved parties. (Rather than attempting to distribute custom packaged internal builds of cuda.) This ticket is a good starting place to see the chain of issues around the ongoing work: https://github.com/NVIDIA/cutlass/pull/2037

Github version: https://github.com/NVIDIA/cutlass I wonder if we search the comments if we can find something referencing this.

It is a law of "AI" that every project must have a name collision with Nvidia: https://github.com/NVIDIA/cutlass

CuTe DSL examples from the MIT Dao-AILab for writing high performance cuda kernels using Python (see https://github.com/NVIDIA/cutlass for more background info).

> The best BLAS implementations are all closed source sadly. NVidia open-sourced CUTLASS [0] some years ago and it achieves pretty competitive performance compared to e.g. the closed-source cuBLAS. Keen observers will notice that Strassen is not used in CUTLASS. [0] https://github.com/NVIDIA/cutlass

Years ago I started a collection of convolution optimization resources: https://github.com/mratsim/laser/wiki/Convolution-optimisati... Also checked and apparently Nvidia Cutlass now supports generic convolutions: https://github.com/NVIDIA/cutlass

I would recommend the course from Oxford (https://people.maths.ox.ac.uk/gilesm/cuda/). Also explore the tutorial section of cutlass (https://github.com/NVIDIA/cutlass/blob/main/media/docs/cute/...) if you want to learn more about high performance gemm. OpenAI triton is another good resource if you want to write relatively performant cuda kernels using python for deep learning (https://openai.com/research/triton)

> we used tensor cores and managed to get back fp32 accuracy with 3 rounds of the things Hey are you referring to 3xTF32 (https://github.com/NVIDIA/cutlass/tree/master/examples/28_am...)? IMO this is a perfect example where proper abstraction could save engineers non-trivial amount of time - imagine a compiler stack which allows 3xTF32 as a normal dtype and subsequent analysis compatible with this special dtype :-)

This is a great post for people who are new to optimizing GPU code. It is interesting to see that the author got this far without interchanging the innermost loop over k to the outermost loop, as is done in CUTLASS (https://github.com/NVIDIA/cutlass). As you can see in this blog post the code ends up with a lot of compile-time constants (e.g. BLOCKSIZE, BM, BN, BK, TM, TN) one way to optimize this code further is to use an auto-tuner to find the optimal value for all of these parameters for your GPU and problem size, for example Kernel Tuner (https://github.com/KernelTuner/kernel_tuner)

Yup, in a semi-related field, NVIDIA has 3xTF32 for cases needing higher precision: https://github.com/NVIDIA/cutlass/discussions/361

So with Tensorcores you use TF32 which is more like FP19-ish and the marketing makes you think you get 8x the performance. But if you want actual FP32 precision you will need something like [1] but then your performance in the Tensorcore path is _only_ 2X faster than the SIMT path. I'll leave the prefix sum for other devs who know more :D https://github.com/NVIDIA/cutlass/blob/master/examples/27_am... //part of nod.ai/shark team

The GEMM instructions are sent to Tensor Cores by the scheduler. The cutlass library is the open source version of cublas: https://github.com/NVIDIA/cutlass