[cudax] CUDA shared memory abstractions proposal

[davebayer]

Nowadays, if you want to create an object living in the shared memory, you can do something like:

__global__ void kernel()
{
  __shared__ int my_shared_obj;
  
  if (threadIdx.x == 0)
  {
    my_shared_obj = 0xbad;
  }
  __sync_threads();

  // ...
}

That works fine for trivially constructible/destructible types. If one wants to use a type that has non-trivial constructor/destructor, he must implement the construction/destruction himself. It makes sense, because objects in shared memory are shared by the whole CTA (block) and there is no simple mechanism to select which does the construction/destruction.

Until now.

This PR introduces several APIs to support managing object lifetime of object in shared memory. It introduces the cudax::static_shared<T> type that statically allocates the memory in shared memory and provides methods for constructing, destructing and obtaining the stored object.

__global__ void kernel()
{
  cudax::static_shared<int> my_shared_obj{0xbedder};
  __syncthreads();
}

The construction/destruction can be done either automatically or manually to allow low level control over when and by which thread is the construction/destruction executed. See the example to get better idea of what the API is capable of. Note, that construction/destruction are collective operations, that must be called by all threads in the block.

For manipulating with shared memory objects, this PR implements cudax::shared_memory_ptr, which is a non-owning pointer type, similar to observer_ptr. It could be used in our APIs that expect pointers to shared memory.

This PR is just a prototype. What is not currently implemented are arrays :)

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.