CUDA Multi-Process Service

NVIDIA's CUDA Multi-Process Service (MPS) enables concurrent kernel execution across multiple processes, improving GPU utilization compared to default time-slicing, according to a technical blog post demonstrating the technique on an Intel Core i9-9900K CPU and NVIDIA GeForce RTX 5080 GPU. The post provides commands to enable and disable MPS on a host machine and launch a PyTorch Docker container with MPS access, using an orchestrated example with a Triton kernel that utilizes a single streaming multiprocessor to show performance gains.

CUDA Multi-Process Service 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. 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. To enable CUDA MPS, please run the following commands on host machine. | 123456 | sudo nvidia-smi -i 0 -c EXCLUSIVE PROCESSexport CUDA VISIBLE DEVICES=0export CUDA MPS PIPE DIRECTORY=/tmp/nvidia-mpsexport CUDA MPS LOG DIRECTORY=/tmp/nvidia-mps-logmkdir -p /tmp/nvidia-mps /tmp/nvidia-mps-lognvidia-cuda-mps-control -d | To disable CUDA MPS, please run the following commands on host machine. | 123 | sudo bash -c 'echo quit | nvidia-cuda-mps-control'sudo nvidia-smi -i 0 -c DEFAULTsudo rm -rf /tmp/nvidia-mps/ /tmp/nvidia-mps-log/ | To launch a PyTorch Docker container that has CUDA MPS access, please run the following command. | 1234567 | 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. 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. | 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145 | python import argparseimport osimport timeimport torchimport torch.multiprocessing as mpimport tritonimport triton.language as tlos.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.jitdef 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. | 1234567891011 | bash $ 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 : 200Total Wall-clock Time : 66.1970 secAggregate System Throughput : 3.02 requests/sec------------------------------------------------------------ | When MPS is enabled, the throughput increases significantly to 17 requests per second, a roughly 6x improvement. | 1234567891011 | bash $ 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 : 200Total Wall-clock Time : 11.7000 secAggregate System Throughput : 17.09 requests/sec------------------------------------------------------------ | CUDA Multi-Process Service