Document the relationship between pipeline::producer_acquire and pipeline_shared_state's StagesCount
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
Quoting @gonzalobg
> For example, we guarantee that here:
>
> ```c++
> __global__ void foo() {
> constexpr unsigned stages_count = 2;
> __shared__ cuda::pipeline_shared_state shared_state;
> auto pipeline = cuda::make_pipeline(group, &shared_state);
> pipeline.producer_acquire();
> ...
> pipeline.producer_commit();
> pipeline.producer_acquire(); // This acquire never blocks
> ...
> ```
>
> the second call to `producer_acquire` never blocks because the `pipeline` has a "capacity" for holding two stages.
>
> If we change the `StageCount` to `1`:
>
> ```c++
> __global__ void foo() {
> constexpr unsigned stages_count = 1;
> __shared__ cuda::pipeline_shared_state shared_state;
> auto pipeline = cuda::make_pipeline(group, &shared_state);
> pipeline.producer_acquire();
> ...
> pipeline.producer_commit();
> pipeline.producer_acquire(); // This acquire always blocks, deadlocks, ... ??
> ...
> ```
>
> AFAICT, the second `producer_acquire` call cannot ever make progress. and that's something that we also guarantee (either by saying that the behavior here is undefined, or defining it to be something else).
>
> Not sure where would be the right place to document this, but I think it should be in the libcudacxx docs.
Contributor guide
Assessment
This issue has not been assessed yet.