Rust-CUDA: Fast GPU code fully in Rust

https://github.com/RDambrosio016/Rust-CUDA

by amadeusine • 5 years ago

269 85 5 years ago

sxp• 5 years ago

"Extremely fast"

When people make claims like this, it would be good if they put the benchmarks on the first page. E.g, how does it compare with https://github.com/gfx-rs/wgpu which lets you target Vulkan, Metal, DX, GL or WASM+WebGPU with rust?

smoldesu• 5 years ago

I hope this doesn't come off as handwaving, but you're kinda comparing apples to oranges here. CUDA has always been in a class of it's own when it comes to GPU compute, for better and worse, so I think the people out there who want to use this will pretty quickly get an idea of who it's for. Benchmarks would be nice, but I don't really think they'd be germane when comparing a proprietary compute system with a generic cross-platform GPU binding.

joe_guy• 5 years ago

If it's not relative to anything, than the word "fast" doesn't have much meaning.

outworlder• 5 years ago

Would have it made any difference had the parent mentioned OpenCL?

rak1507• 5 years ago

What? So you can just claim something is 'fast' with no evidence, but then if it is slower than proprietary things, who cares? That's ridiculous - if you claim something is 'extremely fast' you should back that up with benchmarks.

nynx9• 5 years ago

wgpu is a library for running wgsl on GPUs, not Rust.

zozbot2342• 5 years ago

https://github.com/embarkstudios/rust-gpu would be the closest equivalent, AFAICT.

nextaccountic2• 5 years ago

It turns out that when wgpu is combined with rust-gpu, it can run rust on gpus too

https://github.com/EmbarkStudios/rust-gpu/tree/main/examples...

(on the "builder" directory it builds shaders with the spirv-builder crate)

raggi• 5 years ago

I wrote some OpenCL code recently, wrapped in a rust program using opencl3. My CL code was pretty slow early on, it's a math heavy problem implemented initially with a bunch of for loops.

I largely "optimized" it by smattering the code with #pragma unroll, which was exceptionally effective for the problem at hand, given that I had tons of statically defined loops that could be easily unrolled.

I know rust has a tendency to aggressively inline and flatten everything, but I'm curious about things like this, where it can be so important in this domain to ensure things are unrolled, as the loop conditionals can otherwise introduce pipeline/concurrency issues.

wpietri2• 5 years ago

When they say "extremely fast GPU code" I think they mean relative to Rust's normal home: standard CPUs. So I don't think the claim needs any specific support.

quietbritishjim• 5 years ago

To me, it is clear that "fast GPU code" means GPU code that is faster than other GPU code.

Like a "fast sports car" would obviously have to be faster than other sports cars. You couldn't (truthfully) make that claim if was faster than regular cars but slower than other sports cars.

"Accelerate Rust code by using the GPU" would be a more honest description.

russdpale• 5 years ago

very cool!

gaze• 5 years ago

The writing is better than I might have produced as a first year college student, but this needs copy editing. I might suggest that the word “extremely” should be removed not only from the entire repository but all of the user’s repositories. Fast might be removed as well. Nobody is trying to generate slow GPU code. The salient feature is that one can write the same code as one might write in CUDA with the advantages of Rust’s type system—-which is indeed useful! However, there’s no speed to be gained by using this relative to CUDA.

jacquesm• 5 years ago

But there is relative to plain rust. Besides, even if nobody is trying to write slow GPU code, it's a very easy thing to get subtly wrong resulting in 10's of % speed loss.

imron2• 5 years ago

I might suggest...

Best way to suggest is with a pull request.

ajkjk• 5 years ago

Not if you're critiquing the way some writes English..

rdambrosio7• 5 years ago

there’s no speed to be gained by using this relative to CUDA.

That is not totally true, there are two main things that can make kernels generated by this codegen faster:

- noalias, which is the LLVM equivalent of __restrict__, CUDA can take massive advantage of noalias by using readonly cache if the pointer does not alias. If you don't believe me just take a look at nvidia's blog post: https://developer.nvidia.com/blog/cuda-pro-tip-optimize-poin...

- CUDA builds all the .cu files and links the PTX of them together, this means no LTO happens by default. I do something different where i actually lazily-load every module using dependency graphs, which gives you the benefit of LTO by default. Its not perfect because right now it leaves behind a lot of dead code, but i know how to fix it.

scottlamb• 5 years ago

I might suggest that the word “extremely” should be removed not only from the entire repository but all of the user’s repositories.

You've reminded me of this classic, often (mis?)attributed to Mark Twain: [1]

“Substitute 'damn' every time you're inclined to write 'very;' your editor will delete it and the writing will be just as it should be.”

[1] https://quoteinvestigator.com/2012/08/29/substitute-damn/

dragontamer14• 5 years ago

https://github.com/RDambrosio016/Rust-CUDA/blob/master/guide...

* Missing Atomics -- Gamebreaker IMO. Atomics are absolutely essential when you are dealing with 10,000+ threads on a regular basis. You'll inevitably come across a shared data-structure that requires write-access from each thread, and some coordination mechanism is needed for that. Atomics are one important fit.

Ironic, a few days ago, I argued for the use of Fork-join parallelism in most cases (aka: Kernel launch / synchronized kernel exits). Now I find myself arguing the opposite now that we have a topic here with missing atomics. Like... atomics need to be used very, very rarely, but those rare uses are incredibly important.

* Warp Vote / Match / Reduce / Shuffle missing (Very useful tools for highly-optimized code, but you can write slower code that does the same thing through __shared__ memory just fine)

------

Wait, does this support __shared__ memory at all? Raw access to memory is not really amenable to Rust's programming style, but its absolutely necessary for high-performance GPU programming.

If this is missing __shared__ memory concepts, then the entire idea of "efficient GPU code" is dead IMO. GPU threads can only communicate quickly over __shared__ memory within an OpenCL Workgroup / CUDA Block (A Workgroup or Block is roughly a grouping of 1024 "threads" or SIMD-lanes)

All other forms of communication are incredibly slow. Atomics are maybe the next fastest form of communication, but only across __shared__ memory. Relaxed Atomics to global memory are reasonably performant but once you have either Seq-cst or Acquire/Release semantics (aka: the right memory barriers in the right place), things slow down dramatically in GPU-land.

The big issue is that __shared__ memory is only 64kB in size, its locked down to workgroups / blocks. In NVidia GPUs, the __shared__ memory "eats" into your L1 cache as well (In fact: __shared__ memory can be thought of as programmer-managed cache. The caching heuristics just aren't good enough for high-performance GPU programmers. They want to manually manage that high-speed memory for maximum performance).

nynx9• 5 years ago

Looks like they just haven't gotten around to it. Rust has fine language support for atomics and shared memory.

dragontamer14• 5 years ago

You're lucky I've had this discussion before with other Rust programmers. But I forgot about this issue...

CUDA __shared__ memory is a 64kB block of RAM that is located __INSIDE THE CORE__ and is incredibly fast, but has very peculiar semantics. (Since the memory never "leaves" the core, its "stuck" only on a small set of 1024 "threads". Difficult to manage when you write 30,000+ thread programs but the speed benefits are insane and well worth the trouble)

Rust "shared" memory is thread-to-thread communications that simply doesn't exist in the GPU-world.

-------

Maybe it'd be more appropriate if I used OpenCL terminology instead of CUDA terminology here, because Rust programmers have an unfortunate name conflict. In OpenCL, this 64kB on-core buffer is called __local memory. Does that help?

Whenever I said "__shared__", replace that in your mind with __local instead. Because CUDA __shared__ is NOTHING like Rust-Shared.

rdambrosio7• 5 years ago

As i mentioned, it is an early project, just making the simplest kernel compile was very difficult. Atomics and shared memory are great, but both are very difficult. Atomics need "proper" atomics (i.e. special instructions on sm_70+ and emulated on <sm_70), and shared mem needs some weird codegen support. I will get to both of them. Nevertheless, noalias does cause significant performance speedups in memory bound kernels, see this blogpost: https://developer.nvidia.com/blog/cuda-pro-tip-optimize-poin...

So please do not be surprised that an early project does not contain every single feature of cuda, something thats been around for decades

dragontamer14• 5 years ago

No problem. I understand its a work in progress.

I'd push most strongly for CUDA __shared__ support first along with thread-barriers (CUDA's __syncthreads()), followed by __shared__ atomics. Finally, global atomics + associated memory-barrier stuffs (Ex: seq-cst atomic, acq-release atomic would work but maybe be a bit difficult. Might be easier to support the older-style memory barrier instead?)

--------

EDIT: Alternatively, maybe making a Thrust-like library for Rust is a better step 1? High performance GPU-code is very peculiar and strange. I can't imagine there's a big market for it. It seems like most people writing GPU "code" these days are just Python programmers punching in a Tensorflow, rather than actually trying to write a high-performance GPU thingy.

WithinReason6• 5 years ago

Atomics are absolutely essential

Most problems don't need atomics to solve

If this is missing __shared__ memory concepts, then the entire idea of "efficient GPU code" is dead IMO

Sure, shared memory is great, but not always needed and communicating through global memory can be fast if you stay inside the L1/L2 cache and hide latency.

dragontamer14• 5 years ago

Most problems don't need atomics to solve

How do you handle a global, concurrent memory write and/or read ?? (across many different blocks, maybe even across different grids).

For example: lets say you have a global hash table and 30,000 CUDA-threads are running. How do you insert data into the hash table safely?

Sure, shared memory is great, but not always needed and communicating through global memory can be fast if you stay inside the L1/L2 cache and hide latency.

Scan operations (sum, min, max, AND, OR to name a few) through the parallel prefix pattern (https://en.wikipedia.org/wiki/Prefix_sum). How do you plan to do it, if not through __shared__ memory?

This is a fundamental operation in almost all GPU code I've ever seen. Just pushing / popping to a stack will require a prefix-sum to determine the size across the workgroup.

If you can't do prefix-sum, you won't be able to do effective load-balancing on a GPU. This is something that'd normally take a dozen clock ticks, but if you do it over L2 you're looking at hundreds of clock ticks instead.

------

Sorting networks are also probably best implemented in __shared__ memory... with maybe warp-level intrinsics beating them. (But warp-level programming is much much harder and I prefer to avoid it).

shaklee35• 5 years ago

shared memory and L1 are the same on modern GPUs, but shared memory gives you control of what goes on there. I agree with dragontamer that it's absolutely essentially to have atomics and shared.

zozbot2342• 5 years ago

How would this compare with Accel (which is also built on the cuda ecosystem)?

rdambrosio7• 5 years ago

Accel uses the LLVM PTX backend, which is unusable for serious projects and doesn't work on windows, i have more about it here https://github.com/RDambrosio016/Rust-CUDA/blob/master/guide...

glacambre• 5 years ago

The LLVM PTX backend is still very much WIP and often doesn't have things and/or breaks.

What is missing in your opinion aside from working debug info? I've worked on experimental CUDA support in GNAT (the Ada frontend of GCC), using LLVM's ptx backend to target GPUs and it was mostly working.

m0zg3• 5 years ago

NVIDIA should hire the guy, then hire whoever he says he wants on the team and let them rip. That's what I'd do if I were in their shoes. Viable paths off C/C++ are badly needed, and currently the only real viable path with an ecosystem and community is Rust.

pjmlp3• 5 years ago

NVidia decided for Ada against Rust for their automotive firmware, their cards are designed based on C++'s memory model, and they have a big ISO C++ presence.

m0zg3• 5 years ago

NVIDIA has so much cash these days that they don't have to do just one thing. They can do _all_ the things, at the same time. This realistically would only take a 5-7 person team, including the manager. That's $2M/yr tops, all in, which is peanuts for NVIDIA. Much less if the author is not based in the US.

imtringued• 5 years ago

Yeah do it before he gets the crazy idea of supporting AMD gpus...

nextaccountic2• 5 years ago

Could this maybe support AMD's HIP as well?

rdambrosio7• 5 years ago

Not for the near future, HIP does not seem to be language-agnostic and the codegen is for NVVM IR, not whatever AMD uses. It might be possible to target amdgpu with llvm because all the gpu-unfriendly things are gone in my codegen. So maybe in the future? im not sure

esjeon• 5 years ago

The title says fast, but no benches, but the README puts more emphasis on that it’s a more usable solution than LLVM-PTX + Rust.

I mean, what’s the point of “fast”, if LLVM-PTX is clunky with rust in the first place?

Djrhfbfnsks• 5 years ago

How does the compare with writing GPU code in Julia?

neatze• 5 years ago

n ecosystem of libraries and tools for writing and executing extremely fast GPU code fully in Rust.

Well at least it does not say extremely safe.

127• 5 years ago

Would be really nice to have an actual cross platform GPGPU library. It's really holding every kind of progress back to have only vendor lock-in.

Maybe WebCPU will be capable of compute to the extent that CUDA isn't necessary. https://github.com/UpsettingBoy/gpgpu-rs

DiabloD3• 5 years ago

Since this doesn't target the majority of GPUs, nor the majority of GPUs used in enterprise compute, why continue to target CUDA? If you target SPIR-V, then this can be used with all the existing modern APIs (Vulkan, OpenCL, OpenGL; Microsoft has a SPIR-V to DXIL translator, and third parties have Vulkan/SPIR-V-on-legacy-API support as well, and there is also at least one project that converts LLVM IR to SPIR-V)

... but if you target NVVM alone, you're stuck on the minority of machines that can run CUDA. Even Nvidia has admitted that SPIR-V is the future, since they are both a major contributor to Khronos specs, but also have a top notch SPIR-V compiler that produces better code for Nvidia GPUs than their NVVM compiler.

Rust-CUDA: Fast GPU code fully in Rust

Related Stories