Traditional pose tracking pipelines have a dirty secret: they're bottlenecked by CPU-GPU data transfers, not the actual computation. I built a YOLOv8-Pose + ByteTrack tracker that achieves 420 FPS by keeping the entire pipeline on GPU - TensorRT inference, postprocessing, OKS-based association, and Kalman filtering. No CPU involvement except for final visualization. This is the story of how I adapted ByteTrack for pose estimation and made every operation GPU-native in C++/CUDA.
GitHub Repository: [github.com/naveedprojects/yolo-pose-cpp](https://github.com/naveedprojects/yolo-pose-cpp)
The Problem: Death by Data Transfer
Here's what a typical pose tracking pipeline looks like:
The GPU finishes in 2.1ms, then sits idle while the CPU does busywork. The D2H/H2D copies alone cost 1.1ms - that's 12% of frame time wasted on memory transfers. And the CPU operations? They're sequential, cache-unfriendly, and don't scale.
My goal: eliminate every CPU operation except the final frame display.
The Results: 420 FPS on RTX 3080 Ti
| Model | FP16 FPS | INT8 FPS | Detection (ms) | Track (ms) | Total |
|---|---|---|---|---|---|
| YOLOv8n-pose | 420 | 352 | 1.28 | 0.49 | 2.32 ms |
| YOLOv8s-pose | 327 | 409 | 1.87 | 0.56 | 2.98 ms |
| YOLOv8m-pose | 196 | 309 | 3.89 | 0.64 | 5.10 ms |
| YOLOv11m-pose | 188 | 296 | 4.02 | 0.71 | 5.35 ms |
The tracking overhead is constant ~0.5ms regardless of model size. Whether you're running the tiny YOLOv8n or the larger YOLOv8m, tracking adds the same cost. That's the power of GPU-native execution - the tracker scales with detection count, not model complexity.
Architecture: Zero-Copy Pipeline
The key insight: data should never leave the GPU until you need to display it. Here's the new pipeline:
Notice: one D2H copy at the very end, and only for the tracks we need to visualize. Track state, Kalman covariances, velocity estimates - all persist on GPU memory between frames.
Innovation #1: OKS Instead of IoU
Standard ByteTrack uses bounding box IoU (Intersection over Union) for association. This works for object tracking, but fails for pose tracking. Why?
I replaced IoU with OKS (Object Keypoint Similarity) - the standard metric from COCO pose evaluation. OKS considers individual keypoint positions:
// OKS computation kernel
__device__ float compute_oks(
const float* pose_a, // [17, 3] = x, y, confidence per keypoint
const float* pose_b, // [17, 3]
float scale // sqrt(bbox_area)
) {
// Per-keypoint sigmas from COCO (nose is precise, hip is loose)
constexpr float SIGMAS[17] = {
0.026f, // nose
0.025f, 0.025f, // eyes
0.035f, 0.035f, // ears
0.079f, 0.079f, // shoulders
0.072f, 0.072f, // elbows
0.062f, 0.062f, // wrists
0.107f, 0.107f, // hips
0.087f, 0.087f, // knees
0.089f, 0.089f // ankles
};
float oks_sum = 0.0f;
int visible_count = 0;
for (int k = 0; k < 17; k++) {
float conf_a = pose_a[k * 3 + 2];
float conf_b = pose_b[k * 3 + 2];
// Only count keypoints visible in BOTH poses
if (conf_a > 0.2f && conf_b > 0.2f) {
float dx = pose_a[k * 3 + 0] - pose_b[k * 3 + 0];
float dy = pose_a[k * 3 + 1] - pose_b[k * 3 + 1];
float d_sq = dx * dx + dy * dy;
float sigma = SIGMAS[k];
float denom = 2.0f * scale * scale * sigma * sigma;
oks_sum += expf(-d_sq / denom);
visible_count++;
}
}
return visible_count > 0 ? oks_sum / visible_count : 0.0f;
}The per-keypoint sigmas encode anatomical precision: nose position is precise (sigma=0.026), but hip position has natural variance (sigma=0.107). This makes OKS robust to annotation noise and natural pose variation.
Visibility-Masked Matching
Real-world poses have occlusions. Someone's arm might be hidden behind their body. The visibility mask handles this gracefully:
- •Both keypoints visible: Full contribution to OKS
- •One or both occluded: Skip this keypoint pair
- •Fallback: If <5 keypoints visible, use torso-only matching (shoulders + hips)
This means a person walking behind a desk (legs occluded) still matches correctly based on their visible upper body.
Innovation #2: Third-Order Kalman Filter
Standard trackers use a 4D state: (x, y, vx, vy) for the bounding box center. This assumes constant velocity - fine for cars on highways, terrible for humans.
I implemented a 136-dimensional Kalman state - 8 values per keypoint × 17 keypoints:
// State per keypoint: [x, y, vx, vy, ax, ay, jx, jy]
// Position + Velocity + Acceleration + Jerk
// Total state dimension: 17 keypoints × 8 = 136
__device__ void kalman_predict(
float* state, // [136] mean
float* covariance, // [136] diagonal variances
float dt // time delta (usually 1.0 for frame-to-frame)
) {
for (int k = 0; k < 17; k++) {
int base = k * 8;
float x = state[base + 0];
float y = state[base + 1];
float vx = state[base + 2];
float vy = state[base + 3];
float ax = state[base + 4];
float ay = state[base + 5];
float jx = state[base + 6];
float jy = state[base + 7];
// Third-order motion model (Taylor expansion)
state[base + 0] = x + vx*dt + 0.5f*ax*dt*dt + (1.0f/6.0f)*jx*dt*dt*dt;
state[base + 1] = y + vy*dt + 0.5f*ay*dt*dt + (1.0f/6.0f)*jy*dt*dt*dt;
state[base + 2] = vx + ax*dt + 0.5f*jx*dt*dt;
state[base + 3] = vy + ay*dt + 0.5f*jy*dt*dt;
state[base + 4] = ax * 0.9f + jx*dt; // Acceleration decays
state[base + 5] = ay * 0.9f + jy*dt;
state[base + 6] = jx * 0.85f; // Jerk decays faster
state[base + 7] = jy * 0.85f;
// Variance grows during prediction (process noise)
covariance[base + 0] += 1.0f; // Position uncertainty
covariance[base + 2] += 0.5f; // Velocity uncertainty
covariance[base + 4] += 0.25f; // Acceleration uncertainty
covariance[base + 6] += 0.1f; // Jerk uncertainty
}
}Why does this matter? Consider someone throwing a ball:
- •2nd-order (constant velocity): Arm continues at same speed - wrong prediction
- •3rd-order (with acceleration): Arm decelerates after release - better prediction
- •With jerk decay: Sudden direction changes are captured, then smoothed out
The 0.9 and 0.85 decay factors are crucial. They give the filter 'memory' - it expects acceleration and jerk to persist briefly, then fade. This handles the natural dynamics of human motion.
Innovation #3: Auction Algorithm (Not Hungarian)
The classic solution for assignment problems is the Hungarian algorithm - O(n³) complexity, perfectly optimal. It's also terrible for GPUs.
Hungarian requires sequential row/column operations with global synchronization. GPUs hate sequential. They want thousands of threads doing the same thing in parallel.
I used the Auction Algorithm instead - a parallel-friendly approach where tracks 'bid' on detections:
// Auction algorithm: tracks bid on detections
// Converges in O(n) iterations, each iteration is O(n) parallel work
__global__ void auction_bid_kernel(
const float* cost_matrix, // [num_tracks, num_detections]
float* prices, // [num_detections] - current "price" of each detection
int* assignments, // [num_tracks] - current assignment (-1 = unassigned)
int* bids, // [num_detections] - which track is bidding
float* bid_values, // [num_detections] - bid amount
int num_tracks,
int num_detections,
float epsilon // Minimum bid increment
) {
int track_id = blockIdx.x * blockDim.x + threadIdx.x;
if (track_id >= num_tracks) return;
if (assignments[track_id] >= 0) return; // Already assigned
// Find best and second-best detection for this track
float best_value = -INFINITY;
float second_best = -INFINITY;
int best_det = -1;
for (int d = 0; d < num_detections; d++) {
float value = -cost_matrix[track_id * num_detections + d] - prices[d];
if (value > best_value) {
second_best = best_value;
best_value = value;
best_det = d;
} else if (value > second_best) {
second_best = value;
}
}
if (best_det >= 0 && best_value > -0.5f) { // Threshold for valid match
// Bid = value difference + epsilon (ensures progress)
float bid = best_value - second_best + epsilon;
// Atomic: highest bidder wins
atomicMax((int*)&bid_values[best_det], __float_as_int(bid));
if (bid_values[best_det] == __float_as_int(bid)) {
bids[best_det] = track_id;
}
}
}
// Main loop: bid, assign, repeat until convergence
void auction_solve(/* ... */) {
for (int iter = 0; iter < max_iters; iter++) {
auction_bid_kernel<<<grid, block>>>(cost, prices, assign, bids, vals, T, D, eps);
auction_assign_kernel<<<grid2, block>>>(bids, vals, prices, assign, D);
// Check convergence (all tracks assigned or no valid matches)
if (check_convergence(assign, num_tracks)) break;
}
// Typically converges in 5-10 iterations
}The auction algorithm is approximate - it might not find the globally optimal assignment. But in practice, it finds solutions within 1-2% of optimal, and it's 10x faster on GPU than Hungarian on CPU.
Innovation #4: Spatial Gating
Computing OKS for every track-detection pair is O(T × D × 17) - expensive. With 50 tracks and 30 detections, that's 25,500 keypoint distance calculations per frame.
But most pairs are obviously wrong - a track in the left corner won't match a detection on the right. Spatial gating pre-filters impossible matches:
__global__ void compute_gated_oks_matrix(
const float* track_poses, // [num_tracks, 17, 3]
const float* detection_poses, // [num_detections, 17, 3]
float* cost_matrix, // [num_tracks, num_detections]
int num_tracks,
int num_detections
) {
int t = blockIdx.x;
int d = threadIdx.x;
if (t >= num_tracks || d >= num_detections) return;
// Quick bounding box center distance check
float track_cx = (track_poses[t * 51 + 5*3] + track_poses[t * 51 + 6*3]) / 2; // shoulder midpoint
float track_cy = (track_poses[t * 51 + 5*3+1] + track_poses[t * 51 + 6*3+1]) / 2;
float det_cx = (detection_poses[d * 51 + 5*3] + detection_poses[d * 51 + 6*3]) / 2;
float det_cy = (detection_poses[d * 51 + 5*3+1] + detection_poses[d * 51 + 6*3+1]) / 2;
float pose_height = estimate_pose_height(&track_poses[t * 51]);
float center_dist = sqrtf((track_cx - det_cx)*(track_cx - det_cx) +
(track_cy - det_cy)*(track_cy - det_cy));
// Gate: if centers are > 3x pose height apart, skip OKS computation
if (center_dist > 3.0f * pose_height) {
cost_matrix[t * num_detections + d] = 1.0f; // Max cost = no match
return;
}
// Only compute expensive OKS for plausible pairs
float oks = compute_oks(&track_poses[t * 51], &detection_poses[d * 51], pose_height);
cost_matrix[t * num_detections + d] = 1.0f - oks; // Convert similarity to cost
}In practice, spatial gating reduces the OKS computations from ~1500 to ~200 per frame - a 7.5x reduction with zero loss in matching quality.
Two-Tier Association (ByteTrack Style)
ByteTrack's key insight: don't throw away low-confidence detections. They might be partially occluded people who were tracked in previous frames.
My implementation uses two tiers:
Tier 1: High-Confidence Association (conf > 0.30)
─────────────────────────────────────────────────
1. Compute OKS cost matrix: active_tracks × high_conf_detections
2. Run auction algorithm
3. Matched tracks: Update with Kalman filter, increment hit_count
4. Unmatched tracks: Move to Tier 2
5. Unmatched detections: Initialize as new tracks
Tier 2: Low-Confidence Recovery (0.15 < conf < 0.30)
─────────────────────────────────────────────────
1. Take unmatched tracks from Tier 1
2. Compute OKS cost matrix: unmatched_tracks × low_conf_detections
3. Run auction algorithm with stricter threshold
4. Matched tracks: Update but don't increment hit_count
5. Still unmatched: Increment lost_count, keep predicting
Track Lifecycle:
─────────────────────────────────────────────────
- New detection → Tentative track (hit_count = 1)
- hit_count >= 3 → Confirmed track (visible in output)
- lost_count >= 10 → Delete trackThis handles the common scenario where someone walks behind an obstacle. Their detection confidence drops (partial visibility), but the low-confidence tier keeps the track alive until they reappear.
Persistent GPU State
The final piece: track state never leaves GPU memory. Every frame, we're updating device pointers, not copying data back and forth.
class GPUTracker {
private:
// All state lives on GPU - persists across frames
float* d_track_poses_; // [MAX_TRACKS, 17, 3] - current pose estimates
float* d_kalman_mean_; // [MAX_TRACKS, 136] - Kalman state
float* d_kalman_covariance_; // [MAX_TRACKS, 136] - diagonal covariance
int* d_track_ids_; // [MAX_TRACKS] - unique IDs
int* d_track_states_; // [MAX_TRACKS] - active/lost/tentative
int* d_hit_counts_; // [MAX_TRACKS] - consecutive detections
int* d_lost_counts_; // [MAX_TRACKS] - consecutive misses
// Workspace buffers (reused each frame)
float* d_cost_matrix_; // [MAX_TRACKS, MAX_DETECTIONS]
int* d_assignments_; // [MAX_TRACKS]
float* d_prices_; // [MAX_DETECTIONS]
// Pinned host memory for final output only
float* h_output_poses_; // Pinned for fast D2H
int* h_output_ids_;
public:
void update(const float* d_detections, int num_detections, cudaStream_t stream) {
// Everything happens on GPU in this stream
kalman_predict_kernel<<<...>>>(d_kalman_mean_, d_kalman_covariance_, ...);
compute_gated_oks_kernel<<<...>>>(d_track_poses_, d_detections, d_cost_matrix_, ...);
auction_solve_kernel<<<...>>>(d_cost_matrix_, d_assignments_, ...);
kalman_update_kernel<<<...>>>(d_kalman_mean_, d_detections, d_assignments_, ...);
update_track_states_kernel<<<...>>>(d_track_states_, d_hit_counts_, d_lost_counts_, ...);
// No D2H copies here!
}
std::vector<TrackOutput> getActiveTracks() {
// Only copy when user needs to display results
cudaMemcpyAsync(h_output_poses_, d_track_poses_, ..., cudaMemcpyDeviceToHost);
cudaMemcpyAsync(h_output_ids_, d_track_ids_, ..., cudaMemcpyDeviceToHost);
cudaStreamSynchronize(stream_);
// ... pack into vector
}
};The `getActiveTracks()` call is the only synchronization point - and it only happens when you want to render the visualization. If you're just logging track IDs or feeding into another GPU pipeline, you can skip it entirely.
INT8 Quantization: The Extra 20%
TensorRT supports INT8 quantization, but YOLO's early layers (Conv + SiLU activation) don't quantize well. My solution: partial quantization.
Layer Precision Strategy:
─────────────────────────────────────────────────
Layers 0-4: FP16 (stem and early convolutions)
Layers 5+: INT8 (backbone, neck, head)
Result on YOLOv8m-pose:
- FP16 only: 196 FPS, 24.1 MB engine
- Full INT8: Accuracy drops significantly
- Partial: 309 FPS, 14.2 MB engine (+58% speed, -41% size)The accuracy drop from partial INT8 is <0.5% mAP - negligible for real-time applications. The speed gain is substantial, especially for larger models where INT8 tensor cores really shine.
What I Learned
- •IoU is wrong for poses - OKS captures what actually matters: keypoint positions
- •Data transfer is the enemy - keep everything on GPU, copy only for visualization
- •Hungarian is CPU-optimal, not GPU-optimal - the Auction algorithm parallelizes beautifully
- •Third-order motion models matter - humans accelerate and jerk; constant velocity is a lie
- •Spatial gating is free performance - simple distance checks eliminate 85% of computation
- •Partial INT8 is better than full INT8 - protect the layers that need precision
Performance Breakdown
Where does the time actually go? Here's the profiled breakdown for YOLOv8n-pose at 420 FPS:
| Operation | Time (ms) | % of Frame |
|---|---|---|
| TensorRT inference | 1.28 | 55.2% |
| YOLO decode + NMS | 0.15 | 6.5% |
| OKS matrix (gated) | 0.12 | 5.2% |
| Auction assignment | 0.08 | 3.4% |
| Kalman predict+update | 0.14 | 6.0% |
| Track state management | 0.05 | 2.2% |
| D2H copy (final) | 0.05 | 2.2% |
| OpenCV rendering | 0.45 | 19.4% |
| Total | 2.32 | 100% |
The entire tracking pipeline (OKS + Auction + Kalman + state) is 0.39ms - less than the OpenCV rendering. If you're feeding into another GPU pipeline instead of displaying, you get that 0.45ms back too.
When to Use This
This architecture shines when:
- •High frame rates matter - sports analysis, real-time feedback systems
- •Multi-person scenarios - the GPU parallelism scales with person count
- •Pose matters, not just position - action recognition, gesture control, physical therapy
- •You're already on GPU - feeding into another CUDA pipeline (pose-to-action models)
It's overkill for simple single-person tracking at 30 FPS. But if you need to track 20+ people at 200+ FPS with accurate pose matching, this is the architecture.
Future Work
- •Re-ID features - CNN embeddings for identity matching across camera cuts
- •3D pose lifting - extend to 3D skeleton estimation with temporal consistency
- •Multi-camera fusion - track identities across overlapping camera views
- •Pose forecasting - predict future poses for latency compensation
The fastest code is the code that doesn't run. The second fastest is the code that runs on GPU without ever talking to CPU. This tracker does the latter.