5 ms·

No, I do understand this. Some sort of type which exposes shared semantics will need to be exposed in the rust-cuda crate. Clearly this isn't just type-lev

by nynx 5y ago

No, I do understand this. Some sort of type which exposes __shared__ semantics will need to be exposed in the rust-cuda crate. Clearly this isn't just type-level stuff, but the way rust exposes features like that is usually through a new type. Cuda-rust will probably add (probably does already) some sort of inline assembly the same way that the spirv-rust toolchain does, and through that, a "Shared" type will be able to be implemented.

dragontamer 5y ago
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.
- nynx 5y ago
  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.
  - wizzwizz4 5y ago
    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.)
- wizzwizz4 5y ago
  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();
  - deleted 5y ago
    [deleted]
  - raphlinus 5y ago
    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/46c9ea0c9c7b7c6acfe83ad3dfaccfa21629d658/docs/src/attributes.md#workgroup-shared-memory https://github.com/EmbarkStudios/rust-gpu/blob/46c9ea0c9c7b7...
  - rdambrosio 5y ago
    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.
    - nynx 5y ago
      You might want alignment there as well.
      - dragontamer 5y ago
        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.
rdambrosio 5y ago
> 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
- nynx 5y ago
  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.
  - dragontamer 5y ago
    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)

No, I do understand this. Some sort of type which exposes __shared__ semantics will need to be exposed in the rust-cuda crate. Clearly this isn't just type-lev

by nynx 5y ago

dragontamer 5y ago

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.

nynx 5y ago
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.
- wizzwizz4 5y ago
  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.)
wizzwizz4 5y ago
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();
- deleted 5y ago
  [deleted]
- raphlinus 5y ago
  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/46c9ea0c9c7b7c6acfe83ad3dfaccfa21629d658/docs/src/attributes.md#workgroup-shared-memory https://github.com/EmbarkStudios/rust-gpu/blob/46c9ea0c9c7b7...
- rdambrosio 5y ago
  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.
  - nynx 5y ago
    You might want alignment there as well.
    - dragontamer 5y ago
      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.

rdambrosio 5y ago

> 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

nynx 5y ago
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.
- dragontamer 5y ago
  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)