[QST] How to correctly Initialize a device-scope cuda::barrier?

[osayamenja]

Hey experts,

I want to initialize a cuda::barrier<cuda::thread_scope_device> object that would be visible to all threads in the grid.

Is the code below valid?

I have used this logic elsewhere without any issues whatsoever, but I want to ensure I am not invoking undefined behavior by doing so 😅

#include <cuda/barrier>
#include <cuda_runtime_api.h>

__global__ void foo(cuda::barrier<cuda::thread_scope_device>* devBar){
 // do something fancy...
 __syncthreads();
 if(!threadIdx.x) {
  devBar->arrive_and_wait();
 }
 __syncthreads();
 // do some more fancy stuff...
}
template<cuda::thread_scope t>
__host__ __forceinline__
void init_cuda_barrier(void* const& devBar, const long int& blocks, cudaStream_t stream) {
 static_assert(t == cuda::thread_scope_block || t == cuda::thread_scope_device || t == cuda::thread_scope_system);
 const auto hB = new cuda::barrier<t>{blocks};
 cudaMemcpyAsync(devBar, hB, sizeof(cuda::barrier<t>), cudaMemcpyHostToDevice, stream);
 cudaStreamSynchronize(stream);
 delete hB;
}

int main(){
 constexpr auto t = cuda::thread_scope_device;
 cuda::barrier<t>* devBar;
 auto stream = cudaStreamPerThread;
 // allocate global memory for the barrier object
 cudaMallocAsync(&devBar, sizeof(cuda::barrier<t>), stream);
 const auto blocks = 32; // arbitrarily chosen
 init_cuda_barrier<t>(devBar, blocks, stream);
 foo<<<blocks, 32, 0, stream>>>(devBar);
 cudaStreamSynchronize(stream);
 return 0;
}

[osayamenja]

cc @ericniebler @davebayer

[miscco]

@osayamenja I believe that this is potentially dangerous. Conventionally barriers are put into shared memory and then initialized with the first thread that passes them.

Depending on the constructors used, you could get some issues, because technically a barrier is not copy constructible so cudaMemcpyAsync is technically UB.

To be fully correct you would need to allocate the space on device and then run a kernel to initialize the barrier on device with a single thread.

[wmaxey]

@miscco I agree that a block scope barrier should never be initialized this way on host. I mentioned this over Slack when the discussion popped up there.

I think if this was constrained to just global memory I would be 'okay' with this pattern. System and device scope barriers should be trivially copyable (perhaps true only now and not a guarantee). As you said though, it would be best to have a kernel initialize the barrier, but I think if someone had a good enough reason to avoid the latency involved with spinning up a thread to initialize the barrier on device they could do this. It won't break now, but might in the future. We would have to rev the ABI if we ever decided to change layout again.

After thinking on it more, initializing device scoped barriers on host has an odor. This leaves just system scope for the initialization pattern presented here which is not very useful.

[osayamenja]

@wmaxey @miscco Thanks for the comments, experts! It would be very helpful to users if documentation had an example of the correct initialization pattern for device/system barriers. The current doc only addresses shared memory. I will adopt the recommended single-thread kernel approach.

I assume it's a single thread calling cuda::barrier::init on a barrier object allocated via cudaMalloc right?

[wmaxey]

Yep, that would be correct. That would be a good doc update, adding recommended initialization patterns.