3 ms·

Looks like they just haven't gotten around to it. Rust has fine language support for atomics and shared memory.

by nynx 5y ago

dragontamer 5y ago
You're lucky I've had this discussion before with other Rust programmers. But I forgot about this issue... CUDA __shared__ memory is a 64kB block of RAM that is located __INSIDE THE CORE__ and is incredibly fast, but has very peculiar semantics. (Since the memory never "leaves" the core, its "stuck" only on a small set of 1024 "threads". Difficult to manage when you write 30,000+ thread programs but the speed benefits are insane and well worth the trouble) Rust "shared" memory is thread-to-thread communications that simply doesn't exist in the GPU-world. ------- Maybe it'd be more appropriate if I used OpenCL terminology instead of CUDA terminology here, because Rust programmers have an unfortunate name conflict. In OpenCL, this 64kB on-core buffer is called __local memory. Does that help? Whenever I said "__shared__", replace that in your mind with __local instead. Because CUDA __shared__ is NOTHING like Rust-Shared.
- nynx 5y ago
  That sounds a little complicated to deal with, but I see no reason why either the Rust atomic types or a new type supplied by the rust-cuda library couldn't handle that just fine. I just want to make sure that you realize that Rust does have regular atomics (and that's how other shared memory abstractions are generally implemented underneath).
  - dragontamer 5y ago
    EDIT: Okay, I'm switching to OpenCL terminology. I think this __shared__ thing is getting confusing for ya. ---------- The semantics of __local memory variables are very peculiar, and are prone to memory ordering issues unless coordinated with OpenCL's barrier() function calls. That means the compiler needs to be careful about optimizations and orderings. The compiler's understanding of allowable variable orders must match the programmer's understanding. __local variables can be of any type. In CUDA, you can make new structs or even C++ classes be in __shared__ memory. In OpenCL, you can make *arbitrary* structs be __local. Its entirely freeform use of memory, albeit tiny and 64kB in size. --------- The proper implementation of __local semantics will require compiler support. This isn't something you can just tack on with a type-system. The implications of __local reverberate not only through the type system, but also with optimization and the very understanding of how code gets compiled fundamentally (synchronization and memory orderings). -------- > I just want to make sure that you realize that Rust does have regular atomics (and that's how other shared memory abstractions are generally implemented underneath). And __local memory is nothing like you've ever seen before, unless maybe you're a realtime programmer who has those weird manually-managed L1 CPUs. Because that's what it is: __local memory is a manually managed high-speed RAM. While typical CPU code relies upon the hardware to manage L1 cache for you, GPU programmers regularly manage that region of memory *manually*. Its not easy, but its the only way to reach the highest-levels of performance.
    - 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.)
- DSingularity 5y ago
  The 1024 threads of a warp/block/whatever or just the current threads or what?
  - dragontamer 5y ago
    > The 1024 threads of a warp/block/whatever In CUDA terminology: the block can access __shared__ memory together. Different blocks are locked out of seeing other block's __shared__ memory. A block can be up to 32-warps working together (aka 1024 threads), or it could be as small as 1-warp (aka 32-threads). A warp could be 1-thread in some cases. -- Finally, a grid in CUDA is synchronized across kernel calls. You are 100% certain that all threads of the grid are not running before the kernel_launch<<<x, y, z, stream>>(foo bar)... and you can be 100% certain that all threads of the grid are done after the cudaStreamSynchronize(stream) call. --------- Warps: Largely about very low-level details such as branch divergence. Warps take if/else/for loops together in practice, so you need to think about warps when you think about optimal utilization. In the past 5 years or so, warp-level programming has become more popular, but warp-level stuff is pretty rare and only should be reserved once other, easier, optimizations have taken place. Blocks: Coordinated across __shared__ memory. Because blocks are controlled by the programmer, its easier to think about than warps... but you need to be careful and put the right OpenCL barriers() or CUDA __syncthreads() in the right places. Grid: Coordinated across cudaStreamSynchronize / CUDA-streams. Ad-hoc coordination can be done with atomics to global memory, but this is very slow if memory barriers are involved. (Relaxed atomics are pretty fast though)
    - DSingularity 5y ago
      Thanks!

Looks like they just haven't gotten around to it. Rust has fine language support for atomics and shared memory.

by nynx 5y ago

Looks like they just haven't gotten around to it. Rust has fine language support for atomics and shared memory.

You're lucky I've had this discussion before with other Rust programmers. But I forgot about this issue... CUDA __shared__ memory is a 64kB block of RAM that is located __INSIDE THE CORE__ and is incredibly fast, but has very peculiar semantics. (Since the memory never "leaves" the core, its "stuck" only on a small set of 1024 "threads". Difficult to manage when you write 30,000+ thread programs but the speed benefits are insane and well worth the trouble) Rust "shared" memory is thread-to-thread communications that simply doesn't exist in the GPU-world. ------- Maybe it'd be more appropriate if I used OpenCL terminology instead of CUDA terminology here, because Rust programmers have an unfortunate name conflict. In OpenCL, this 64kB on-core buffer is called __local memory. Does that help? Whenever I said "__shared__", replace that in your mind with __local instead. Because CUDA __shared__ is NOTHING like Rust-Shared.

nynx 5y ago
That sounds a little complicated to deal with, but I see no reason why either the Rust atomic types or a new type supplied by the rust-cuda library couldn't handle that just fine. I just want to make sure that you realize that Rust does have regular atomics (and that's how other shared memory abstractions are generally implemented underneath).
- dragontamer 5y ago
  EDIT: Okay, I'm switching to OpenCL terminology. I think this __shared__ thing is getting confusing for ya. ---------- The semantics of __local memory variables are very peculiar, and are prone to memory ordering issues unless coordinated with OpenCL's barrier() function calls. That means the compiler needs to be careful about optimizations and orderings. The compiler's understanding of allowable variable orders must match the programmer's understanding. __local variables can be of any type. In CUDA, you can make new structs or even C++ classes be in __shared__ memory. In OpenCL, you can make *arbitrary* structs be __local. Its entirely freeform use of memory, albeit tiny and 64kB in size. --------- The proper implementation of __local semantics will require compiler support. This isn't something you can just tack on with a type-system. The implications of __local reverberate not only through the type system, but also with optimization and the very understanding of how code gets compiled fundamentally (synchronization and memory orderings). -------- > I just want to make sure that you realize that Rust does have regular atomics (and that's how other shared memory abstractions are generally implemented underneath). And __local memory is nothing like you've ever seen before, unless maybe you're a realtime programmer who has those weird manually-managed L1 CPUs. Because that's what it is: __local memory is a manually managed high-speed RAM. While typical CPU code relies upon the hardware to manage L1 cache for you, GPU programmers regularly manage that region of memory *manually*. Its not easy, but its the only way to reach the highest-levels of performance.
  - 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.)
DSingularity 5y ago
The 1024 threads of a warp/block/whatever or just the current threads or what?
- dragontamer 5y ago
  > The 1024 threads of a warp/block/whatever In CUDA terminology: the block can access __shared__ memory together. Different blocks are locked out of seeing other block's __shared__ memory. A block can be up to 32-warps working together (aka 1024 threads), or it could be as small as 1-warp (aka 32-threads). A warp could be 1-thread in some cases. -- Finally, a grid in CUDA is synchronized across kernel calls. You are 100% certain that all threads of the grid are not running before the kernel_launch<<<x, y, z, stream>>(foo bar)... and you can be 100% certain that all threads of the grid are done after the cudaStreamSynchronize(stream) call. --------- Warps: Largely about very low-level details such as branch divergence. Warps take if/else/for loops together in practice, so you need to think about warps when you think about optimal utilization. In the past 5 years or so, warp-level programming has become more popular, but warp-level stuff is pretty rare and only should be reserved once other, easier, optimizations have taken place. Blocks: Coordinated across __shared__ memory. Because blocks are controlled by the programmer, its easier to think about than warps... but you need to be careful and put the right OpenCL barriers() or CUDA __syncthreads() in the right places. Grid: Coordinated across cudaStreamSynchronize / CUDA-streams. Ad-hoc coordination can be done with atomics to global memory, but this is very slow if memory barriers are involved. (Relaxed atomics are pretty fast though)
  - DSingularity 5y ago
    Thanks!