19 ms·
Rust-CUDA: Fast GPU code fully in Rust
- esjeon 5y agoThe 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?
- sxp 5y 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 https://github.com/gfx-rs/wgpu which lets you target Vulkan, Metal, DX, GL or WASM+WebGPU with rust?
- smoldesu 5y agoI 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 5y agoIf it's not relative to anything, than the word "fast" doesn't have much meaning.
- seeekr 5y ago"Fast", to me, from a software development perspective can still be meaningful, in the sense of knowing what techniques, patterns, paths, ... enable performant execution, and providing easy and straightforward paths for the user along those. Which, ultimately, leads to high performance in most ways the user will apply the provided framework (in a more general sense of the word). Hope that makes sense. And it must be OK to claim "fast" as a goal, from the early stages of a project, even before it may be possible to create any meaningful benchmarks. As long as it's discernable for the intended audience the precise stage of development or maturity the project is currently at. Which, I believe, the project in question is communicating just fine ("still in early development").
- outworlder 5y agoWould have it made any difference had the parent mentioned OpenCL?
- rak1507 5y agoWhat? 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.
- nynx 5y agowgpu is a library for running wgsl on GPUs, not Rust.
- zozbot234 5y agohttps://github.com/embarkstudios/rust-gpu https://github.com/embarkstudios/rust-gpu would be the closest equivalent, AFAICT.
- deleted 5y ago[deleted]
- nextaccountic 5y agoIt 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/runners/wgpu https://github.com/EmbarkStudios/rust-gpu/tree/main/examples... (on the "builder" directory it builds shaders with the spirv-builder crate)
- raggi 5y agoI 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.
- wpietri 5y agoWhen 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 5y agoTo 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.
- wpietri 5y agoI agree it could be read both ways. I'm just saying what I think their intent is.
- russdpale 5y agovery cool!
- gaze 5y agoThe 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 5y agoBut 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.
- imron 5y ago> I might suggest... Best way to suggest is with a pull request.
- ajkjk 5y agoNot if you're critiquing the way some writes English..
- imron 5y agoDepends how it’s worded. If someone with copy editing experience made constructive suggestions for the readme and other documentation, it may well be appreciated. A PR also gives the maintainers a way to discuss wording if there are points of disagreement - far more so than a post on HN that they may or may not even see.
- rdambrosio 5y 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-pointer-aliasing/ 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.
- dragontamer 5y agohttps://github.com/RDambrosio016/Rust-CUDA/blob/master/guide/src/features.md 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).
- nynx 5y agoLooks like they just haven't gotten around to it. Rust has fine language support for atomics and shared memory.
- dragontamer 5y ago
- zozbot234 5y agoHow would this compare with Accel (which is also built on the cuda ecosystem)?
- rdambrosio 5y agoAccel 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/src/faq.md#why-not-use-rustc-with-the-llvm-ptx-backend https://github.com/RDambrosio016/Rust-CUDA/blob/master/guide...
- glacambre 5y 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.
- m0zg 5y agoNVIDIA 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.
- pjmlp 5y agoNVidia 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.
- m0zg 5y agoNVIDIA 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.
- pjmlp 5y agoWith what ROI? This isn't charity.
- wongarsu 5y agoEnable more people to build more GPU-accelerated software, which in turn drives sales of their GPUs. A cuda-specific library like this also strengthens their monopoly on GPU compute. It's hard to put in numbers because it depends on the software that will be built. A single marginally popular software adding or improving GPU-acceleration could easily increase GPU sales by $2M/year
- pjmlp 5y agoThey already made it when CUDA 3.0 introduced PTX instead of being C focused like most Khronos stuff. NVidia is not to blame when others fail to execute and aren't able to provide the same polyglot level and IDE integrations as NVidia. They don't have any issues trying to keep up with demand, on the contrary, there aren't enough of them to sell.
- nextaccountic 5y agoCould this maybe support AMD's HIP as well?
- rdambrosio 5y agoNot 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
- booleancoercion 5y agosus
- Djrhfbfnsks 5y agoHow does the compare with writing GPU code in Julia?
- neatze 5y 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 5y agoWould 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 https://github.com/UpsettingBoy/gpgpu-rs
- DiabloD3 5y agoSince 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.