File size: 2,006 Bytes
8ae5fc5 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 | ---
grand_parent: Extended API
parent: Barriers
---
# `cuda::barrier::init`
Defined in header `<cuda/barrier>`:
```cuda
template <cuda::thread_scope Scope,
typename CompletionFunction = /* unspecified */>
class barrier {
public:
// ...
__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`.
## Example
```cuda
#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](https://godbolt.org/z/jG8se6Kd8){: .btn }
[`cuda::thread_scope`]: ./thread_scopes.md
[thread.barrier.class paragraph 12]: https://eel.is/c++draft/thread.barrier.class#12
[coalesced threads]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#coalesced-group-cg
[`concurrentManagedAccess` property]: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp_116f9619ccc85e93bc456b8c69c80e78b
[`hostNativeAtomicSupported` property]: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp_1ef82fd7d1d0413c7d6f33287e5b6306f
|