> 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

If memory is exposed in CUDA more like allocators than pre-allocated static memory, look into using the alloc crate with custom allocators that hand out global or shared memory.

No. This is a fundamental misunderstanding of __shared__ memory and how kernels work.

__shared__ memory is a specification like .bss in ELF files. Its something that must be known at compile-time, well-before runtime. The GPU kernel, when launched, will reserve some __shared__ memory for itself.

----

Different kernels may use different chunks of that 64kB for themselves. For example, if Kernel Foo uses 30kB of __shared__ memory and is currently running, and Kernel Bar only uses 20kB of __shared__ memory, the GPU will allow Foo + Bar to run simultaneously.

The fundamental "unit of compute" is the OpenCL workgroup / CUDA block. The GPU will run many, many kernels (even different kernels) _ON THE SAME CORE_, as long as there's enough registers, and __shared__ memory available for them.

-----

"Thread local" variables compile into registers in practice. GPUs have ~256 registers per core, so if Kernel Foo uses 40 registers, and Kernel Bar uses 70, the GPU SMs (Symmetric Multiprocessors, what you'd call a "core" in the CPU world) could run 4 copies of Foo + 1 copy of Bar.

This is because Foo#1 will eventually run into a memory-latency issue (VRAM read/writes have well over 100-cycles of latency, maybe 300+ cycles on older GPUs). Instead of waiting for this memory operation to complete, the GPU will switch to another workgroup (like Foo#2, Foo#3, or Bar#1) to ensure that the GPU-cores stay utilized.

__shared__ memory works kinda like these registers, they're divy'd out at runtime by treating the OpenCL workgroups / CUDA blocks as a unit.

------

So __shared__ memory has to be preallocated by this model. Its an important per-unit resource that is tracked by the GPU at the lowest level, so that multiple kernels could be run concurrently on the same cores (GPUs are like SMT or Hyperthreading: capable of running 8+ kernels per core as long as you have enough registers / ___shared__ memory to launch all kernels)

Gotcha. I assumed that was the case, but I was seeing some other stuff that implied otherwise elsewhere in this thread.

Yeah, in that case, it'd probably have to be a transformation on top of static variables or something like that.

In the Lisp world... it is ambiguous if functions or macros are called at compile time or runtime.

I'm not sure if this is kosher in the world of Rust, but...

    static RAY_STACK: [Shared<Ray>; 2000] = [Shared::new(Ray::default()); 2000];

This could still work, if Shared::new(...) were a compile-time function. Or a language-extension that looked like a compile-time function.

EDIT: You wouldn't be allowed to have Shared::new inside of a loop or a recursive function though. But as long as you had assurances that any such Shared::new instance ran exactly once throughout the code, it might work?

Or maybe that's too ugly. "static" probably captures the idea better

The best way to do it is probably the way rust-gpu does it: https://github.com/EmbarkStudios/rust-gpu/blob/main/docs/src...

The entry point of the kernel would supply any objects that have special properties.

fwiw, and I think you know this, shared memory does not have to by preallocated. dynamic shared memory allows you to allocate at kernel launch time.

Yeah, I'm not 100% sure what to say in English though, lol.

There's compile-time and runtime. But there's also kernel-launch time? Dynamic shared memory is done before kernel-launch, possibly during cpu-runtime but before gpu-runtime.

-------

Things get crazier when you see OpenCL paradigms like... #define constants during CPU-runtime, invoke the OpenCL compiler (under the assumption that the compiler will now optimize the constant into the code directly), and then kernel-launch.

Because programmers expect arbitrary types to become __shared__ (ex: if you are writing a Raytracer, you probably want your rays to be stored into __shared__ memory).

So the programmer would write:

    struct Vec3{
        float x;
        float y;
        float z;
    }

    struct Ray{
        Vec3 origin;
        Vec3 direction;
        int bouncesRemaining;
    };

And then maybe in one function...

    __local struct Ray raystack[2000];
    // As rays bounce, they get new origins / directions, 
    // and may spawn new rays as needed.

    // If the raystack overflows, transfer the rays to global memory.

And then maybe my OpenCL kernel operates over these rays in parallel, tracking wherever they go.

I can't guarantee that it'd be implemented like this, but I could see it working like this in Rust.

  struct Ray {
      origin: Vec3,
      dir: Vec3,
      remaining_bounces: u32,
  }

and then using it like this:

  static RAY_STACK: [Shared<Ray>; 2000] = [Shared::new(Ray::default()); 2000];

or

  static RAY_STACK: Shared<[Ray; 2000]> = Shared::default();

I guess it could also be done like this:

  #[shared]
  static RAY_STACK: [Ray; 2000] = ...;

but that's not really the rust way.

This seems more likely than my proposals, but it does mean you won't be able to use a particular part of the `__shared__` region for more than one thing. (Mine are broken by function calls, so yours is still better.)

That's easy enough in Rust.

  struct Vec3(f32, f32, f32);
  struct Ray {
      origin: Vec3,
      direction: Vec3,
      bounces_remaining: i32
  }

and then something like (made-up API):

  let shared = rust_cuda::shared();
  let mut raystack = shared.alloc_zeroed::<[Ray; 2000]>();

or maybe:

  let mut raystack: Shared<[Ray; 2000]> = [Default::default(); 2000].into();

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)
  }

For the raw version, the codegen internally intercepts the call to the intrinsic and declares an extern global in the shared addrspace, which is what libnvvm wants you to do, basically like

  __shared__ int foo[5];

Dynamic shared mem is a bit more weird because if you query the ptr for the dynamic smem it yields the same ptr every time.

You might want alignment there as well.

Both alignment and unalignment actually.

GPUs are weird. Prime numbers to 'unalign' data so that you minimize bank conflicts is a common optimization trick.

GPUs don't have one memory load/store unit. They are incredibly parallel and have like 32 load/store units that try to operate in parallel.

If all your data is aligned, then bank#0 gets more requests than bank#31. (Thread#0 accesses memory 800. Thread#1 accesses memory 832. Threas#2 accesses 864... Woops you just hammered one bank and now 31 of your memory banks are sitting around doing nothing, while bank#0 is doing all the work sequentially)

Unalignment means more read/writes are sent to bank#31, and fewer to bank#0, better balancing the load across your parallel load/store units.

Here be dragons, and this person tames them. This is insane, actually. I’m guessing it would be cool if gpu automagically scrambled memory so you didn't have to manually unalign it?

When you're doing "uint32_t array[thread_idx.x]" sorts of things, you'll notice that your threads are all lined up with the array. So you're in perfect bank-alignment.

With "array[thread_idx.x]" kind of access, Thread#0 accesses array[0], Thread#1 accesses array[1]... etc. etc.

array[0] might map to memory location #0x8001200, which will probably be bank#0. array[1] might map to #0x8001204, which would be bank#1. Etc. etc. (I forget exactly how many bytes per bank, but... you get the gist).

At the end of the day, all your array[] accesses from Thread#0 through Thread#1023 of your workgroup/block will be perfectly balanced and perfectly spread out between all banks.

--------

So really, the "lesson" is to just organize your data in arrays as much as possible. GPUs are really, really good at simple array reads/writes.

That's not always possible of course. You should only "shuffle" the banks if you know for certain that one bank is going to be hit more than the other banks.

--------

It really comes down to the size of the object you made an array out of. If you have a large object for some reason, maybe array[0] and array[1], array[2], etc. etc. will all map to bank#0.

This is done in rust-gpu[1] by the `spirv(workgroup)` attribute on the kernel function signature.

[1]: https://github.com/EmbarkStudios/rust-gpu/blob/46c9ea0c9c7b7...