A CTA parked on a counted barrier cannot be preempted

Accelerated Computing CUDA CUDA Programming and Performance
1

If a CTA has some warps parked on a counted barrier while another participant never arrives, the block stays resident and cannot be preempted. Reclaiming that tenant then fails either way:

  • terminate_client does remove it — and every other process on the GPU dies with cudaErrorIllegalAddress. 4/4 untouched co-tenants, 3/3 runs.

  • A device-side assert(0) cannot be delivered into that context at all. cudaLaunchKernel returns 0, the kernel never runs, the tenant keeps spinning — it fails silently while reporting success.

The damaging form is specifically a blocking barrier with a thread-count operand: __syncthreads() and bar.arrive 0, N wedge just as hard and are reclaimed with no collateral at all.

Reproducer, controls and full measurements: GitHub - Yeosangho/mps-counted-barrier-preempt-repro: CUDA MPS: a CTA parked on a counted barrier cannot be preempted; reclaiming it poisons every co-tenant. · GitHub

docker run --rm -it --gpus '"device=0"' --ipc=host \
  -v "$PWD/scripts:/w" -w /w nvcr.io/nvidia/pytorch:25.10-py3 \
  bash -c './build.sh && ./run_preempt.sh counted terminate 2 0'

Measured on B200, driver 580.126.20, CUDA MPS with a private control daemon.

Is this something that can be fixed? As it stands there is no reclaim path that is both effective and safe: the safe one does nothing, and the documented graceful one takes the whole GPU down with it. If a fix is not possible, is there any supported way to reclaim a tenant whose kernel cannot be preempted?