CUDA Multi-Process Service

Introduction

CUDA time-slicing is a software-based GPU sharing mechanism that allows multiple workloads or containers to multiplex and interleave on a single physical NVIDIA GPU. It works by rapidly switching the execution context of the hardware between different applications over a timeframe, making it look like the processes are running concurrently. But at any time point, only one process is actually executing on the GPU. It is the default mode of operation for NVIDIA GPUs. Consequently, if a process consumes very little GPU resource for a very long time, other processes may experience significant delays in accessing the GPU, resulting in GPU underutilizations and poor application performances.

To address this issue, NVIDIA provides the CUDA Multi-Process Service (MPS), which allows multiple CUDA applications to share a GPU more efficiently by enabling concurrent kernel execution and faster context switching across different processes. In this blog post, I would like to demonstrate how to use CUDA MPS to improve GPU utilization and application performance using an orchestrated example.

CUDA Multi-Process Service

The example will be executed on a Linux operating system with an Intel Core i9-9900K CPU and an NVIDIA GeForce RTX 5080 GPU via a Docker container.

Enabling CUDA Multi-Process Service

To enable CUDA MPS, please run the following commands on host machine.

sudo nvidia-smi -i 0 -c EXCLUSIVE_PROCESS
export CUDA_VISIBLE_DEVICES=0
export CUDA_MPS_PIPE_DIRECTORY=/tmp/nvidia-mps
export CUDA_MPS_LOG_DIRECTORY=/tmp/nvidia-mps-log
mkdir -p /tmp/nvidia-mps /tmp/nvidia-mps-log
nvidia-cuda-mps-control -d

Disabling CUDA Multi-Process Service

To disable CUDA MPS, please run the following commands on host machine.

sudo bash -c 'echo quit | nvidia-cuda-mps-control'
sudo nvidia-smi -i 0 -c DEFAULT
sudo rm -rf /tmp/nvidia-mps/* /tmp/nvidia-mps-log/*

Running a PyTorch Docker Container with CUDA MPS Enabled

To launch a PyTorch Docker container that has CUDA MPS access, please run the following command.

docker run -it --rm --gpus all --ipc=host \
  --user $(id -u):$(id -g) \
  -v /tmp/nvidia-mps:/tmp/nvidia-mps \
  -v /tmp/nvidia-mps-log:/tmp/nvidia-mps-log \
  --ulimit memlock=-1 --ulimit stack=67108864 \
  -v $(pwd):/mnt -w /mnt \
  nvcr.io/nvidia/pytorch:26.07-py3

Then CUDA MPS access in the Docker container can be enabled or disabled from host machine.

CUDA Time-Slicing VS Multi-Process Service

In the following example, I orchestrated a CUDA kernel that only utilizes a single Streaming Multiprocessor (SM) on GPU. Normally, in a single-process multi-stream application, we can launch multiple such kernels currently being executed on multiple streams to maximize the utilization of SMs on GPU. However, in a multi-process single-stream application, due to time-slicing, only one kernel can be executed at a time, resulting in GPU underutilization. With CUDA MPS, the low-utilization kernels from multiple processes can execute concurrently, improving overall GPU utilization.

The orchestration in this example intends to maximize the effect of CUDA MPS over time-slicing. In a real-world application, it is rare to see such underutilization of GPU resources and consequently the effect of CUDA MPS can be much less pronounced.

mps_triton_balanced.py

import argparse
import os
import time
import torch
import torch.multiprocessing as mp
import triton
import triton.language as tl

os.environ["CUDA_MPS_PIPE_DIRECTORY"] = "/tmp/nvidia-mps"
os.environ["CUDA_MPS_LOG_DIRECTORY"] = "/tmp/nvidia-mps-log"
os.environ["PYTORCH_CUDA_ALLOC_CONF"] = "expandable_segments:True"

torch.set_num_threads(1)
os.environ["OMP_NUM_THREADS"] = "1"


@triton.jit
def single_sm_compute_kernel(
    output_ptr,
    N_ELEMENTS: tl.constexpr,
    N_LOOPS: tl.constexpr,
    BLOCK_SIZE: tl.constexpr,
):
    pid = tl.program_id(axis=0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    acc = tl.zeros((BLOCK_SIZE, ), dtype=tl.float32)

    for i in range(N_LOOPS):
        f_i = i.to(tl.float32)
        acc += tl.sin(acc + f_i) * tl.cos(acc - f_i)

    tl.store(output_ptr + offs, acc, mask=offs < N_ELEMENTS)


def check_mps_status():
    mps_control_pipe = "/tmp/nvidia-mps/control"
    if os.path.exists(mps_control_pipe):
        return "ACTIVE (MPS control pipe present)"
    return "INACTIVE (MPS control pipe missing)"


def worker_task(rank, num_requests, n_loops, results_queue, barrier):
    device = torch.device("cuda:0")
    torch.cuda.set_device(device)

    n_elements = 1024
    output = torch.zeros(n_elements, device=device, dtype=torch.float32)
    grid = (1, )  # Strictly 1 SM grid execution

    # Synchronize start across all workers
    barrier.wait()

    start_time = time.perf_counter()
    for _ in range(num_requests):
        single_sm_compute_kernel[grid](output,
                                       N_ELEMENTS=n_elements,
                                       N_LOOPS=n_loops,
                                       BLOCK_SIZE=1024)
    torch.cuda.synchronize()

    elapsed = time.perf_counter() - start_time
    results_queue.put((rank, num_requests, elapsed))


def main():
    parser = argparse.ArgumentParser(
        description=
        "Single-SM Triton Kernel Concurrency Benchmark for NVIDIA MPS")
    parser.add_argument(
        "--workers",
        type=int,
        default=16,
        help="Number of parallel multiprocessing client workers")
    parser.add_argument(
        "--requests",
        type=int,
        default=50,
        help=
        "Number of sequential kernel invocation requests per worker process")
    parser.add_argument(
        "--loops",
        type=int,
        default=160000,
        help="Number of internal arithmetic compute loops per kernel invocation"
    )
    args = parser.parse_args()

    print("=" * 60)
    print(f"NVIDIA MPS Status : {check_mps_status()}")
    print("=" * 60)

    try:
        mp.set_start_method('spawn', force=True)
    except RuntimeError:
        pass

    # PRE-WARM/COMPILE KERNEL IN MAIN TO PREVENT TRITON CACHE RACE CONDITIONS
    print("Pre-compiling Triton kernel in main process...")
    device = torch.device("cuda:0")
    dummy_output = torch.zeros(1024, device=device, dtype=torch.float32)
    single_sm_compute_kernel[(1, )](dummy_output,
                                    N_ELEMENTS=1024,
                                    N_LOOPS=100,
                                    BLOCK_SIZE=1024)
    torch.cuda.synchronize()
    del dummy_output
    torch.cuda.empty_cache()

    print(
        f"Launching Single-SM Concurrency Test: {args.workers} Workers, {args.requests} Requests Each, {args.loops} Loops"
    )

    results_queue = mp.Queue()
    barrier = mp.Barrier(args.workers)
    workers = []

    wall_start = time.perf_counter()
    for rank in range(args.workers):
        p = mp.Process(target=worker_task,
                       args=(rank, args.requests, args.loops, results_queue,
                             barrier))
        p.start()
        workers.append(p)

    total_reqs = 0
    for _ in range(args.workers):
        _, reqs, _ = results_queue.get()
        total_reqs += reqs

    for p in workers:
        p.join()

    wall_elapsed = time.perf_counter() - wall_start

    print("-" * 60)
    print(f"Total Client Requests Completed : {total_reqs}")
    print(f"Total Wall-clock Time           : {wall_elapsed:.4f} sec")
    print(
        f"Aggregate System Throughput     : {total_reqs / wall_elapsed:.2f} requests/sec"
    )
    print("-" * 60)


if __name__ == "__main__":
    main()

With 8 workers running jobs simultaneously, when MPS is disabled, the throughput is only 3 requests per second.

$ python mps_triton_balanced.py --workers 8 --requests 25 --loops 160000
============================================================
NVIDIA MPS Status : INACTIVE (MPS control pipe missing)
============================================================
Pre-compiling Triton kernel in main process...
Launching Single-SM Concurrency Test: 8 Workers, 25 Requests Each, 160000 Loops
------------------------------------------------------------
Total Client Requests Completed : 200
Total Wall-clock Time           : 66.1970 sec
Aggregate System Throughput     : 3.02 requests/sec
------------------------------------------------------------

When MPS is enabled, the throughput increases significantly to 17 requests per second, a roughly 6x improvement.

$ python mps_triton_balanced.py --workers 8 --requests 25 --loops 160000
============================================================
NVIDIA MPS Status : ACTIVE (MPS control pipe present)
============================================================
Pre-compiling Triton kernel in main process...
Launching Single-SM Concurrency Test: 8 Workers, 25 Requests Each, 160000 Loops
------------------------------------------------------------
Total Client Requests Completed : 200
Total Wall-clock Time           : 11.7000 sec
Aggregate System Throughput     : 17.09 requests/sec
------------------------------------------------------------

References

• NVIDIA MPS Documentation
• Scaling Small LLMs with NVIDIA MPS