RTX 4090 Support - Compilation Issues with ThunderKittens

[MIracleyin]

Labels: question, enhancement, help wanted

---

Issue Description

I'm unable to compile the low-latency-llama demo on RTX 4090 GPUs due to ThunderKittens compatibility issues. The project seems designed primarily for H100+ architectures.

Environment

• GPU: RTX 4090
• CUDA: 12.4
• OS: Linux
• Python: 3.12

Compilation Errors

error: barrier is not a template
error: identifier "semaphore" is undefined  
error: name followed by "::" must be a class or namespace name (move<T>::lds, etc.)
static_assert(NUM_PAGES == 13, "NUM_PAGES must be 13"); // Fails on RTX 4090

Root Cause

ThunderKittens uses Hopper/Blackwell-specific features not available on RTX 4090:

• TMA operations
• Advanced barrier/semaphore primitives
• Architecture-specific memory limits (RTX 4090: 100KB shared memory vs H100: 227KB)

Questions

1. Is RTX 4090 officially supported?
2. Are there plans to add RTX 4090 support?
3. Would contributions for RTX 4090 compatibility be welcome?

Workaround

I've created a basic CUDA test that compiles and runs successfully on RTX 4090, confirming the environment is correct. The issue is specifically with ThunderKittens advanced features.

Potential Solutions

• Conditional compilation for different architectures
• Fallback implementations using standard CUDA operations
• Architecture-specific configuration constants
• Clear documentation of supported GPUs

RTX 4090 is widely used in research/development, so adding support would significantly expand the user base. Happy to contribute if there's interest!

Additional Info

• Basic CUDA compilation works fine
• Python bindings work correctly
• Issue is specifically with ThunderKittens library features
• Already fixed some Makefile and config issues locally

[LiuYang328]

I also have this error situation

[OscarSavNS]

I have this same error on A100 PCIe, and image runpod/pytorch:2.4.0-py3.11-cuda12.4.1-devel-ubuntu22.04.

[moonlockwood]

I also get this error. It will build for H100 but not A100 or 4090.

[nmstoker]

Based on the allowed values for the GPU environment variable, which are just H100 or B200, it looks to me as if the 4090 wouldn't be supported. See here:

export GPU=H100 # options are {H100, B200}, else defaults to B200
cd demos/low-latency-llama
make

[OscarSavNS]

Based on the allowed values for the GPU environment variable, which are just H100 or B200, it looks to me as if the 4090 wouldn't be supported. See here:

export GPU=H100 # options are {H100, B200}, else defaults to B200
cd demos/low-latency-llama
make

Thanks a recent commit unfortunately, based on it seems like it won't be supported: 7309cec

I had a big dig through the TK code, trying to see if it could be generalised to AMD, but the code is heavily optimised for Hopper and Blackwell. Which is totally fine, I would love for this to be generalised to other targets, but you can only expect so much. Would love to work on generalising it.