| | --- |
| | grand_parent: Extended API |
| | parent: Synchronization Primitives |
| | --- |
| | |
| | # `cuda::pipeline` |
| |
|
| | Defined in header `<cuda/pipeline>`: |
| |
|
| | ```cuda |
| | template <cuda::thread_scope Scope> |
| | class cuda::pipeline { |
| | public: |
| | pipeline() = delete; |
| | |
| | __host__ __device__ ~pipeline(); |
| | |
| | pipeline& operator=(pipeline const&) = delete; |
| | |
| | __host__ __device__ void producer_acquire(); |
| | |
| | __host__ __device__ void producer_commit(); |
| | |
| | __host__ __device__ void consumer_wait(); |
| | |
| | template <typename Rep, typename Period> |
| | __host__ __device__ bool consumer_wait_for( |
| | cuda::std::chrono::duration<Rep, Period> const& duration); |
| | |
| | template <typename Clock, typename Duration> |
| | __host__ __device__ |
| | bool consumer_wait_until( |
| | cuda::std::chrono::time_point<Clock, Duration> const& time_point); |
| | |
| | __host__ __device__ void consumer_release(); |
| | |
| | __host__ __device__ bool quit(); |
| | }; |
| | ``` |
| |
|
| | The class template `cuda::pipeline` provides a coordination mechanism which |
| | can sequence [asynchronous operations], such as [`cuda::memcpy_async`], into |
| | stages. |
| |
|
| | A thread interacts with a _pipeline stage_ using the following pattern: |
| | 1. Acquire the pipeline stage. |
| | 2. Commit some operations to the stage. |
| | 3. Wait for the previously committed operations to complete. |
| | 4. Release the pipeline stage. |
| |
|
| | For [`cuda::thread_scope`]s other than `cuda::thread_scope_thread`, a |
| | [`cuda::pipeline_shared_state`] is required to coordinate the participating |
| | threads. |
| |
|
| | _Pipelines_ can be either _unified_ or _partitioned_. |
| | In a _unified pipeline_, all the participating threads are both producers and |
| | consumers. |
| | In a _partitioned pipeline_, each participating thread is either a producer or |
| | a consumer. |
| |
|
| | ## Template Parameters |
| |
|
| | | `Scope` | The scope of threads participating in the _pipeline_. | |
| |
|
| | ## Member Functions |
| |
|
| | | (constructor) [deleted] | `cuda::pipeline` is not constructible. | |
| | | [(destructor)] | Destroys the `cuda::pipeline`. | |
| | | `operator=` [deleted] | `cuda::pipeline` is not assignable. | |
| | | [`producer_acquire`] | Blocks the current thread until the next _pipeline stage_ is available. | |
| | | [`producer_commit`] | Commits operations previously issued by the current thread to the current _pipeline stage_. | |
| | | [`consumer_wait`] | Blocks the current thread until all operations committed to the current _pipeline stage_ complete. | |
| | | [`consumer_wait_for`] | Blocks the current thread until all operations committed to the current _pipeline stage_ complete or after the specified timeout duration. | |
| | | [`consumer_wait_until`] | Blocks the current thread until all operations committed to the current _pipeline stage_ complete or until specified time point has been reached. | |
| | | [`consumer_release`] | Release the current _pipeline stage_. | |
| | | [`quit`] | Quits current thread's participation in the _pipeline_. | |
| |
|
| | ## Notes |
| |
|
| | A thread role cannot change during the lifetime of the pipeline object. |
| |
|
| | ## Example |
| |
|
| | ```cuda |
| | #include <cuda/pipeline> |
| | #include <cooperative_groups.h> |
| | |
| | // Disables `pipeline_shared_state` initialization warning. |
| | #pragma diag_suppress static_var_with_dynamic_init |
| | |
| | template <typename T> |
| | __device__ void compute(T* ptr); |
| | |
| | template <typename T> |
| | __global__ void example_kernel(T* global0, T* global1, cuda::std::size_t subset_count) { |
| | extern __shared__ T s[]; |
| | auto group = cooperative_groups::this_thread_block(); |
| | T* shared[2] = { s, s + 2 * group.size() }; |
| | |
| | // Create a pipeline. |
| | constexpr auto scope = cuda::thread_scope_block; |
| | constexpr auto stages_count = 2; |
| | __shared__ cuda::pipeline_shared_state<scope, stages_count> shared_state; |
| | auto pipeline = cuda::make_pipeline(group, &shared_state); |
| | |
| | // Prime the pipeline. |
| | pipeline.producer_acquire(); |
| | cuda::memcpy_async(group, shared[0], |
| | &global0[0], sizeof(T) * group.size(), pipeline); |
| | cuda::memcpy_async(group, shared[0] + group.size(), |
| | &global1[0], sizeof(T) * group.size(), pipeline); |
| | pipeline.producer_commit(); |
| | |
| | // Pipelined copy/compute. |
| | for (cuda::std::size_t subset = 1; subset < subset_count; ++subset) { |
| | pipeline.producer_acquire(); |
| | cuda::memcpy_async(group, shared[subset % 2], |
| | &global0[subset * group.size()], |
| | sizeof(T) * group.size(), pipeline); |
| | cuda::memcpy_async(group, shared[subset % 2] + group.size(), |
| | &global1[subset * group.size()], |
| | sizeof(T) * group.size(), pipeline); |
| | pipeline.producer_commit(); |
| | pipeline.consumer_wait(); |
| | compute(shared[(subset - 1) % 2]); |
| | pipeline.consumer_release(); |
| | } |
| | |
| | // Drain the pipeline. |
| | pipeline.consumer_wait(); |
| | compute(shared[(subset_count - 1) % 2]); |
| | pipeline.consumer_release(); |
| | } |
| | |
| | template void __global__ example_kernel<int>(int*, int*, cuda::std::size_t); |
| | ``` |
| |
|
| | [See it on Godbolt](https://godbolt.org/z/zc41bWvja){: .btn } |
| |
|
| |
|
| | [asynchronous operations]: ../asynchronous_operations.md |
| | [`cuda::memcpy_async`]: ../asynchronous_operations/memcpy_async.md |
| |
|
| | [`cuda::thread_scope`]: ../thread_scopes.md |
| | [`cuda::pipeline_shared_state`]: ./pipeline_shared_state.md |
| |
|
| | [(destructor)]: ./pipeline/destructor.md |
| | [`producer_acquire`]: ./pipeline/producer_acquire.md |
| | [`producer_commit`]: ./pipeline/producer_commit.md |
| | [`consumer_wait`]: ./pipeline/consumer_wait.md |
| | [`consumer_wait_for`]: ./pipeline/consumer_wait.md |
| | [`consumer_wait_until`]: ./pipeline/consumer_wait.md |
| | [`consumer_release`]: ./pipeline/consumer_release.md |
| | [`quit`]: ./pipeline/quit.md |
| |
|
| |
|
