File size: 2,753 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 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 |
---
grand_parent: Extended API
parent: Synchronization Primitives
---
# `cuda::pipeline_shared_state`
Defined in header `<cuda/pipeline>`:
```cuda
template <cuda::thread_scope Scope, cuda::std::uint8_t StagesCount>
class cuda::pipeline_shared_state {
public:
__host__ __device__
pipeline_shared_state();
~pipeline_shared_state() = default;
pipeline_shared_state(pipeline_shared_state const&) = delete;
pipeline_shared_state(pipeline_shared_state&&) = delete;
};
```
The class template `cuda::pipeline_shared_state` is a storage type used to
coordinate the threads participating in a `cuda::pipeline`.
## Template Parameters
| `Scope` | A [`cuda::thread_scope`] denoting a scope including all the threads participating in the `cuda::pipeline`. `Scope` cannot be `cuda::thread_scope_thread`. |
| `StagesCount` | The number of stages for the _pipeline_. |
## Member Functions
| [(constructor)] | Constructs a `cuda::pipeline_shared_state`. |
| (destructor) [implicitly declared] | Destroys the `cuda::pipeline_shared_state`. |
| `operator=` [deleted] | `cuda::pipeline_shared_state` is not assignable. |
## NVCC `__shared__` Initialization Warnings
When using libcu++ with NVCC, a `__shared__` `cuda::pipeline_shared_state` will
lead to the following warning because `__shared__` variables are not
initialized:
```
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/pipeline>
// Disables `pipeline_shared_state` initialization warning.
#pragma diag_suppress static_var_with_dynamic_init
__global__ void example_kernel(char* device_buffer, char* sysmem_buffer) {
// Allocate a 2 stage block scoped shared state in shared memory.
__shared__ cuda::pipeline_shared_state<cuda::thread_scope_block, 2> pss0;
// Allocate a 2 stage block scoped shared state in device memory.
auto* pss1 = new cuda::pipeline_shared_state<cuda::thread_scope_block, 2>;
// Construct a 2 stage device scoped shared state in device memory.
auto* pss2 =
new (device_buffer) cuda::pipeline_shared_state<cuda::thread_scope_device, 2>;
// Construct a 2 stage system scoped shared state in system memory.
auto* pss3 =
new (sysmem_buffer) cuda::pipeline_shared_state<cuda::thread_scope_system, 2>;
}
```
[See it on Godbolt](https://godbolt.org/z/M9ah7r1Yx){: .btn }
[`cuda::thread_scope`]: ../thread_scopes.md
[(constructor)]: ./pipeline_shared_state/constructor.md
|