Sorry for the month of silence, that’s on me. You posted two substantial updates and then had to ping to get an answer, after doing most of the experimental work yourself. Not the turnaround you should be getting.
Two corrections on the summary posted just above mine, since both points would send you the wrong way:
- A10 is not “effectively unusable” here. Your own Jun 28 table has A10 at 1,408 RPS at 8 threads against T4’s 858. That is the fastest configuration measured anywhere in this thread.
- CUDA Graphs are not off the table for dynamic shapes. There is a supported per-shape pattern, in section 3 below.
1. The single-thread row is the answer
| GPU |
1 thread, no graphs |
1 thread, with graphs |
Delta |
| A10 |
244 |
409 |
+68% |
| T4 |
268 |
250 |
-7% |
At one thread there are no sibling worker threads to convoy against, so lock contention cannot explain that A10 gap. The only thing CUDA Graphs take away at one thread is per-kernel launch cost. A10 gaining 68% from that means your A10 path was spending most of its wall clock launching kernels instead of running them. That is the definition of an enqueue-bound workload: each launch costs the CPU and driver roughly 5 to 15 microseconds, and when kernels are only a few microseconds of GPU work, launching becomes the bottleneck.
T4 losing 7% is the same fact seen from the other side. Its kernels run long enough to hide the launch cost already, so graphs buy nothing and the replay bookkeeping costs a little.
So A10 is not worse at concurrency. It finishes each kernel sooner, which means the same launch rate starves it first. That is also the honest answer to your Jun 22 point: the +10% A10 lead at one thread looked too small to matter, but it was small precisely because A10 was already launch-capped at one thread. Take the launches away and it is 409 against 250, much closer to the roughly 2x FP16 gap you were expecting from the spec sheet. I framed this in June as “T4 hides the CPU cost”, which pointed the right way but explained it badly.
The rest follows. Once you are launch-limited, extra host threads just add more callers to the same driver launch path, and that is where your pthread_rwlock_wrlock band comes from. The locks are a symptom of the launch rate, not a separate A10 defect. It also explains your single-submitter result: you removed the contention, which is why Nsight showed less locking, but the total number of launches never changed and you moved all pre and post-processing onto one thread. Same launch bill, worse CPU distribution.
2. Confirm it per network, one command
Before restructuring anything, get an objective read on which of your networks are actually enqueue-bound:
trtexec --loadEngine=detector.plan --noCudaGraph --shapes=images:1x3x640x640
trtexec enables CUDA graphs by default, so --noCudaGraph gives you the ungraphed path. Compare the reported Enqueue Time against GPU Compute Time. Enqueue close to or above compute means that network is enqueue-bound and section 3 applies. Compute dominating means that network was never your problem.
3. CUDA Graphs with dynamic shapes
The constraint is that a captured graph is pinned to one input size and one context state, not that dynamic shapes are excluded. The documented way around it is one execution context per captured graph, with the contexts sharing device memory through createExecutionContextWithoutDeviceMemory() so N contexts do not cost N times the activation memory.
Build a graph cache keyed on shape:
// once per distinct shape
IExecutionContext* ctx = engine->createExecutionContextWithoutDeviceMemory();
ctx->setDeviceMemoryV2(sharedActivations, sharedActivationsSize);
ctx->setInputShape("images", shape);
ctx->enqueueV3(stream); // flushes the deferred shape update
cudaGraph_t graph;
cudaGraphExec_t instance;
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
ctx->enqueueV3(stream);
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&instance, graph, 0);
graphCache[shapeKey] = instance;
Then cudaGraphLaunch on a cache hit, and only a first-seen shape pays capture. That bare enqueueV3() before cudaStreamBeginCapture is required rather than optional: TensorRT defers the shape-change work, and capturing before it is flushed records the wrong thing.
⚠️ The captured graph also records the activation memory address and the input and output buffer addresses. If you pool buffers per stream, every (shape, buffer set) pair needs its own captured graph. Getting this wrong gives you undefined behavior rather than a clean error.
If the detector’s shape space is too wide for a cache, letterbox or pad to a small set of canonical sizes first. Most detection pipelines already resize to a fixed input and only vary batch, and bucketing batch to powers of two gets you down to a handful of graphs.
4. When a network cannot be captured at all
For an enqueue-bound network that cannot be graphed, make each launch do more work: raise the batch per enqueueV3() call so kernel duration grows against that fixed 5 to 15 microsecond launch cost. The docs give this as the direct alternative to graphs for the same problem. Fewer and larger inferences beat more and smaller ones on A10 specifically, because A10 is the card that runs out of work first.
One measurement would settle the rest: the Enqueue Time and GPU Compute Time from that --noCudaGraph run on your dynamic-shape detector. If enqueue dominates, the graph cache is worth the bookkeeping. If compute dominates, that network was never the bottleneck and we should look at the pipeline around it instead.
I will put a synthetic dynamic-shape detector with a per-shape graph cache on an A10 here and post the numbers in this thread within the next week, so you have something to compare against rather than taking the pattern on faith.
Best, Atharva
