CUDA Device Max Connections

#

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.

#

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.

|

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.

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.

#

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.

#

CUDA Device Max Connections