NVIDIA / NVIDIA/cccl

[BUG]: semaphores are not fair

Open
#450 8 comments 1 reaction 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.