[cuda] make LambdaRank gradient kernel bit-deterministic when deterministic=true

[maxwbuckley]

Summary

CUDALambdarankNDCG's gradient kernel scatters per-pair gradient/hessian contributions into shared memory via atomicAdd_block. Different runs interleave those atomics in different orders, and FP add is non-associative, so the per-slot sums (and the trees built from them) are not bit-identical run-to-run. Today the deterministic config flag is a no-op on CUDA — only the runtime warning fires.

This PR makes deterministic=true actually deliver bit-deterministic LambdaRank training on CUDA. The flag dispatches to two new kernels that:

1. Compute sum_lambdas via cub::BlockReduce::Sum (fixed reduction tree, deterministic across runs).
2. Parallelize over output slots. For slot s with rank r in the descending sort, only pairs (i, r) for i < r and (r, j) for j > r touch s; the thread iterates these in (i, r)-then-(r, j) order — mirroring CPU's flat (i, j) iteration for slot s's contributions — and accumulates into a score_t (float) variable. No atomics; the accumulation order is fixed by the loop, not by the hardware.

Verification

Run-to-run determinism on the existing 50-case parity sweep, with deterministic=True:

state	before	after
bit-identical run-to-run	48 / 50	50 / 50
non-deterministic	2 / 50 (lambdarank_500, lambdarank_1500 at ~6e-8 max|Δ|)	0 / 50

The deterministic and existing non-deterministic kernels agree to within 1e-8 on the same input, so existing CPU/CUDA divergence is unchanged.

Performance impact (RTX 5090, 100 rounds, gpu_use_dp=true)

config	det=False	det=True	slowdown
items=50, groups=10	0.98s	1.10s	+12%
items=200, groups=20	1.02s	1.41s	+38%
items=500, groups=10	0.99s	2.79s	+182%
items=1500, groups=4	1.05s	4.99s	+376%

Slowdown grows with items_per_query because Phase B is O(items × (items + truncation_level)) per query. Default truncation_level=30 keeps the per-slot inner loops small enough that this is acceptable as an opt-in mode.

Test plan

• Two parametrized regression tests in test_dual.py covering both the BitonicArgSort_1024 (items ≤ 1024) and BitonicArgSort_2048 (1024 < items ≤ 2048) code paths.
• Full 50-case CUDA-determinism sweep: 50/50 bit-identical run-to-run.
• CPU/CUDA parity sweep on lambdarank_500 / lambdarank_1500: divergence vs CPU unchanged from the existing non-deterministic kernel (within 1e-8).
• Independent reviewer to sanity-check the (i, r)/(r, j) slot decomposition matches CPU's flat (i, j) iteration order for slot s's contributions.

🤖 Generated with Claude Code

[BelixRogner]

Thanks Max — and Claude Code. Real solution to a real problem: deterministic=true was silently a no-op for LambdaRank on CUDA (warning fires but the kernel was still non-deterministic via atomicAdd_block).

I worked through the slot-decomposition argument and it checks out — the (i, r) loop bounded by min(r, num_items_i) and the (r, j) loop gated by r < num_items_i correctly mirror CPU's flat for i in range(num_items_i): for j in range(i+1, query_item_count) from the perspective of slot s's contributions. With a score_t (float) accumulator iterating in the same sequence as CPU, the per-slot sums are bit-identical to CPU's. The cub::BlockReduce for sum_lambdas is run-to-run deterministic on the GPU (fixed reduction tree), and the ~1 ulp it can differ from CPU only scales norm_factor — well below the existing CPU/CUDA divergence floor. Consider the "independent reviewer to sanity-check the slot decomposition" box checked.

Verification: 50/50 bit-identical run-to-run on the parity sweep, kernels agree to within 1e-8 on the same input. Performance tradeoff (12–376% slowdown depending on items_per_query) is fine since this is opt-in via deterministic=true.

Two small things before merge:

1. Lint — Static Analysis/lint fails on this PR:
   • cpplint: src/objective/cuda/cuda_rank_objective.cpp:19 Do not indent within a namespace. [whitespace/indent_namespace] — the constructor body after the : CUDALambdaRankObjectiveInterface<...>(config), deterministic_(config.deterministic) {} line.
   • ruff check: tests/python_package_test/test_dual.py:4 — I001 Import block is un-sorted — same pattern hitting your other open PRs. pre-commit run --all-files should fix both.

Once that's in, this is ready to merge.

[BelixRogner]

Quick rebase nudge — #7 and #8 just landed on master and touched files this PR also modifies (cuda_algorithms.hpp / cuda_best_split_finder.cu / etc.), so this branch now shows a merge conflict on GitHub. One more git merge master && git push round should clear it; CI is otherwise green (the apparent failures yesterday were all environmental — dask socket flakes, a cancelled job rolling up, and a Boost-headers wheel-build issue, none of which were touching the actual PR content).

Ready to merge as soon as the conflict's resolved.

[BelixRogner]

Thank you, Max — and thank you, Claude Code (separately 🙂). Friendly rebase nudge 👉

master has moved a long way since this branch — the collapse onto latest upstream plus a whole stack of CUDA merges this cycle (#17, #21, #22, #23/#26, #24, #25, #27) — so GitHub now shows this one CONFLICTING. The change itself reviewed well earlier; it's purely stale.

Could you rebase onto current master when you get a chance? And good news — you've got write access to the repo now (accept the invite at https://github.com/BelixRogner/ExaBoost/invitations if you haven't), so once it's rebased and green you can merge it yourself. 🚀

[maxwbuckley]

Rebased onto current master. 👋

Please re-review — the branch was rebuilt rather than replayed, so it isn't byte-for-byte what you reviewed before.

What changed. The old branch carried 23 commits, ~18 of which have since landed on master via other PRs. Because master was collapsed onto upstream, those commits survived under different SHAs, so a straight replay tried to re-apply them and conflicted. I rebuilt the branch as master + only the two commits that are actually this PR:

• src/objective/cuda/cuda_rank_objective.{cpp,cu,hpp}
• tests/python_package_test/test_dual.py

(10 files → 4, 23 commits → 2.) The LambdaRank determinism change itself is unchanged.

Validation (RTX 5090, CUDA 13.2, built from source): full test_dual.py = 121 passed, 1 skipped. cpplint clean on all three changed C++ files with the repo's filters; no new compiler warnings in the touched files.

Rebased and validated with Claude Code; the test run above is a real GPU run, not CI (CI has no GPU, so it never exercises these paths).

[BelixRogner]

Verdict: mergeable as-is pending gates. Applies cleanly to current master (the rank objective files were untouched by last night's 27 commits).

The key acceptance criterion for us: zero cost when deterministic=false — confirmed. The dispatch is a host-side branch in LaunchGetGradientsKernel; the non-deterministic kernels are byte-identical to before. Determinism here is a verification tool, not a production requirement, and this PR treats it exactly that way.

Design review:

• Phase A (cub::BlockReduce for sum_lambdas) with a fixed reduction tree gives run-to-run bit-identity; the honest "~1 ulp vs CPU's sequential sum, scales only norm_factor" note is appreciated — this PR promises run-to-run determinism, not CPU parity, and says so.
• Phase B per-slot accumulation order: I checked the claim against CPU's flat (i, j) loop — for slot s, CPU's contributions arrive as (0,s), (1,s), …, (s-1,s), then (s,s+1)…(s,n), which is exactly your (i, r)-then-(r, j) order, and the score_t accumulator matches CPU's per-slot precision. Correct.
• The inverse-permutation trick makes the shared-memory variant O(n) per slot. The _Sorted_Deterministic variant (queries >2048 items) instead scans all pairs per slot with if (slot != high && slot != low) continue; — that's O(n²) per slot, O(n³/blockDim) per query. Fine for a debug tool, but add a comment stating the cost cliff, and note the inverse-rank trick as the upgrade path if anyone ever profiles it.

Nits:

• #include <cub/block/block_reduce.cuh> should sit with the other system/library includes (cpplint ordering); our ROCm branch will need a hipcub shim here — no action for this PR, just flagging.
• Full-training bit-identity also depends on the histogram pipeline being run-to-run stable; the 3-run test is empirical evidence, not proof. Tomorrow's gate will re-run it on the bench GPU (different SM count than your dev box — a good additional data point).

CI: the regular/Windows failures look inherited from the dd55ece1 rebase base (the dtype-test reconcile landed on master after) — should clear on re-merge; please confirm.

md5 locks: unaffected (default deterministic=false, and no locked benchmark uses lambdarank).

Gate runs tomorrow: both new determinism tests (50 and 1500 items/query) x3 runs, plus a timing sanity check that deterministic=false throughput is unchanged.

🤖 Review drafted with Claude Code