From 104e28447b0af2c74cb75601093753e0b1e0fa23 Mon Sep 17 00:00:00 2001 From: kferjaoui Date: Thu, 29 Oct 2026 09:52:15 +0100 Subject: [PATCH] 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 | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/include/aare/cf/clusterfinder_kernel.cuh b/include/aare/cf/clusterfinder_kernel.cuh index ef7b43e4..028b093d 100644 --- a/include/aare/cf/clusterfinder_kernel.cuh +++ b/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) - dst[idx] = static_cast(lround(val)); + dst[idx] = static_cast(round(val)); else dst[idx] = static_cast(val); idx++;