CUDA: round, not lround, in ClusterFinderCUDA copy-out

lround(float) has no hardware path on device and compiles to a
software routine. Only photon threads run it, so the warp diverges
and the f32 9x9 kernel was 1.8x slower than f64 on the A100.
round(float) is a few instructions and rounds half away from zero,
like std::lround on the CPU path, so cluster values do not change.

A100, 9x9, streams=1, 20k frames: kernel 32.1 -> 14.7 us/frame
(f64: 17.5); FPS 31,200 -> 41,400, now limited by D2H. 3x3 kernel
8.0 -> 6.7 us/frame.

include/aare/cf/clusterfinder_kernel.cuh

@@ -10,7 +10,7 @@ namespace aare::cuda::detail {
 /// Stencil arithmetic precision. Set both aliases to double for an f64 build.
 /// float is safe here because the device pedestal stores centered moments
 /// against a host-computed baseline: f32 and f64 agree to 3e-7 in cluster
-/// count, and f32 cuts the 9x9 kernel time by ~40%.
+/// count, and the f32 kernel is faster (fewer registers, higher occupancy).
 using COMPUTE_TYPE = float;
 /// Device pedestal storage and running variance update. See COMPUTE_TYPE.
 using DEVICE_PED_TYPE = float;
@@ -279,8 +279,11 @@ __global__ void find_clusters_in_single_frame(
 #pragma unroll
         for (int ic = -col_radius; ic <= col_radius; ++ic) {
             COMPUTE_TYPE val = shmem[shmem_tid + ir * TILE_W + ic];
+            // round, not lround: lround(float) is emulated in software on
+            // device (~2x the 9x9 kernel time); round(float) is a few
+            // instructions. Both round half away from zero, like the CPU path.
             if constexpr (std::is_integral_v<CT>)
-                dst[idx] = static_cast<CT>(lround(val));
+                dst[idx] = static_cast<CT>(round(val));
             else
                 dst[idx] = static_cast<CT>(val);
             idx++;