Booth 0.6.1

Kia ora, G'day and hello!

Here is Booth 0.6.1! I was going to add this piecemeal over time but ended up not doing that lmao.

All of ggml-cuda compiles down through Booth now, with full SASS coverage (chip dependent), woohoo! It's been great sitting with a bunch of machine dumps all day.

A massive thanks to the folks at NAK, who've been pointing me in the right direction with a lot of this. When I say I'm reverse engineering NVIDIA GPUs, it's less me reversing ALL GPUs from scratch and more building on top of other folks' work, because my specific generation was the gap. I've also sent a few contributions to NAK, which has been great, except I've had to do a crash course in Rust.

Also a big thanks to Hot Aisle (https://hotaisle.xyz) because I was able to compile ggml to AMD too.

I've moved the MLIR frontend out for now. Ya boy is tired, and I think I need to contribute to MLIR over in LLVM for a bit so I can understand what's going on with it. It was a huge amount of code to think about every time I fiddled with my own IR.

Elsewhere:

• kath cuda builds a whole program the way nvcc does, host half and all, with clean-room cuda.h, cuda_runtime.h, cuda_fp16.h, cuda_bf16.h and cublas_v2.h.
• kath cuda -c compiles one .cu to one relocatable object, with -fPIC for shared libraries, so many translation units link into one library.
• --nvidia-cubin puts every kernel of a module into one cubin, spills to the local frame when a kernel runs out of registers, and lowers local memory, half and bfloat16 conversions, votes, membar, nanosleep and wmma.
• ldmatrix and mma.sync reach SASS as LDSM, HMMA and IMMA.
• Explicit template instantiation, templates handed out through function pointers, and launches through a function-pointer local all resolve to the right kernel.

Kindest,
Zane