perf(cuda): reduce redundant per-token synchronizations

[hai-pilgrim]

Summary

Removes unnecessary CUDA synchronization points in the inference pipeline to reduce per-token latency. Inspired by upstream PR ggml-org#20793 which demonstrated measurable gains by reducing redundant syncs between async operations and graph execution.

All changes are scheduling-only — no computation changes, bit-exact output preserved.

Sync points changed

1. ggml_cuda_mul_mat_id — remove stream sync after H2D copy (line ~2420)

• Before: cudaMemcpyAsync (H2D) followed by cudaStreamSynchronize, then get_rows_cuda kernel on the same stream
• After: Removed the sync
• Why safe: Both the async copy and the subsequent kernel are on the same CUDA stream. CUDA stream semantics guarantee that operations submitted to the same stream execute in order. The sync was blocking the host thread unnecessarily without providing any correctness benefit.

2. ggml_backend_cuda_split_buffer_set_tensor — deduplicate sync loop (lines ~951-953)

• Before: After issuing N async H2D copies (one per device) all on cudaStreamPerThread, a loop called cudaStreamSynchronize(cudaStreamPerThread) once per device count — syncing the same stream N times
• After: Single cudaStreamSynchronize(cudaStreamPerThread) call
• Why safe: cudaStreamPerThread is a per-thread implicit stream. All copies are submitted to the same stream regardless of target device, so one sync drains the entire stream. The N-1 extra syncs were no-ops.

3. ggml_backend_cuda_split_buffer_get_tensor — deduplicate sync loop (lines ~990-992)

• Before: Same pattern as feat: add --remap-developer-role flag for developer→system translation #2 but for D2H copies
• After: Single cudaStreamSynchronize(cudaStreamPerThread) call
• Why safe: Same reasoning as feat: add --remap-developer-role flag for developer→system translation #2.

What is NOT changed

• cudaDeviceSynchronize in ggml_cuda_set_peer_access — this is a rare state-transition path for multi-GPU peer access that correctly requires device-wide synchronization
• cudaStreamSynchronize after D2H copy of ids_host in mul_mat_id (line ~2399) — host code reads the copied data immediately, so the sync is required
• All buffer interface syncs (set_tensor, get_tensor, memset_tensor, cpy_tensor, clear) — these are called from the generic backend API which expects synchronous completion

Compatibility

• Single GPU (RTX 3090 / SM86): Fully compatible, primary beneficiary
• Multi-GPU / split tensors: Split buffer paths still sync correctly (single sync drains all pending copies)
• CUDA graphs: No changes to graph capture/launch path
• CPU/GPU hybrid: No changes to cross-backend copy or event-based synchronization

Test plan

• Run inference with a MoE model (exercises mul_mat_id path)
• Run inference with tensor splitting across GPUs (exercises split buffer paths)
• Compare perplexity output before/after to verify bit-exactness
• Benchmark t/s on RTX 3090 to measure latency improvement

🤖 Generated with Claude Code

[marksverdhei]

Closing — head branch was deleted from remote. Reopen with a fresh branch if the work is still relevant.