rdambrosio

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.

There is one issue i dont know how to solve, its generic kernels. It's a bit impossible to do without doing some seriously weird handling of it in rustc. I can monomorphize ahead of time but monomorphization from the CPU that talks from the CPU rustc to GPU rustc is just... it seems a bit impossible to do. Which is why a thrust-like library could only be made for specific types like f32, f64, etc.

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

That is basically what im going to do. It will break down approximately like this:

  pub fn get_shared_mem_ptr<const Bytes: usize>() -> \*mut u8 {
    __nvvm_get_shared_mem_ptr(Bytes)
  }

  __shared__ int foo[5];

Cuda-rust will probably add (probably does already) some sort of inline assembly Already done and used a lot in cuda_std, it uses normal asm!.

My idea for shared memory is that it will always be unsafe and the only thing i can do is: - expose raw low level "give me a pointer to 70 bytes of shared mem" function because this needs codegen support.

- Then expose a higher level "give me an array of 50 u16s in shared memory"

- Make as much of it as possible usable behind an abstraction, like block reduce, stuff like that.

But i want to make it clear that shared memory will always be unsafe, shared memory semantics are literally impossible to statically prove. I dont think this is the end of the world, high performance GPU programming will always require people to really know what they're doing

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...

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