Rust CUDA Project

Looks like a dead end. Why CUDA? There should be some way to use Rust for GPU programming in general fashion, without being tied to Nvidia.

There's no cross-vendor API which exposes the full power of the hardware. For example, you can use Vulkan to do compute on the GPU, but it doesn't expose all of the features that CUDA exposes, and you need to do the legwork yourself reimplementing all of the well optimized libraries (like e.g. cublas or cudnn) that you get for free with CUDA.

CUDA is the easiest-to-use and most popular GPGPU framework. I agree that it's unfortunate there aren't good alternatives! As kouteiheika pointed out, you can use Vulkan (Or OpenCL), but they are not as pleasant.

Because others so far have failed to deliver anything worthwhile using, with the same tooling ecosystem as CUDA.

While I agree, that CUDA is the best in class API for GPU programming, OpenCL, Vulkan compute shaders and Sycl are alternatives that are usable. I'm for example, using compute shaders for writing GPGPU algorithms that work on Mac, AMD, Intel and Nvidia. It works ok. The debugging experience and ecosystem sucks compared to CUDA, but being able to run the algorithms across platforms is a huge advantage over CUDA.

To deliver, you need to make Rust target the GPU in a general way, like some IR, and then may be compile that into GPU machine code for each GPU architecture specifically.

So this project is a dead end, because it's them who are these "others" - they are developing it and they are doing it wrong.

How are you writing compute shaders that work on all platforms, including Mac? Are you just writing Vulkan and relying on MoltenVK?

AFAIK, the only solution that actually works on all major platforms without additional compatibility layers today is OpenCL 1.2 - which also happens to be officially deprecated on MacOS, but still works for now.

why do you need to run across all those platforms? what's the cost benefit for doing so?

Make a compiler that takes Rust and compiles into some IR, then another compiler that compiles that IR into GPU machine code. Then it can work and that's going to be your API (what you developed in Rust).

That's the whole point of what's missing. Not some wrapper around CUDA.

It defeats the purpose. Easy to use should be something in Rust, not CUDA.

No they aren't, because they lack the polyglot support from CUDA and as you acknowledge the debugging experience and ecosystem sucks.

And is stuck with C99, versus C++20, Fortran, Julia, Haskell, C#, anything else someone feels like targeting PTX with.

Yes, MoltenVK works fine. Alternatively, you can also use WebGPU (there are C++ and Rust native libs) which is a simpler but more limiting API.

Plus IDE support, Nsight level debugging, GPU libraries, yes most likely bound to fail unless NVidia, like it happened with other languages sees enough business value to give an helping hand.

They are already using Rust in Dynamo, even though the public API is Python.

Well it really depends on the kind of work you're doing. My (non-AI) software allows users to run my algorithms on whatever server-side GPU or local device they have. This is a big advantage IMO.

Technically, OpenCL can also include inline PTX assembly in kernels (unlike any compute shader API I've ever seen), which is relevant for targeting things like tensor cores. You're absolutely right about the language limitation, though.

WebGPU has no support for tensor cores (or their Apple Silicon equivalents). Vulkan has an Nvidia extension for it, is there any way to make MoltenVK use simdgroup_matrix instructions in compute shaders?

AFAIK, MoltenVK doesn't. Dawn (Google's C++ WebGPU implementation) does have some experimental support for it [0][1].

[0] https://issues.chromium.org/issues/348702031

[1] https://github.com/gpuweb/gpuweb/issues/4195

interesting! Can you say more about what kind of algorithms your software runs?

You can look to https://github.com/Rust-GPU/rust-gpu/ for vulkan.

The purpose is to get shit done.

My work is primarily about the processing of medical images (which usually are large 3D images). Doing this on the GPU, can be up to 10-20x faster.

At which point why bother, PTX is CUDA.

Generally, the reason to bother with this approach is if you have a project that only needs tensor cores in a tiny part of the code and otherwise benefits from the cross platform nature of OpenCL, so you have a mostly shared codebase with a small vendor-specific optimization in a kernel or two. I've been in that situation and do find that approach valuable, but I'll be the first to admit the modern GPGPU landscape is full of unpleasant compromises whichever way you look.

But what about that wants to be multi-platform instead of picking one an specializing, probably picking up some more optimizations along the way?

That's a good example of of what it should be.