Link Search Menu Expand Document


Defined in header <cuda/barrier>:

template <cuda::thread_scope Scope,
          typename CompletionFunction = /* unspecified */>
class barrier {
  // ...

  __host__ __device__
  friend void init(cuda::std::barrier* bar,
                   cuda::std::ptrdiff_t expected,
                   CompletionFunction cf = CompletionFunction{});

The friend function cuda::barrier::init may be used to initialize an cuda::barrier that has not been initialized.

When using libcu++ with NVCC, __shared__ cuda::barrier will not have its constructors run because __shared__ variables are not initialized. cuda::barrier::init should be use to properly initialize such a cuda::barrier.

An NVCC diagnostic warning about the ignored constructor will be emitted:

warning: dynamic initialization is not supported for a function-scope static
__shared__ variable within a __device__/__global__ function

It can be silenced using #pragma diag_suppress static_var_with_dynamic_init.


#include <cuda/barrier>

// Disables `pipeline_shared_state` initialization warning.
#pragma diag_suppress static_var_with_dynamic_init

__global__ void example_kernel() {
  __shared__ cuda::barrier<cuda::thread_scope_block> bar;
  init(&bar, 1);

See it on Godbolt