# CUDA Multi-Process Service > Source: > Published: 2026-09-08 05:31:09.615155+00:00 # 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