| 1 | /* |
| 2 | * Shared 4096-bin single-CTA TopK helpers. |
| 3 | */ |
| 4 | |
| 5 | #ifndef TOPK_HISTOGRAM_4096_CUH_ |
| 6 | #define TOPK_HISTOGRAM_4096_CUH_ |
| 7 | |
| 8 | #include <cuda_fp16.h> |
| 9 | #include <cuda_runtime.h> |
| 10 | #include <cstdint> |
| 11 | |
| 12 | namespace vllm { |
| 13 | namespace topk_histogram_4096 { |
| 14 | |
| 15 | constexpr uint32_t kBlockSize = 1024; |
| 16 | constexpr uint32_t RADIX = 256; |
| 17 | constexpr uint32_t kMaxTies = 1024; |
| 18 | static_assert(kMaxTies <= kBlockSize, |
| 19 | "tie_handle requires kMaxTies <= kBlockSize"); |
| 20 | constexpr uint32_t kWarpSize = 32; |
| 21 | constexpr uint32_t kNumWarps = kBlockSize / kWarpSize; |
| 22 | |
| 23 | // Register path |
| 24 | constexpr uint32_t kHist4096VecsPerThread = 4; |
| 25 | constexpr uint32_t kHist4096MaxLen = |
| 26 | kHist4096VecsPerThread * 4 * kBlockSize; // 16384 |
| 27 | |
| 28 | struct alignas(16) MatchBin { |
| 29 | uint32_t bin, above_count, equal_count; |
| 30 | }; |
| 31 | struct alignas(8) Tie { |
| 32 | uint32_t idx; |
| 33 | float score; |
| 34 | }; |
| 35 | |
| 36 | __device__ __forceinline__ void load_float4_predicated(const float* ptr, |
| 37 | int base, int seq_len, |
| 38 | float& v0, float& v1, |
| 39 | float& v2, float& v3) { |
| 40 | uint32_t r0, r1, r2, r3; |
| 41 | const int p0 = (base < seq_len); |
| 42 | const int p1 = (base + 1 < seq_len); |
| 43 | const int p2 = (base + 2 < seq_len); |
| 44 | const int p3 = (base + 3 < seq_len); |
| 45 | asm volatile( |
| 46 | "{\n" |
| 47 | " .reg .pred pr0, pr1, pr2, pr3;\n" |
| 48 | " setp.ne.u32 pr0, %4, 0;\n" |
| 49 | " setp.ne.u32 pr1, %5, 0;\n" |
| 50 | " setp.ne.u32 pr2, %6, 0;\n" |
| 51 | " setp.ne.u32 pr3, %7, 0;\n" |
| 52 | " mov.u32 %0, 0xFF800000;\n" |
| 53 | " mov.u32 %1, 0xFF800000;\n" |
| 54 | " mov.u32 %2, 0xFF800000;\n" |
| 55 | " mov.u32 %3, 0xFF800000;\n" |
| 56 | " @pr0 ld.global.cg.u32 %0, [%8];\n" |
| 57 | " @pr1 ld.global.cg.u32 %1, [%8+4];\n" |
| 58 | " @pr2 ld.global.cg.u32 %2, [%8+8];\n" |
| 59 | " @pr3 ld.global.cg.u32 %3, [%8+12];\n" |
| 60 | "}\n" |
| 61 | : "=r"(r0), "=r"(r1), "=r"(r2), "=r"(r3) |
| 62 | : "r"(p0), "r"(p1), "r"(p2), "r"(p3), "l"(ptr)); |
| 63 | v0 = __uint_as_float(r0); |
| 64 | v1 = __uint_as_float(r1); |
| 65 | v2 = __uint_as_float(r2); |
| 66 | v3 = __uint_as_float(r3); |
| 67 | } |
| 68 | |
| 69 | // converts the float32 score to a 32-bit ordered unsigned integer — the full |
| 70 | // precision key for radix sorting |
| 71 | __device__ __forceinline__ auto convert_to_uint32_v2(float x) -> uint32_t { |
| 72 | uint32_t bits = __float_as_uint(x); |
| 73 | return (bits & 0x80000000u) ? ~bits : (bits | 0x80000000u); |
| 74 | } |
| 75 | |
| 76 | // Converts each score to a 12-bit bin (FP16 sign-magnitude -> top 12 bits -> |
| 77 | // bin 0-4095) |
| 78 | template <uint32_t kBits> |
| 79 | __device__ __forceinline__ uint32_t extract_coarse_bin_N(float x) { |
| 80 | __half h = __float2half_rn(x); |
| 81 | uint16_t bits = __half_as_ushort(h); |
| 82 | uint16_t key = (bits & 0x8000) ? static_cast<uint16_t>(~bits) |
| 83 | : static_cast<uint16_t>(bits | 0x8000); |
| 84 | return key >> (16 - kBits); |
| 85 | } |
| 86 | |
| 87 | // running sum within each warp — thread 0 gets its own value, thread 1 gets |
| 88 | // thread 0 + thread 1, thread 2 gets threads 0+1+2, etc. |
| 89 | __device__ __forceinline__ uint32_t warp_inclusive_sum(uint32_t lane, |
| 90 | uint32_t v) { |
| 91 | #pragma unroll |
| 92 | for (uint32_t o = 1; o < 32; o *= 2) { |
| 93 | uint32_t n = __shfl_up_sync(0xFFFFFFFF, v, o); |
| 94 | if (lane >= o) v += n; |
| 95 | } |
| 96 | return v; |
| 97 | } |
| 98 | |
| 99 | // Returns the sum of a value across all 32 threads in the warp, and every |
| 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 | // ============================================================================ |
| 119 | // Tie refinement (single CTA): 4-round radix-256 topK on the full FP32 ordered |
| 120 | // key Each round narrows by 8 bits until ties are fully resolved |
| 121 | // ============================================================================ |
| 122 | |
| 123 | template <uint32_t TopK> |
| 124 | __device__ void tie_handle(const Tie* ties, uint32_t num_ties, |
| 125 | uint32_t num_above, int32_t* output, void* _smem) { |
| 126 | struct TS { |
| 127 | alignas(128) uint32_t counter; |
| 128 | alignas(128) MatchBin match; |
| 129 | uint32_t histogram[RADIX]; |
| 130 | uint32_t warp_sum[kNumWarps]; |
| 131 | }; |
| 132 | auto* s = static_cast<TS*>(_smem); |
| 133 | const auto tx = threadIdx.x; |
| 134 | const auto li = tx % kWarpSize, wi = tx / kWarpSize; |
| 135 | |
| 136 | // Each thread loads one tie element. |
| 137 | const bool has = tx < num_ties; |
| 138 | const auto tie = has ? ties[tx] : Tie{0, 0.0f}; |
| 139 | const uint32_t key = convert_to_uint32_v2(tie.score); |
| 140 | |
| 141 | bool active = has; // tracks whether this thread's tie is still a candidate. |
| 142 | uint32_t remain = |
| 143 | TopK - num_above; // decreases each round as ties are resolved. |
| 144 | uint32_t wpos = TopK; // wpos will hold the final output position. |
| 145 | s->counter = 0; |
| 146 | __syncthreads(); |
| 147 | |
| 148 | // The 4-round radix loop - each round narrows by 8 bits until ties are fully |
| 149 | // resolved |
| 150 | #pragma unroll |
| 151 | for (int r = 0; r < 4; r++) { |
| 152 | uint32_t sh = 24 - r * 8; // round 0: bits 31-24, round 1: 23-16, etc. |
| 153 | uint32_t bin = (key >> sh) & 0xFF; // this tie's 8-bit bin for this round |
| 154 | |
| 155 | // Step 1: Build 256-bin histogram. |
| 156 | if (tx < RADIX) s->histogram[tx] = 0; |
| 157 | __syncthreads(); |
| 158 | if (active) atomicAdd(&s->histogram[bin], 1); |
| 159 | __syncthreads(); |
| 160 | |
| 161 | // Step 2: Prefix scan to find threshold |
| 162 | uint32_t hv = 0, wi2 = 0; |
| 163 | if (tx < RADIX) { |
| 164 | hv = s->histogram[tx]; |
| 165 | wi2 = warp_inclusive_sum(li, hv); |
| 166 | if (li == kWarpSize - 1) s->warp_sum[wi] = wi2; |
| 167 | } |
| 168 | __syncthreads(); |
| 169 | |
| 170 | if (tx < RADIX) { |
| 171 | auto tmp = (li < RADIX / kWarpSize) ? s->warp_sum[li] : 0; |
| 172 | auto tot = warp_reduce_sum_full(tmp); |
| 173 | auto inter = warp_reduce_sum_full(li < wi ? tmp : 0); |
| 174 | auto above = tot - (inter + wi2); |
| 175 | if (above < remain && above + hv >= remain) { |
| 176 | s->match = {tx, above, remain - above}; |
| 177 | } |
| 178 | } |
| 179 | __syncthreads(); |
| 180 | |
| 181 | // Step 3: Scatter |
| 182 | auto [thr, na, _] = s->match; // threshold bin, num above, unused |
| 183 | if (active) { |
| 184 | if (bin > thr) { |
| 185 | wpos = num_above + |
| 186 | atomicAdd(&s->counter, 1); // above -> place in output directly |
| 187 | active = false; |
| 188 | } else if (bin < thr) |
| 189 | active = false; // below -> discard |
| 190 | else if (r == 3) |
| 191 | wpos = TopK - atomicAdd(&s->match.equal_count, |
| 192 | -1u); // last round: place remaining |
| 193 | } |
| 194 | remain -= na; |
| 195 | if (!remain) break; // all ties resolved early |
| 196 | } |
| 197 | // Final write |
| 198 | if (wpos < TopK) output[wpos] = tie.idx; |
| 199 | } |
| 200 | |
| 201 | // Extended tie_handle for TopK > kBlockSize (e.g. TopK=2048). |
| 202 | // tie_handle assumes 1 tie per thread (max 1024). |
| 203 | // This version handles 2 ties per thread via kPerThread=2 |
| 204 | template <uint32_t TopK> |
| 205 | __device__ void tie_handle_large(const Tie* ties, uint32_t num_ties, |
| 206 | uint32_t num_above, int32_t* output, |
| 207 | void* _smem) { |
| 208 | static_assert(TopK > kBlockSize); |
| 209 | struct TS { |
| 210 | alignas(128) uint32_t counter; |
| 211 | alignas(128) MatchBin match; |
| 212 | uint32_t histogram[RADIX]; |
| 213 | uint32_t warp_sum[kNumWarps]; |
| 214 | }; |
| 215 | auto* s = static_cast<TS*>(_smem); |
| 216 | const auto tx = threadIdx.x; |
| 217 | const auto li = tx % kWarpSize; |
| 218 | const auto wi = tx / kWarpSize; |
| 219 | |
| 220 | constexpr uint32_t kPerThread = (TopK + kBlockSize - 1) / kBlockSize; |
| 221 | Tie my_ties[kPerThread]; |
| 222 | uint32_t keys[kPerThread]; |
| 223 | bool active[kPerThread]; |
| 224 | |
| 225 | for (uint32_t e = 0; e < kPerThread; e++) { |
| 226 | uint32_t idx = e * kBlockSize + tx; |
| 227 | if (idx < num_ties) { |
| 228 | my_ties[e] = ties[idx]; |
| 229 | keys[e] = convert_to_uint32_v2(ties[idx].score); |
| 230 | active[e] = true; |
| 231 | } else { |
| 232 | my_ties[e] = {0, 0.0f}; |
| 233 | keys[e] = 0; |
| 234 | active[e] = false; |
| 235 | } |
| 236 | } |
| 237 | |
| 238 | uint32_t remain = TopK - num_above; |
| 239 | s->counter = 0; |
| 240 | __syncthreads(); |
| 241 | |
| 242 | for (int r = 0; r < 4; r++) { |
| 243 | uint32_t sh = 24 - r * 8; |
| 244 | if (tx < RADIX) { |
| 245 | s->histogram[tx] = 0; |
| 246 | } |
| 247 | __syncthreads(); |
| 248 | |
| 249 | for (uint32_t e = 0; e < kPerThread; e++) { |
| 250 | if (active[e]) { |
| 251 | atomicAdd(&s->histogram[(keys[e] >> sh) & 0xFF], 1); |
| 252 | } |
| 253 | } |
| 254 | __syncthreads(); |
| 255 | |
| 256 | uint32_t hv = 0; |
| 257 | if (tx < RADIX) { |
| 258 | hv = s->histogram[tx]; |
| 259 | auto wi2 = warp_inclusive_sum(li, hv); |
| 260 | if (li == kWarpSize - 1) { |
| 261 | s->warp_sum[wi] = wi2; |
| 262 | } |
| 263 | } |
| 264 | __syncthreads(); |
| 265 | if (tx < RADIX) { |
| 266 | auto tmp2 = (li < RADIX / kWarpSize) ? s->warp_sum[li] : 0; |
| 267 | auto total = warp_reduce_sum_full(tmp2); |
| 268 | auto inter = warp_reduce_sum_full(li < wi ? tmp2 : 0); |
| 269 | auto wi2 = warp_inclusive_sum(li, hv); |
| 270 | auto above = total - (inter + wi2); |
| 271 | if (above < remain && above + hv >= remain) { |
| 272 | s->match = { |
| 273 | .bin = tx, .above_count = above, .equal_count = remain - above}; |
| 274 | } |
| 275 | } |
| 276 | __syncthreads(); |
| 277 | |
| 278 | auto thr = s->match.bin; |
| 279 | auto na = s->match.above_count; |
| 280 | |
| 281 | for (uint32_t e = 0; e < kPerThread; e++) { |
| 282 | if (!active[e]) { |
| 283 | continue; |
| 284 | } |
| 285 | uint32_t bin = (keys[e] >> sh) & 0xFF; |
| 286 | if (bin > thr) { |
| 287 | uint32_t wpos = num_above + atomicAdd(&s->counter, 1); |
| 288 | if (wpos < TopK) { |
| 289 | output[wpos] = my_ties[e].idx; |
| 290 | } |
| 291 | active[e] = false; |
| 292 | } else if (bin < thr) { |
| 293 | active[e] = false; |
| 294 | } else if (r == 3) { |
| 295 | uint32_t wpos = TopK - atomicAdd(&s->match.equal_count, -1u); |
| 296 | if (wpos < TopK) { |
| 297 | output[wpos] = my_ties[e].idx; |
| 298 | } |
| 299 | } |
| 300 | } |
| 301 | |
| 302 | num_above += na; |
| 303 | remain -= na; |
| 304 | __syncthreads(); |
| 305 | s->counter = 0; |
| 306 | __syncthreads(); |
| 307 | } |
| 308 | } |
| 309 | |
| 310 | // ============================================================================ |
| 311 | // Register-based single-CTA fast path for seq_len <= 16384 |
| 312 | // 4 float4 per thread × 1024 threads = 16384 elements max |
| 313 | // Uses 4096-bin (12-bit) histogram for better precision |
| 314 | // ============================================================================ |
| 315 | |
| 316 | template <uint32_t TopK, uint32_t HIST_BITS> |
| 317 | struct Histogram4096Smem { |
| 318 | static constexpr uint32_t HIST_BINS = 1 << HIST_BITS; |
| 319 | static constexpr uint32_t TIE_CAPACITY = TopK > kMaxTies ? TopK : kMaxTies; |
| 320 | alignas(128) uint32_t counter_gt; |
| 321 | alignas(128) uint32_t counter_eq; |
| 322 | MatchBin match; |
| 323 | uint32_t warp_sum[kNumWarps]; |
| 324 | union { |
| 325 | uint32_t histogram[HIST_BINS]; |
| 326 | Tie tie_buffer[TIE_CAPACITY]; |
| 327 | }; |
| 328 | }; |
| 329 | |
| 330 | template <uint32_t TopK, uint32_t HIST_BITS, |
| 331 | uint32_t VECS_PER_THREAD = kHist4096VecsPerThread, |
| 332 | bool UsePredicatedLoads = false> |
| 333 | __device__ void histogram_4096_topk(const float* __restrict__ scores, |
| 334 | int32_t* __restrict__ output, |
| 335 | uint32_t length, void* _smem) { |
| 336 | constexpr uint32_t HIST_BINS = 1 << HIST_BITS; |
| 337 | constexpr uint32_t ITEMS_PER_THREAD = HIST_BINS / kBlockSize; |
| 338 | static_assert(HIST_BINS >= kBlockSize, |
| 339 | "HIST_BITS must give >= kBlockSize bins"); |
| 340 | |
| 341 | using Smem = Histogram4096Smem<TopK, HIST_BITS>; |
| 342 | auto* smem = static_cast<Smem*>(_smem); |
| 343 | const auto tx = threadIdx.x; |
| 344 | const auto lane_id = tx % kWarpSize; |
| 345 | const auto warp_id = tx / kWarpSize; |
| 346 | |
| 347 | // Phase 1: Load all data into RF + build histogram |
| 348 | float4 |
| 349 | vecs[VECS_PER_THREAD]; // 4 vectors x 4 floats = 16 elements per thread |
| 350 | if constexpr (ITEMS_PER_THREAD >= 4) { |
| 351 | // Zero the histogram (SMEM writes) |
| 352 | for (uint32_t i = 0; i < ITEMS_PER_THREAD / 4; i++) |
| 353 | reinterpret_cast<uint4*>( |
| 354 | smem->histogram)[tx * (ITEMS_PER_THREAD / 4) + i] = |
| 355 | make_uint4(0, 0, 0, 0); |
| 356 | } else { |
| 357 | if (tx < HIST_BINS) smem->histogram[tx] = 0; |
| 358 | } |
| 359 | if (tx == 0) { |
| 360 | smem->counter_gt = 0; |
| 361 | smem->counter_eq = 0; |
| 362 | } |
| 363 | if constexpr (UsePredicatedLoads) { |
| 364 | const bool row_aligned = (reinterpret_cast<uintptr_t>(scores) & 0xFu) == 0; |
| 365 | #pragma unroll |
| 366 | for (uint32_t v = 0; v < VECS_PER_THREAD; v++) { |
| 367 | const uint32_t base = (tx + v * kBlockSize) * 4; |
| 368 | if (base < length) { |
| 369 | if (row_aligned && base + 3 < length) { |
| 370 | vecs[v] = *reinterpret_cast<const float4*>(scores + base); |
| 371 | } else { |
| 372 | load_float4_predicated(scores + base, static_cast<int>(base), |
| 373 | static_cast<int>(length), vecs[v].x, vecs[v].y, |
| 374 | vecs[v].z, vecs[v].w); |
| 375 | } |
| 376 | } |
| 377 | } |
| 378 | } else { |
| 379 | #pragma unroll |
| 380 | for (uint32_t v = 0; v < VECS_PER_THREAD; v++) { |
| 381 | const uint32_t base = (tx + v * kBlockSize) * 4; |
| 382 | if (base < length) { |
| 383 | vecs[v] = *reinterpret_cast<const float4*>(scores + base); |
| 384 | } |
| 385 | } |
| 386 | } |
| 387 | __syncthreads(); |
| 388 | |
| 389 | // Build histogram from RF via atomic adds into the shared histogram |
| 390 | bool done = false; |
| 391 | #pragma unroll |
| 392 | for (uint32_t v = 0; v < VECS_PER_THREAD && !done; v++) { |
| 393 | const float* elems = reinterpret_cast<const float*>(&vecs[v]); |
| 394 | #pragma unroll |
| 395 | for (uint32_t e = 0; e < 4 && !done; e++) { |
| 396 | const uint32_t idx = (tx + v * kBlockSize) * 4 + e; |
| 397 | if (idx >= length) { |
| 398 | done = true; |
| 399 | } else { |
| 400 | atomicAdd(&smem->histogram[extract_coarse_bin_N<HIST_BITS>(elems[e])], |
| 401 | 1); |
| 402 | } |
| 403 | } |
| 404 | } |
| 405 | __syncthreads(); |
| 406 | |
| 407 | // Phase 2: Prefix scan to find threshold bin |
| 408 | // Multi-element scan (4096 bins: 4 per thread) |
| 409 | uint32_t orig[ITEMS_PER_THREAD]; |
| 410 | uint32_t local_sum = 0; |
| 411 | |
| 412 | // Step 1: Each thread sums its 4 bins |
| 413 | #pragma unroll |
| 414 | for (uint32_t i = 0; i < ITEMS_PER_THREAD; i++) { |
| 415 | orig[i] = smem->histogram[tx * ITEMS_PER_THREAD + i]; |
| 416 | local_sum += orig[i]; |
| 417 | } |
| 418 | |
| 419 | // Step 2: Warp-level inclusive prefix sum on local_sum |
| 420 | const auto warp_inc = warp_inclusive_sum(lane_id, local_sum); |
| 421 | if (lane_id == kWarpSize - 1) smem->warp_sum[warp_id] = warp_inc; |
| 422 | __syncthreads(); |
| 423 | |
| 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 |
| 428 | prefix += |
| 429 | warp_inc - local_sum; // exclusive prefix within this thread's position |
| 430 | |
| 431 | // Step 4: Find threshold - scan 4 bins, accumulate prefix |
| 432 | #pragma unroll |
| 433 | for (uint32_t i = 0; i < ITEMS_PER_THREAD; i++) { |
| 434 | prefix += orig[i]; |
| 435 | const auto above = length - prefix; // elements in bins ABOVE this one |
| 436 | if (above < TopK && above + orig[i] >= TopK) { |
| 437 | smem->match = {.bin = tx * ITEMS_PER_THREAD + i, |
| 438 | .above_count = above, |
| 439 | .equal_count = orig[i]}; |
| 440 | } |
| 441 | } |
| 442 | |
| 443 | __syncthreads(); |
| 444 | |
| 445 | // Phase 3: Scatter from registers |
| 446 | const auto [thr_bin, num_above, num_equal] = smem->match; |
| 447 | const bool need_tie = (num_equal + num_above > TopK); |
| 448 | |
| 449 | done = false; |
| 450 | #pragma unroll |
| 451 | for (uint32_t v = 0; v < VECS_PER_THREAD && !done; v++) { |
| 452 | const float* elems = reinterpret_cast<const float*>(&vecs[v]); |
| 453 | #pragma unroll |
| 454 | for (uint32_t e = 0; e < 4 && !done; e++) { |
| 455 | const uint32_t idx = (tx + v * kBlockSize) * 4 + e; |
| 456 | if (idx >= length) { |
| 457 | done = true; |
| 458 | } else { |
| 459 | const uint32_t bin = extract_coarse_bin_N<HIST_BITS>(elems[e]); |
| 460 | if (bin > thr_bin) { |
| 461 | output[atomicAdd(&smem->counter_gt, 1)] = |
| 462 | idx; // above -> output directly |
| 463 | } else if (bin == thr_bin) { |
| 464 | const auto pos = atomicAdd(&smem->counter_eq, 1); |
| 465 | if (!need_tie) { |
| 466 | if (pos + num_above < TopK) { |
| 467 | output[pos + num_above] = idx; // all fit |
| 468 | } |
| 469 | } else { |
| 470 | if (pos < TopK) { |
| 471 | smem->tie_buffer[pos] = {idx, elems[e]}; // store for refirement |
| 472 | } |
| 473 | } |
| 474 | } |
| 475 | // else: bin < thr_bin - discard (not in top-k) |
| 476 | } |
| 477 | } |
| 478 | } |
| 479 | |
| 480 | // Phase 4: Tie-breaking |
| 481 | if (!need_tie) return; |
| 482 | __syncthreads(); |
| 483 | |
| 484 | // Fast warp-ballot tie-breaking for small tie counts |
| 485 | const uint32_t num_ties = min(num_equal, static_cast<uint32_t>(TopK)); |
| 486 | const uint32_t topk_remain = |
| 487 | TopK - num_above; // pick exactly remaining elements to fill topK |
| 488 | |
| 489 | auto is_greater = [](const Tie& a, const Tie& b) { |
| 490 | return (a.score > b.score) || (a.score == b.score && a.idx < b.idx); |
| 491 | }; |
| 492 | |
| 493 | if (num_ties <= kWarpSize) { |
| 494 | // <=32 ties - Use warp ballot |
| 495 | // All-to-all comparison in one __ballot_sync. 32 ties x 32 warps = 1024 |
| 496 | // comparisons in one instruction per warp. O(1) work. |
| 497 | const auto lane_id = tx % kWarpSize; |
| 498 | const auto warp_id = tx / kWarpSize; |
| 499 | if (lane_id >= num_ties || warp_id >= num_ties) return; |
| 500 | const uint32_t mask = (1ull << num_ties) - 1u; |
| 501 | const auto tie = smem->tie_buffer[lane_id]; // each lane holds one tie |
| 502 | const auto target = |
| 503 | smem->tie_buffer[warp_id]; // each warp evaluates one candidate |
| 504 | const bool pred = |
| 505 | is_greater(tie, target); // compare all ties against target |
| 506 | const auto rank = static_cast<uint32_t>( |
| 507 | __popc(__ballot_sync(mask, pred))); // count how many are greater |
| 508 | if (lane_id == 0 && rank < topk_remain) { |
| 509 | output[num_above + rank] = target.idx; // place at correct position |
| 510 | } |
| 511 | } else if (num_ties <= |
| 512 | kWarpSize * |
| 513 | 2) { // TODO (roberto): try to refactor this with <=32 case |
| 514 | // Same idea but each thread handles 2 tie elements |
| 515 | const auto lane_id = tx % kWarpSize; |
| 516 | const auto warp_id = tx / kWarpSize; |
| 517 | const auto lane1 = lane_id + kWarpSize; |
| 518 | const auto warp1 = warp_id + kWarpSize; |
| 519 | const auto invalid = Tie{0xFFFFFFFF, -__FLT_MAX__}; |
| 520 | const auto tie0 = smem->tie_buffer[lane_id]; |
| 521 | const auto tie1 = lane1 < num_ties ? smem->tie_buffer[lane1] : invalid; |
| 522 | if (warp_id < num_ties) { |
| 523 | const auto target = smem->tie_buffer[warp_id]; |
| 524 | const auto r0 = |
| 525 | __popc(__ballot_sync(0xFFFFFFFF, is_greater(tie0, target))); |
| 526 | const auto r1 = |
| 527 | __popc(__ballot_sync(0xFFFFFFFF, is_greater(tie1, target))); |
| 528 | if (lane_id == 0 && r0 + r1 < topk_remain) |
| 529 | output[num_above + r0 + r1] = target.idx; |
| 530 | } |
| 531 | if (warp1 < num_ties) { |
| 532 | const auto target = smem->tie_buffer[warp1]; |
| 533 | const auto r0 = |
| 534 | __popc(__ballot_sync(0xFFFFFFFF, is_greater(tie0, target))); |
| 535 | const auto r1 = |
| 536 | __popc(__ballot_sync(0xFFFFFFFF, is_greater(tie1, target))); |
| 537 | if (lane_id == 0 && r0 + r1 < topk_remain) |
| 538 | output[num_above + r0 + r1] = target.idx; |
| 539 | } |
| 540 | } else { |
| 541 | // Large tie count: fall back to 4-round radix-256 sort |
| 542 | if constexpr (TopK <= kBlockSize) { |
| 543 | tie_handle<TopK>(smem->tie_buffer, num_ties, num_above, output, smem); |
| 544 | } else { |
| 545 | tie_handle_large<TopK>(smem->tie_buffer, num_ties, num_above, output, |
| 546 | smem); |
| 547 | } |
| 548 | } |
| 549 | } |
| 550 | |
| 551 | template <uint32_t TopK, uint32_t HIST_BITS, |
| 552 | uint32_t VECS_PER_THREAD = kHist4096VecsPerThread> |
| 553 | __device__ __noinline__ void histogram_4096_topk_predicated( |
| 554 | const float* __restrict__ scores, int32_t* __restrict__ output, |
| 555 | uint32_t length, void* _smem) { |
| 556 | histogram_4096_topk<TopK, HIST_BITS, VECS_PER_THREAD, true>(scores, output, |
| 557 | length, _smem); |
| 558 | } |
| 559 | |
| 560 | } // namespace topk_histogram_4096 |
| 561 | } // namespace vllm |
| 562 | |
| 563 | #endif // TOPK_HISTOGRAM_4096_CUH_ |