[CI/Build] Fix topk histogram build on SM75 (#46550)
Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>
Mohammad Miadh Angkad committed
Jun 24, 2026 at 15:51 UTC
191826ec612dc6648b176ed4e94c3e6e551e9c09
1 file changed
+13
-4
csrc/libtorch_stable/topk_histogram_4096.cuh
+13
-4
@@ -97,13 +97,22 @@ __device__ __forceinline__ uint32_t warp_inclusive_sum(uint32_t lane,
97
}
98
99
// Returns the sum of a value across all 32 threads in the warp, and every
100
-// thread gets the same result SM90+ PTX instruction that does a hardware
101
-// warp-wide reduction in a single instruction w.r.t. warp::reduce_sum(), which
102
-// uses a __shfl_xor_sync butterfly tree (5 shuffles for 32 lanes)
100
+// thread gets the same result. SM80+ uses redux.sync.add.u32, a single PTX
101
+// instruction for hardware warp-wide reduction. Older targets use the
102
+// __shfl_xor_sync butterfly tree, like warp::reduce_sum() (5 shuffles for 32
103
+// lanes).
104
__device__ __forceinline__ uint32_t warp_reduce_sum_full(uint32_t v) {
105
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800)
106
uint32_t r;
107
asm("redux.sync.add.u32 %0, %1, 0xFFFFFFFF;" : "=r"(r) : "r"(v));
108
return r;
109
+#else
110
+ #pragma unroll
111
+ for (uint32_t mask = kWarpSize >> 1; mask > 0; mask >>= 1) {
112
+ v += __shfl_xor_sync(0xFFFFFFFF, v, mask);
113
+ }
114
+ return v;
115
+#endif
116
}
117
118
// ============================================================================
@@ -412,7 +421,7 @@ __device__ void histogram_4096_topk(const float* __restrict__ scores,
421
if (lane_id == kWarpSize - 1) smem->warp_sum[warp_id] = warp_inc;
422
__syncthreads();
423
415
- // Step 3: Inter-warp prefix via redux.sync
424
+ // Step 3: Inter-warp prefix across warp sums.
425
const auto tmp = smem->warp_sum[lane_id];
426
uint32_t prefix = warp_reduce_sum_full(
427
lane_id < warp_id ? tmp : 0); // sum of all prior warps