1

In OpenCL I could use __local whenever I wanted to manipulate subgroup memory. Analogically CUDA has __shared__ keyword. Does Vulkan have something equivalent? I cannot see anything in the subgroup tutorial https://www.khronos.org/blog/vulkan-subgroup-tutorial although I do see that they mention shared memory, but they never actually explain how to initialize it.

alagris
  • 1,838
  • 16
  • 31
  • Does this answer your question? [Vulkan compute shader for parallel sum reduction](https://stackoverflow.com/questions/68586498/vulkan-compute-shader-for-parallel-sum-reduction) – krOoze Aug 22 '21 at 16:43

1 Answers1

1

Vulkan supports shared buffers. They are better described here

https://www.khronos.org/opengl/wiki/Compute_Shader#Shared_variables

An example of usage might look as follows

layout (local_size_x = 32) in;

layout(std430, set = 0, binding = 1) buffer SomeBuffer{
    int some_ints[];
};

shared int[32] shared_ints;

void main(){
   shared_ints[gl_LocalInvocationID.x] = some_ints[gl_GlobalInvocationID.x];
}
alagris
  • 1,838
  • 16
  • 31
  • The word is workgroup then? Workgroup - block of threads that gets dispatched, can share data through shared memory, synchronise with barriers; And subgroups are waves(AMD) or warps(nvidia), part of a workgroups, where all threads execute same instruction at a time. If you want to move data within subgroup with extension you use shuffle or broadcast. – user369070 Aug 24 '21 at 15:51
  • This terminology is so confusing. Do you perhaps know of some place that describes it in detail? – alagris Aug 24 '21 at 16:51
  • https://en.wikipedia.org/wiki/Single_instruction,_multiple_threads – user369070 Aug 25 '21 at 15:42