| name | cuda-performance |
| description | Use when reviewing, optimizing, or debugging CUDA performance, especially LHCb Allen trigger software or GPU-resident reconstruction code. |
CUDA Performance Skill
Measure before optimizing. A useful change is localized, validated, and compatible with the execution model.
Required Context
Identify GPU model, driver/CUDA, build, exact benchmark command, input dataset/event count, baseline throughput or latency, target kernel/algorithm, memory layout, stream behavior, correctness tolerance, and profiling evidence.
For Allen, also identify sequence, algorithm, build directory, HLT1 throughput impact, Allen-managed buffers, stream ownership, and whether the change adds allocations, synchronization, or external libraries.
Constraints
- Do not accidentally use the default stream.
- Do not add
cudaDeviceSynchronize() or device allocation in hot paths.
- Do not introduce unmanaged memory ownership or deployment-conflicting dependencies.
- Preserve Allen SoA conventions and existing event/block assumptions.
- Do not change physics semantics to gain speed.
Checklist
Check indexing, bounds, coalescing, shared memory, register pressure, occupancy, warp divergence, atomics, host-device transfers, launch count, hidden synchronization, temporary allocations, data layout conversions, and numerical precision changes.
Procedure
- Record baseline.
- Identify the dominant compute, memory, launch, transfer, or synchronization cost.
- Make one meaningful change.
- Validate against reference behavior with stated tolerances.
- Re-measure with the same command and input.
- Report delta and remaining bottleneck.
Output
Diagnosis:
- ...
Likely bottleneck:
- ...
Recommended change:
- ...
Allen compatibility:
- ...
Correctness validation:
- ...
Benchmark:
- GPU:
- Command:
- Input:
- Baseline:
- New result:
- Delta:
Risks:
- ...
