CUDA Device Max Connections

09-26-202609-26-2026 blog 10 minutes read (About 1514 words) visits

Introduction

To maximize GPU utilization, it is common to have multiple workers processing tasks concurrently on GPU. However, by default, the GPU hardware concurrency is limited, no matter how much software concurrency is implemented. As a consequence, GPU might still be under utilized, even if at the software level the concurrency appears high in the implementation. `CUDA_DEVICE_MAX_CONNECTIONS` is an environment variable that can be set to control the number of hardware concurrency on GPU.

In this blog post, I would like to quickly discuss the importance of setting CUDA\_DEVICE\_MAX\_CONNECTIONS for maximizing GPU concurrency and overall utilization.

CUDA Device Max Connections

In CUDA programming, a CUDA stream is an abstraction which allows the programmer to express a sequence of operations. The developer could create multiple streams to enable concurrent execution of different tasks on the GPU, thereby improving overall utilization and performance. CUDA kernels launched in different streams can run concurrently, subject to hardware limitations and resource availability, such as the number of available Streaming Multiprocessors. There is one key factor that the developer might overlook, which is the `CUDA_DEVICE_MAX_CONNECTIONS` environment variable that controls the maximum number of concurrent connections to the GPU. If this variable is not set appropriately, no matter how many CUDA streams are created, how lightweight the kernels are on each stream, the GPU concurrency will still be limited.

In the following example, we created 32 CUDA streams to run concurrent tasks on GPU. The inference performances are benchmarked and profiling traces are collected.

multi\_stream\_runner.cpp

| 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132 | #include "multi\stream.h"#include "cuda\_utils.h"#include "cuda\_worker.h"#include <chrono>#include <cstdlib>#include <iostream>#include <set>#include <string>#include <thread>#include <torch/torch.h>#include <vector>int run\_multi\_stream(int argc, char\\ argv){ char const\ max\_conn\_env = std::getenv("CUDA\_DEVICE\_MAX\_CONNECTIONS"); std::string max\_conn\_str = max\_conn\_env ? max\_conn\_env : "default"; std::string trace\_filename = "multi\_stream\_trace\_max\_conn\" \+ max\_conn\_str + ".json"; if (argc == 3 && std::string(argv\[1\]) == "--trace-file") { trace\_filename = argv\[2\]; } else if (argc != 1) { std::cerr << "Usage: " << argv\[0\] << " \[--trace-file <path-to-trace.json>\]" << std::endl; return EXIT\_FAILURE; } std::cout << "CUDA\_DEVICE\_MAX\_CONNECTIONS = " << (max\_conn\_env ? max\_conn\_env : "Not Set (Defaults to 8)") << std::endl; int const num\_threads = 32; int const queries\_per\_thread = 200; unsigned long long const cycles = 1000000ULL; int const total\_queries = num\_threads \ queries\_per\_thread; std::vector<cudaStream\_t> streams(num\_threads); for (int thread\_index = 0; thread\_index < num\_threads; ++thread\_index) { CHECK\_CUDA\_ERROR(cudaStreamCreateWithFlags(&streams\[thread\_index\], cudaStreamNonBlocking)); } at::ThreadLocalState tls\_state; std::cout << "\\n--- Phase 1: Measuring Pure System Throughput (QPS) ---" << std::endl; auto start\_time = std::chrono::high\_resolution\_clock::now(); std::vector<std::thread> benchmark\_workers; for (int thread\_index = 0; thread\_index < num\_threads; ++thread\_index) { benchmark\_workers.emplace\_back(enqueue\_delay\_kernels, streams\[thread\_index\], queries\_per\_thread, cycles, tls\_state); } for (auto& worker : benchmark\_workers) { worker.join(); } for (cudaStream\_t stream : streams) { CHECK\_CUDA\_ERROR(cudaStreamSynchronize(stream)); } auto end\_time = std::chrono::high\_resolution\_clock::now(); std::chrono::duration<double\> elapsed = end\_time - start\_time; double elapsed\_seconds = elapsed.count(); double throughput = total\_queries / elapsed\_seconds; std::cout << "Total Queries Processed : " << total\_queries << std::endl; std::cout << "Elapsed Time : " << elapsed\_seconds << " seconds" << std::endl; std::cout << "Pure System Throughput : " << throughput << " queries/second" << std::endl; std::cout << "\\n--- Phase 2: Collecting Profiling Trace ---" << std::endl; torch::autograd::profiler::ProfilerConfig profiler\_config( torch::autograd::profiler::ProfilerState::KINETO, /\report\_input\_shapes=\/false, /\profile\_memory=\/false, /\with\_stack=\/false, /\with\_flops=\/false, /\with\_modules=\/false); std::set<torch::autograd::profiler::ActivityType> activities = { torch::autograd::profiler::ActivityType::CUDA}; torch::autograd::profiler::prepareProfiler(profiler\_config, activities); torch::autograd::profiler::enableProfiler(profiler\_config, activities); std::vector<std::thread> profile\_workers; for (int thread\_index = 0; thread\_index < num\_threads; ++thread\_index) { profile\_workers.emplace\_back(enqueue\_delay\_kernels, streams\[thread\_index\], queries\_per\_thread, cycles, tls\_state); } for (auto& worker : profile\_workers) { worker.join(); } std::this\_thread::sleep\_for(std::chrono::milliseconds(500)); auto profiler\_result = torch::autograd::profiler::disableProfiler(); for (cudaStream\_t stream : streams) { CHECK\_CUDA\_ERROR(cudaStreamSynchronize(stream)); } if (profiler\_result) { profiler\_result->save(trace\_filename); std::cout << "Saved PyTorch Profiler trace to: " << trace\_filename << std::endl; } for (cudaStream\_t stream : streams) { CHECK\_CUDA\_ERROR(cudaStreamDestroy(stream)); } return 0;} | | ------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------ | ------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- |

By varying `CUDA_DEVICE_MAX_CONNECTIONS`, we can control the maximum number of concurrent connections to the GPU device, which affects the performance of multi-stream workloads.

| 12345678910111213141516171819202122232425262728293031323334353637383940414243444546474849505152535455565758596061626364656667686970 | $ for connections in 1 2 4 8 16 32; do CUDA\DEVICE\_MAX\_CONNECTIONS="$connections" ./build/multi\_stream \\ \--trace-file "build/max\_conn\${connections}.json"doneCUDA\_DEVICE\_MAX\_CONNECTIONS = 1\--- Phase 1: Measuring Pure System Throughput (QPS) ---Total Queries Processed : 6400Elapsed Time : 1.66418 secondsPure System Throughput : 3845.73 queries/second\--- Phase 2: Collecting Profiling Trace ---USDT:2026-09-24 03:17:03 553:553 SyncActivityProfilerHandler.cpp:39\] profiler\_startUSDT:2026-09-24 03:17:04 553:553 SyncActivityProfilerHandler.cpp:46\] profiler\_stopSaved PyTorch Profiler trace to: build/max\_conn\_1.jsonCUDA\_DEVICE\_MAX\_CONNECTIONS = 2\--- Phase 1: Measuring Pure System Throughput (QPS) ---Total Queries Processed : 6400Elapsed Time : 1.13147 secondsPure System Throughput : 5656.38 queries/second\--- Phase 2: Collecting Profiling Trace ---USDT:2026-09-24 03:17:06 626:626 SyncActivityProfilerHandler.cpp:39\] profiler\_startUSDT:2026-09-24 03:17:07 626:626 SyncActivityProfilerHandler.cpp:46\] profiler\_stopSaved PyTorch Profiler trace to: build/max\_conn\_2.jsonCUDA\_DEVICE\_MAX\_CONNECTIONS = 4\--- Phase 1: Measuring Pure System Throughput (QPS) ---Total Queries Processed : 6400Elapsed Time : 0.438797 secondsPure System Throughput : 14585.3 queries/second\--- Phase 2: Collecting Profiling Trace ---USDT:2026-09-24 03:17:08 699:699 SyncActivityProfilerHandler.cpp:39\] profiler\_startUSDT:2026-09-24 03:17:09 699:699 SyncActivityProfilerHandler.cpp:46\] profiler\_stopSaved PyTorch Profiler trace to: build/max\_conn\_4.jsonCUDA\_DEVICE\_MAX\_CONNECTIONS = 8\--- Phase 1: Measuring Pure System Throughput (QPS) ---Total Queries Processed : 6400Elapsed Time : 0.35048 secondsPure System Throughput : 18260.7 queries/second\--- Phase 2: Collecting Profiling Trace ---USDT:2026-09-24 03:17:11 772:772 SyncActivityProfilerHandler.cpp:39\] profiler\_startUSDT:2026-09-24 03:17:11 772:772 SyncActivityProfilerHandler.cpp:46\] profiler\_stopSaved PyTorch Profiler trace to: build/max\_conn\_8.jsonCUDA\_DEVICE\_MAX\_CONNECTIONS = 16\--- Phase 1: Measuring Pure System Throughput (QPS) ---Total Queries Processed : 6400Elapsed Time : 0.142692 secondsPure System Throughput : 44851.8 queries/second\--- Phase 2: Collecting Profiling Trace ---USDT:2026-09-24 03:17:12 845:845 SyncActivityProfilerHandler.cpp:39\] profiler\_startUSDT:2026-09-24 03:17:13 845:845 SyncActivityProfilerHandler.cpp:46\] profiler\_stopSaved PyTorch Profiler trace to: build/max\_conn\_16.jsonCUDA\_DEVICE\_MAX\_CONNECTIONS = 32\--- Phase 1: Measuring Pure System Throughput (QPS) ---Total Queries Processed : 6400Elapsed Time : 0.0756491 secondsPure System Throughput : 84601.2 queries/second\--- Phase 2: Collecting Profiling Trace ---USDT:2026-09-24 03:17:14 918:918 SyncActivityProfilerHandler.cpp:39\] profiler\_startUSDT:2026-09-24 03:17:15 918:918 SyncActivityProfilerHandler.cpp:46\] profiler\_stopSaved PyTorch Profiler trace to: build/max\_conn\_32.json | | ----------------------------------------------------------------------------------------------------------------------------------- | ---------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- |

The system throughputs benchmarked and the profiling traces collected for different values of `CUDA_DEVICE_MAX_CONNECTIONS` are summarized in the table below.

| CUDA\_DEVICE\_MAX\_CONNECTIONS | Number of CUDA Streams | System Throughput (QPS) | Perfetto Trace | | ------------------------------ | ---------------------- | ----------------------- | --------------------------------------------------------------------------------------------------------------------------------------------------------- | | 1 | 32 | 3,845.73 | Trace | | 2 | 32 | 5,656.38 | Trace | | 4 | 32 | 14,585.30 | Trace | | 8 | 32 | 18,260.70 | Trace | | 16 | 32 | 44,851.80 | Trace | | 32 | 32 | 84,601.20 | Trace |

We could see that the system throughput nearly doubles as `CUDA_DEVICE_MAX_CONNECTIONS` is doubled, indicating a strong correlation between the number of allowed CUDA connections and the overall system performance. By examining the Perfetto traces, we could see that despite the very lightweight kernel, there are lots of bubbles in CUDA stream which are not caused by CPU launch overhead, if `CUDA_DEVICE_MAX_CONNECTIONS` is not the same as the number of CUDA streams.

Technically, each CUDA stream is associated with a hardware queue on GPU, and the number of hardware queues is configured by the `CUDA_DEVICE_MAX_CONNECTIONS` environment variable. By default, `CUDA_DEVICE_MAX_CONNECTIONS` is set to `8`. Therefore, in our application, if `CUDA_DEVICE_MAX_CONNECTIONS` is not set, the system will be significantly underutilized.

We could check what hardware queue each CUDA stream is mapped to by examining the `stream` and the `channel` attributes of CUDA kernels. For example, in the Perfetto trace of `CUDA_DEVICE_MAX_CONNECTIONS=1`, all CUDA streams are mapped to the same hardware queue `0`.

Miscellaneous

AMD GPUs have similar concepts of hardware queues and stream-to-queue mapping, which can be controlled through environment variables specific to the ROCm platform. In the case of AMD GPUs, GPU\_MAX\_HW\_QUEUES specifies the maximum number of hardware queues available for mapping streams and `hsa_queue` is the hardware queue associated with a particular stream that can be checked from the Perfetto trace attributes. Note that the default value of `GPU_MAX_HW_QUEUES` is `4`, which means the maximum GPU concurrency is very limited unless this environment variable is increased.