▲A Gentle Introduction to CUDA PTX(philipfabianek.com)
49 points by ashvardanian 3 days ago | 3 comments
▲the_panopticon 2 hours ago
Very interesting. It sounds like tuning at the PTX level can increase workload efficiencies, such as quote "Specifically, we employ customized PTX (Parallel Thread Execution) instructions" from the DeepSeek folks https://arxiv.org/abs/2412.19437.
▲shetaye 2 hours ago
Agreed! The gulf between pure-C++ CUDA and PTX is getting larger with these optimizations. My understanding is that Deepseek used PTX instructions that either had no corresponding C++ implemented (like `wgmma` mentioned in the article) or uncommon permutations of modifiers (`LD.Global.NC.L1::no_allocate.L2::256b`).
▲saagarjha 1 hour ago
They didn’t employ custom PTX instructions; they used existing ones in ways they were not designed to be used.
▲saagarjha 1 hour ago
It’s really not true anymore that PTX is forward compatible. There’s a subset that is but any of the new interesting interfaces that have been added are not forward compatible and change in each microarchitectural revision. Most of the reason you’d drop down to PTX anyway is to use those; otherwise compilers are fairly good these days and it’s rarely the case you’ll see PTX unless you’re profiling.
▲the__alchemist 1 hour ago
Is this analogy valid: Writing PTX is like writing assembly instead of a higher-level language (C, C++, rust etc) for CPU code? E.g. normally the higher level code compiles to it, but you can do optimizations by going lower?

For context, like the opening paragraph in the article goes into, I generate PTX code regularly, but have no idea what the actual code in the PTX file means!

I'm curious about the forward compatibility the article goes into. I only experience that to a point: Code compiled on Cuda 12 does not seem to work on machines with Cuda 13.

▲philipfabianek 1 hour ago
Indeed, this is one way to think about it. However, PTX is an instruction set for a virtual machine, not the actual hardware. The true, hardware-specific assembly is called SASS (Streaming Assembly) and the PTX code is translated into SASS by the GPU driver (using ptxas) in a final compilation step. Unlike SASS, PTX is (mostly) forward compatible.

I don't know the details about your CUDA 12 vs. 13 issue but I think it is not about hardware compatibility but more about the software stack. An application linked against CUDA 12 libraries and might not work with CUDA 13 libraries.