5 ms·

> Atomics are absolutely essential Most problems don't need atomics to solve > If this is missing shared memory concepts, then the entire idea of "efficie

> Atomics are absolutely essential Most problems don't need atomics to solve > If this is missing __shared__ memory concepts, then the entire idea of "efficient GPU code" is dead IMO Sure, shared memory is great, but not always needed and communicating through global memory can be fast if you stay inside the L1/L2 cache and hide latency.

dragontamer 5y ago
> Most problems don't need atomics to solve How do you handle a global, concurrent memory write and/or read ?? (across many different blocks, maybe even across different grids). For example: lets say you have a global hash table and 30,000 CUDA-threads are running. How do you insert data into the hash table safely? > Sure, shared memory is great, but not always needed and communicating through global memory can be fast if you stay inside the L1/L2 cache and hide latency. Scan operations (sum, min, max, AND, OR to name a few) through the parallel prefix pattern (https://en.wikipedia.org/wiki/Prefix_sum https://en.wikipedia.org/wiki/Prefix_sum). How do you plan to do it, if not through __shared__ memory? This is a fundamental operation in almost all GPU code I've ever seen. Just pushing / popping to a stack will require a prefix-sum to determine the size across the workgroup. If you can't do prefix-sum, you won't be able to do effective load-balancing on a GPU. This is something that'd normally take a dozen clock ticks, but if you do it over L2 you're looking at hundreds of clock ticks instead. ------ Sorting networks are also probably best implemented in __shared__ memory... with maybe warp-level intrinsics beating them. (But warp-level programming is much much harder and I prefer to avoid it).
- floatboth 5y ago
  Global hash table on GPUs sounds cursed and perverted. What you’re meant to do with them is run the same computation on all pixels in parallel independently! :P
- WithinReason 5y ago
  As I said, these features are often not needed. You can implement e.g. a neural network library without needing atomic operations. > How do you plan to do it, if not through __shared__ memory? Can't you use __shared__ memory the same way you use workgroup barriers and global memory? Might be slower, but good caching should make it comparable, which should be the case of prefix sum (you read right after writing, so should get good cache hit probability).
  - dragontamer 5y ago
    > As I said, these features are often not needed. Global parallel hash table is probably the fastest way to implement collision detection. Its pretty fundamental to manipulation of 3d space, be it graphics, physics or other such simulations. Which of course, run _GREAT_ on GPUs. -------------- I've written "Inner Join" on a GPU for fun. Yes, the SQL operator. Its pretty fast. Databases probably can run on GPUs and parallelize easily. But any database would need globally consistent reads/writes. Sure, GPUs have less RAM than a CPU, but GPU-RAM is way faster so that might actually be a net benefit if your data is between 200MB and 4GB in size. Use your imagination. Anywhere you'd use an atomic on a CPU is where you might use an atomic on a GPU.
shaklee3 5y ago
shared memory and L1 are the same on modern GPUs, but shared memory gives you control of what goes on there. I agree with dragontamer that it's absolutely essentially to have atomics and shared.
- WithinReason 5y ago
  How is it "absolutely essential" if you just want to sum 2 large arrays for example?
  - shaklee3 5y ago
    I'm talking about realistic kernels, which aren't that simple. sure, we can also say that multiplying two arrays doesn't need it. but everything from gemm to FFT does.
    - deleted 5y ago
      [deleted]
    - WithinReason 5y ago
      How are atomics absolutely essential for gemm?
      - shaklee3 5y ago
        we are talking about shared memory
        WithinReason 5y ago
        That makes less sense, since there is literally no algorithm that can't be implemented without shared memory.
        shaklee3 5y ago
        I'm not sure you were reading this thread. obviously anything can be written without shared memory, but they will be much, much slower, and using a GPU becomes less appealing. the entire purpose of the article and project is that it's fast, but it can't be anywhere near as fast as most cuda apps until it supports shared memory (not worth arguing about atomics).

> Atomics are absolutely essential Most problems don't need atomics to solve > If this is missing __shared__ memory concepts, then the entire idea of "efficie

dragontamer 5y ago

> Most problems don't need atomics to solve How do you handle a global, concurrent memory write and/or read ?? (across many different blocks, maybe even across different grids). For example: lets say you have a global hash table and 30,000 CUDA-threads are running. How do you insert data into the hash table safely? > Sure, shared memory is great, but not always needed and communicating through global memory can be fast if you stay inside the L1/L2 cache and hide latency. Scan operations (sum, min, max, AND, OR to name a few) through the parallel prefix pattern (https://en.wikipedia.org/wiki/Prefix_sum https://en.wikipedia.org/wiki/Prefix_sum). How do you plan to do it, if not through __shared__ memory? This is a fundamental operation in almost all GPU code I've ever seen. Just pushing / popping to a stack will require a prefix-sum to determine the size across the workgroup. If you can't do prefix-sum, you won't be able to do effective load-balancing on a GPU. This is something that'd normally take a dozen clock ticks, but if you do it over L2 you're looking at hundreds of clock ticks instead. ------ Sorting networks are also probably best implemented in __shared__ memory... with maybe warp-level intrinsics beating them. (But warp-level programming is much much harder and I prefer to avoid it).

floatboth 5y ago
Global hash table on GPUs sounds cursed and perverted. What you’re meant to do with them is run the same computation on all pixels in parallel independently! :P
WithinReason 5y ago
As I said, these features are often not needed. You can implement e.g. a neural network library without needing atomic operations. > How do you plan to do it, if not through __shared__ memory? Can't you use __shared__ memory the same way you use workgroup barriers and global memory? Might be slower, but good caching should make it comparable, which should be the case of prefix sum (you read right after writing, so should get good cache hit probability).
- dragontamer 5y ago
  > As I said, these features are often not needed. Global parallel hash table is probably the fastest way to implement collision detection. Its pretty fundamental to manipulation of 3d space, be it graphics, physics or other such simulations. Which of course, run _GREAT_ on GPUs. -------------- I've written "Inner Join" on a GPU for fun. Yes, the SQL operator. Its pretty fast. Databases probably can run on GPUs and parallelize easily. But any database would need globally consistent reads/writes. Sure, GPUs have less RAM than a CPU, but GPU-RAM is way faster so that might actually be a net benefit if your data is between 200MB and 4GB in size. Use your imagination. Anywhere you'd use an atomic on a CPU is where you might use an atomic on a GPU.

shaklee3 5y ago

shared memory and L1 are the same on modern GPUs, but shared memory gives you control of what goes on there. I agree with dragontamer that it's absolutely essentially to have atomics and shared.

WithinReason 5y ago
How is it "absolutely essential" if you just want to sum 2 large arrays for example?
- shaklee3 5y ago
  I'm talking about realistic kernels, which aren't that simple. sure, we can also say that multiplying two arrays doesn't need it. but everything from gemm to FFT does.
  - deleted 5y ago
    [deleted]
  - WithinReason 5y ago
    How are atomics absolutely essential for gemm?
    - shaklee3 5y ago
      we are talking about shared memory
      - WithinReason 5y ago
        That makes less sense, since there is literally no algorithm that can't be implemented without shared memory.
        shaklee3 5y ago
        I'm not sure you were reading this thread. obviously anything can be written without shared memory, but they will be much, much slower, and using a GPU becomes less appealing. the entire purpose of the article and project is that it's fast, but it can't be anywhere near as fast as most cuda apps until it supports shared memory (not worth arguing about atomics).