A Gentle Introduction to CUDA PTX

https://philipfabianek.com/posts/cuda-ptx-introduction/

59•ashvardanian•4mo ago

Comments

the_panopticon•4mo 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•4mo 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•4mo ago

They didn’t employ custom PTX instructions; they used existing ones in ways they were not designed to be used.

the__alchemist•4mo 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•4mo 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.

neuroelectron•4mo ago

That's not much different than a modern CPU with an OS on top; where you have the OS doing some of the scheduling then the CPU is splitting up the instructions into microinstructions and then scheduling them again in finer detail (hyperthreading and such). Seems to me there must be a C-level syntax and compiler so you're not manually splitting up individual adds and such and is still capable of optimizing the math effectively. But if that were true, we wouldn't have AAA game studios going to NVidia to optimize their game engines for each individual game.

checker659•4mo ago

https://en.wikipedia.org/wiki/Intermediate_representation

saagarjha•4mo 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.

LegalArgumentException: From Courtrooms to Clojure – Sen [video]

US moves to deport 5-year-old detained in Minnesota

If you lose your passport in Austria, head for McDonald's Golden Arches

Show HN: Mermaid Formatter – CLI and library to auto-format Mermaid diagrams

RFCs vs. READMEs: The Evolution of Protocols

Kanchipuram Saris and Thinking Machines

Chinese chemical supplier causes global baby formula recall

I've used AI to write 100% of my code for a year as an engineer

Looking for 4 Autistic Co-Founders for AI Startup (Equity-Based)

AI-native capabilities, a new API Catalog, and updated plans and pricing

What changed in tech from 2010 to 2020?

From Human Ergonomics to Agent Ergonomics

Advanced Inertial Reference Sphere

Toyota Developing a Console-Grade, Open-Source Game Engine with Flutter and Dart

Typing for Love or Money: The Hidden Labor Behind Modern Literary Masterpieces

Show HN: A longitudinal health record built from fragmented medical data

CoreWeave's $30B Bet on GPU Market Infrastructure

Creating and Hosting a Static Website on Cloudflare for Free

"The Stanford scam proves America is becoming a nation of grifters"

Elon Musk on Space GPUs, AI, Optimus, and His Manufacturing Method

X (Twitter) is back with a new X API Pay-Per-Use model

Zlob.h 100% POSIX and glibc compatible globbing lib that is faste and better

Show HN: Deterministic signal triangulation using a fixed .72% variance constant

Scientists Discover Levitating Time Crystals You Can Hold, Defy Newton’s 3rd Law

When Michelangelo Met Titian

Solving NYT Pips with DLX

Baldur's Gate to be turned into TV series – without the game's developers

Interview with 'Just use a VPS' bro (OpenClaw version) [video]

EchoJEPA: Latent Predictive Foundation Model for Echocardiography

Disablling Go Telemetry