feat(benchmark): per-node cost and bottleneck attribution for a run

--benchmark <path> reports cumulative CPU and wall time per node and
names the node pacing the run. The pacing node is located from sampled
channel occupancy, not from time-in-node: backpressure inflates
time-in-node for everything downstream of the real bottleneck, so the
obvious measure names the victim rather than the cause.

Sampling starts with the network and stops before it is destroyed.
Channel fill is instantaneous and everything has drained by shutdown, so
a single read at the end reports an idle pipeline however congested it
was.

kill -USR1 dumps the table from a running or wedged process. Channel
occupancy identifies a stalled node -- full input, empty output --
without a debug build or a debugger, which is the difference between
diagnosing the AR-004 hang in seconds and reproducing it under gdb.

Two knobs this exposes for measurement rather than sets: SAE_CV_THREADS,
because OpenCV's TBB arena and KPN's thread-per-node are two schedulers
unaware of each other on the same cores; and SAE_CUDA_BLOCKING_SYNC,
because the default spin-wait held the embedder thread at 99.7% user
time while nvidia-powerd cut the GPU's clock from 1005 to 210 MHz.
Neither default changes until a measurement says it should.

TRACES: VR-015 | PR-004

src/backends/trt_backend.cpp

@@ -19,6 +19,7 @@
 
 #include <NvInfer.h>
 #include <cuda_runtime_api.h>
+#include <cstdlib>
 #include <opencv2/dnn.hpp>
 #include <opencv2/imgproc.hpp>
 
@@ -45,6 +46,40 @@ inline void check_cuda(cudaError_t e, const char* what) {
         throw CudaError(std::string(what) + ": " + cudaGetErrorString(e));
 }
 
+/// TRACES: VR-015 | PR-004
+/// Select how a thread waits for the GPU. Must run before the CUDA context is
+/// created, so every engine constructor calls it and the first one wins.
+///
+/// The default (`cudaDeviceScheduleAuto`) spin-waits: `cudaStreamSynchronize`
+/// burns the calling thread's CPU for the whole of the device's work. Measured
+/// here, the embedder thread sat at 99.7% *user* time with 0.5 s of system time
+/// across 183 s — i.e. no blocking syscalls at all — while the GPU ran flat out.
+///
+/// On this laptop that is not merely wasted CPU. `nvidia-powerd` arbitrates one
+/// power budget across CPU and GPU, and the GPU's ceiling was observed dropping
+/// from 20 W idle to 15 W under our load, with the SM clock *falling* from
+/// 1005 MHz to 210 MHz once work started. Spinning may therefore be buying
+/// watts away from the device the pipeline is actually waiting on.
+///
+/// SAE_CUDA_BLOCKING_SYNC=1 switches to a blocking wait so the A/B needs no
+/// rebuild. Default is unchanged until the measurement says otherwise.
+inline void configure_cuda_sync_once() {
+    static const bool done = [] {
+        const char* env = std::getenv("SAE_CUDA_BLOCKING_SYNC");
+        if (env && env[0] == '1') {
+            cudaError_t e = cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+            std::cerr << "[cuda] sync policy: BlockingSync"
+                      << (e == cudaSuccess ? "" : " (FAILED — context already created)")
+                      << "\n";
+        } else {
+            std::cerr << "[cuda] sync policy: default (spin) — "
+                         "set SAE_CUDA_BLOCKING_SYNC=1 to compare\n";
+        }
+        return true;
+    }();
+    (void)done;
+}
+
 class TrtLogger : public nvinfer1::ILogger {
 public:
     void log(Severity sev, const char* msg) noexcept override {
@@ -109,6 +144,7 @@ public:
                                       (output_is_fp16_ ? 2 : 4);
         check_cuda(cudaMalloc(&d_input_,  in_bytes),  "cudaMalloc input");
         check_cuda(cudaMalloc(&d_output_, out_bytes), "cudaMalloc output");
+        configure_cuda_sync_once();
         check_cuda(cudaStreamCreate(&stream_), "cudaStreamCreate");
 
         context_->setTensorAddress(input_name_.c_str(),  d_input_);
@@ -293,6 +329,7 @@ public:
             out_elem_counts_[oi] = count;
         }
 
+        configure_cuda_sync_once();
         check_cuda(cudaStreamCreate(&stream_), "cudaStreamCreate");
 
         std::cerr << "[TrtScrfd] loaded: " << engine_path
@@ -462,6 +499,7 @@ public:
         const std::size_t out_count = static_cast<std::size_t>(kWindow);
         check_cuda(cudaMalloc(&d_input_,  in_count  * 4), "cudaMalloc input");
         check_cuda(cudaMalloc(&d_output_, out_count * 4), "cudaMalloc output");
+        configure_cuda_sync_once();
         check_cuda(cudaStreamCreate(&stream_), "cudaStreamCreate");
         context_->setTensorAddress(input_name_.c_str(),  d_input_);
         context_->setTensorAddress(output_name_.c_str(), d_output_);