[BUG]: semaphores are not fair
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
### Is this a duplicate?
- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Something else
### Component
libcu++
### Describe the bug
libcu++ semaphores are not fair, causing the following reproducer to hang (https://cuda.godbolt.org/z/7vGxhaocj):
```c++
#include
#include
#include
struct semaphore_lock {
cuda::binary_semaphore s{1};
__device__ void lock() {
s.acquire();
}
__device__ void unlock() {
s.release();
}
};
__device__ semaphore_lock l{};
__device__ int mask = 0;
__global__ void reproducer() {
l.lock();
bool cont = false;
do {
l.unlock();
cont = atomicAdd_block(&mask, threadIdx.x) == 0;
l.lock();
} while (cont);
l.unlock();
}
```
The standard does not require them to be fair, but that's pretty low QoI, particularly in our platform in which high latency differences between threads to different memories can penalize some threads CAS forever.
### How to Reproduce
.
### Expected behavior
The expected behavior is for the above program to never hang.
We could achieve this by making the semaphores fair, such that the threads acquire the semaphore in FIFO order.
This can be implemented by adding a `ticket`, and having the threads take a ticket (atomic add) and only `acquire` the semaphore if there are resources available _and it is their turn_.
This would turn the lock implemented on top of the semaphore into a ticket lock, which will never hang.
### Reproduction link
_No response_
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Research direction
Start with the libcu++ cuda/semaphore component and run the provided CUDA Godbolt reproducer to observe the hang. Investigate the semaphore acquisition behavior and evaluate the proposed ticket-based FIFO approach. Done means the reproducer no longer hangs because semaphore acquisition is fair, with appropriate coverage added for this behavior.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc
- Issue type
- Bug
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 30/100