Rust-CUDA: Fast GPU code fully in Rust
<a href="https:&#x2F;&#x2F;github.com&#x2F;RDambrosio016&#x2F;Rust-CUDA&#x2F;blob&#x2F;master&#x2F;guide&#x2F;src&#x2F;features.md" rel="nofollow">https:&#x2F;&#x2F;github.com&#x2F;RDambrosio016&#x2F;Rust-CUDA&#x2F;blob&#x2F;master&#x2F;guide...</a><p>* Missing Atomics -- Gamebreaker IMO. Atomics are absolutely essential when you are dealing with 10,000+ threads on a regular basis. You&#x27;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.<p>Ironic, a few days ago, I argued for the use of Fork-join parallelism in most cases (aka: Kernel launch &#x2F; 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.<p>* Warp Vote &#x2F; Match &#x2F; Reduce &#x2F; 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)<p>------<p>Wait, does this support __shared__ memory at all? Raw access to memory is not really amenable to Rust&#x27;s programming style, but its absolutely necessary for high-performance GPU programming.<p>If this is missing __shared__ memory concepts, then the entire idea of &quot;efficient GPU code&quot; is dead IMO. GPU threads can only communicate quickly over __shared__ memory within an OpenCL Workgroup &#x2F; CUDA Block (A Workgroup or Block is roughly a grouping of 1024 &quot;threads&quot; or SIMD-lanes)<p>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&#x2F;Release semantics (aka: the right memory barriers in the right place), things slow down dramatically in GPU-land.<p>The big issue is that __shared__ memory is only 64kB in size, its locked down to workgroups &#x2F; blocks. In NVidia GPUs, the __shared__ memory &quot;eats&quot; into your L1 cache as well (In fact: __shared__ memory can be thought of as programmer-managed cache. The caching heuristics just aren&#x27;t good enough for high-performance GPU programmers. They want to manually manage that high-speed memory for maximum performance).