4 ms·

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 han

• 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.)
      • 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();