[cuda] hybrid level-batched tree growth: leaf-wise trees at depth-wise speed#32
Merged
Conversation
Grow whole levels while the leaf budget cannot bind (every positive-gain frontier leaf will be split by leaf-wise growth anyway), then hand off to the classic one-split-at-a-time loop with all frontier candidates cached. Split decisions are node-local, so the resulting trees match leaf-wise growth whenever the budget does not bind (verified: identical quality and leaf counts on fraud/year/covtype; bit-stable under quantized gradients). Per level this costs two device synchronizations (one batched split-info readback, one batched best-split readback) instead of three per split, and all of a level's histogram/split kernels are enqueued back to back: - data partition: deferred split mode (per-slot split-info buffer, level-end FinishSplitBatch readback; CopyDataIndicesKernel launched from host-known args), struct data_indices pointers written against the main index array instead of the per-split scratch, and an event ordering back-to-back splits against the async scratch->main copy (latent race in the classic loop, masked by its per-split syncs) - best split finder: optional per-pair synchronization, batched per-leaf best-split readback (SyncAllLeafBestSplitsToHost) - learner: TrainLevelWisePrefix + EnqueuePairBestSplitSearch; classic loop refactored onto the same pair-search helper Speed on covtype (1023 leaves, depth 10, 100 iterations, RTX 5090): non-quantized 7.8s -> 4.5s, quantized 12.1s -> 6.8s. Gated by EXABOOST_HYBRID_GROWTH (default on; single-GPU, non-categorical, no by-node feature sampling, no forced splits). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…/subtract/find/sync)
Replace the per-sibling-pair kernel launches of the hybrid level-wise prefix
with batched per-level variants: one launch each for histogram construct, fix,
subtract, find-best-split and sync-best-split covers ALL pairs of a level
(pairs indexed by a grid dimension), instead of ~7 tiny launches per pair.
Design:
- New CUDAHybridPairDescriptor (cuda_leaf_splits.hpp): per-pair leaf-splits
struct pointers, leaf indices/sizes, validity flags (mirrors the per-pair
construct/find gating exactly) and quantized histogram bit widths. The
learner builds one host array per level and uploads it with a single H2D.
- Histogram constructor: ConstructHistogramsForLevel launches batched
construct (blockIdx.z = pair; grid sized for the level's largest pair,
blocks beyond a pair's own rows exit early), batched fix and batched
subtract (blockIdx.y = pair) in order on pipeline_streams_[0], then records
subtract_done_events_[0]. Quantized per-pair bit widths are block-uniform
runtime branches; each mixed-bit pair (parent >16 bit, larger child <=16)
gets its own region of the 64->32-bit compaction buffer.
- Best split finder: FindBestSplitsForLevel waits on the subtract event, then
one batched find launch (grid = tasks x pairs x {smaller,larger}) writing a
per-pair region of cuda_best_split_info_ (resized to 2*num_tasks*max_pairs),
and one batched sync launch reducing each region with the exact same block
size / reduction order as the per-pair kernel (bit-identical results) into
the disjoint per-leaf best-split slots.
- Kernel bodies are shared: the existing per-pair kernels and the new batched
kernels call the same __device__ __forceinline__ helpers (shared-memory
histograms are declared by the wrapping kernel and passed in), so no
histogram math is duplicated.
- Gated by EXABOOST_HYBRID_BATCH_KERNELS (default on, "0" -> per-pair path)
and clean fallback for unsupported shapes: sparse rows, large-bin
partitions, compact column view, global-memory find, extra_trees, L1,
path smoothing, per-node feature selection, categorical features,
num_tasks > 1024.
Also fixes two pre-existing races in the (uncommitted) multi-pipeline
per-pair path that made quantized training nondeterministic:
- the 64->32-bit histogram compaction scratch buffer was shared by pairs
running concurrently on different pipeline streams (now one region per
pipeline);
- the smaller leaf's FindBestSplits waited only on the construct-done event,
which is recorded BEFORE FixHistogramKernel writes the most-frequent bin;
it now waits on subtract-done (recorded after fix + subtract).
With both fixes the per-pair fallback is deterministic again and produces the
BIT-IDENTICAL quantized covtype model as the batched path (md5 5f4e7bdfff1e).
Acceptance battery (100 trees, RTX PRO 6000 Blackwell; batched | fallback):
covtype 1023/10 quant : q 0.91952 lv 291.6 4.62s | 0.91952 291.6 6.74s (ref 0.9195/291.6; bit-identical md5, deterministic)
covtype 1023/10 : q 0.91887 lv 165.2 2.97s | 0.92029 166.9 4.29s (ref 0.915+-0.01 / 165+-10; was 4.5s)
fraud 63/6 : q 0.97469 lv 37.2 0.19s | 0.97469 37.2 0.22s (ref exact)
year 1023/10 : rmse 8.994 lv 670.1 2.74s | 8.994 670.1 3.32s (ref 8.99+-0.02 / 670+-10)
higgs 1023/10 : q 0.83208 lv 931.9 5.42s | 0.83208 931.9 5.86s (ref exact; was 6.2s)
Speed targets covtype <=2.5s / higgs <=4.0s missed (2.97s / 5.42s): nsys shows
the search phase is no longer the bottleneck. covtype: per-split PARTITION
kernels now dominate (SplitTreeStructure/Split/AggregateBlockOffset/SplitInner/
GenBitVector/UpdateDataIndex/CopyDataIndices = ~66% GPU time, ~7 launches per
split) plus ~2 cudaDeviceSynchronize per split. higgs: the leaf-wise TAIL
dominates (mean 931.9 leaves vs 512-leaf prefix -> ~420 one-at-a-time splits
per tree with per-pair search + FindBestFromAllSplits + device syncs).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The prefix previously stopped once the next level could approach the budget (2x headroom), forcing e.g. higgs (932-leaf trees, 1023 budget) through ~420 one-at-a-time tail splits per tree. When the finished tree never exhausts num_leaves, batching every fitting level is exactly leaf-wise-equivalent, so stop only when a level would overflow the budget. higgs 100-tree train: 5.42s -> 3.59s (quant 3.11s), year 2.74s -> 2.18s; trees unchanged (identical quality and leaf counts, covtype quant still bit-identical). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… kernels)
Batch the apply phase of the hybrid level-wise prefix: one launch per KERNEL
FAMILY per level instead of ~7 launches + 2 cudaDeviceSynchronize per split
(the syncs hid in LaunchSplitInnerKernel after AggregateBlockOffsetKernel and
SplitInnerKernel, ~554ms host time / 30 iters on covtype). The level's splits
partition disjoint index regions with host-known bounds (previous level's
FinishSplitBatch readback), so all updates are independent.
Design:
- CUDAHybridApplyDescriptor (cuda_data_partition.hpp): per-split region
(leaf_data_start, num_data), children indices, best-split-info + leaf-splits
struct pointers, per-split grid size and block-offset region start, and
runtime versions of the GenDataToLeftBitVector/UpdateDataIndexToLeafIndex
template flags (bit_type, missing/mfb flags, min_is_max, use_min_bin, ...).
Built host-side by CUDADataPartition::SplitLevelBatched (mirroring the host
metadata math of LaunchGenDataToLeftBitVectorKernel), one H2D per level.
- Batched partition kernels (cuda_data_partition.cu), all on cuda_streams_[0]
in dependency order, sized (max split blocks) x (num splits) with early-exit
for blocks beyond a split's own region:
* HybridGenBitVectorUpdateLeafIndexBatchKernel: FUSES the bit-vector and
data-index-to-leaf-index kernels (both derive from the same
(index, bin) load, halving the apply phase's dominant memory traffic);
decision logic replicates the templated kernels verbatim (including the
MISSING_IS_ZERO || MFB_IS_ZERO quirk) as block-uniform runtime branches.
* HybridAggregateBlockOffsetBatchKernel: one block per split, generalizes
AggregateBlockOffsetKernel0 to any block count; explicitly zeroes its
region's slot 0 (region layouts change between levels).
* HybridSplitInnerBatchKernel / HybridCopyDataIndicesBatchKernel: scatter
into the split leaf's own [leaf_data_start, +num_data) window of the
out-indices scratch (mirroring the main array layout), then one batched
region copy back. The stable partition is block-size invariant, so the
fixed 1024-thread block reproduces CalcBlockDim's results bit-exactly;
per-split grid dims are host-computed into the descriptor.
* HybridSplitTreeStructureBatchKernel: replica of SplitTreeStructureKernel
<false, USE_GRAD_DISCRETIZED> with point_structs_at_main fixed true,
writing deferred split-info slot k (FinishSplitBatch layout unchanged).
Per-split scratch becomes per-split regions of level-sized buffers: the
uint16 bit vector and out-indices use the region windows; the block offset
buffers grow to the level's total ceil(n/1024)+1 slots (classic-path resize
sites guarded against shrinking them back).
- CUDATree::SplitBatch (cuda_tree.{hpp,cpp,cu}): one SplitBatchKernel launch
(one 32-thread block per split) does all tree-structure updates of a level;
host mirrors (num_leaves_, leaf_depth_, branch features) update in the same
order the per-split loop would. Sibling splits' parent child-pointer
updates are race-free (node indices >= 0 never equal ~leaf codes).
- Learner: batched branch in TrainLevelWisePrefix builds both descriptor
arrays per level (right children take consecutive leaf indices, exactly as
the per-split loop assigns them); readback bookkeeping unchanged.
ZeroHistForLeaf calls dropped (no-op since BeforeTrain zeroes cuda_hist_).
- Gated by EXABOOST_HYBRID_BATCH_APPLY (default on) inside TrainLevelWisePrefix
only; "0" restores the per-split deferred loop (verified: reproduces HEAD
timings and the bit-exact covtype quant model). Tail loop and classic
leaf-wise growth untouched. Hybrid preconditions guarantee support
(single GPU, no categorical splits); bagging shares the same index array
(verified: batched == fallback model md5 with bagging_fraction 0.7, quant).
Acceptance (100 trees, RTX PRO 6000 Blackwell; batched | HEAD reference):
covtype 1023/10 quant : q 0.91952 lv 291.6 md5 5f4e7bdfff1e (bit-exact) 1.24s | 4.62s
covtype 1023/10 : q 0.91819 lv 165.7 1.11s | 3.03s (target <=2.0 met)
fraud 63/6 : q 0.97469 lv 37.2 0.13s | 0.19s (target <=0.15 met)
year 1023/10 : rmse 8.994 lv 670.1 0.98s | 2.18s (target <=1.5 met)
higgs 1023/10 : q 0.83208 lv 931.9 2.49s | 3.59s (target <=2.6 met)
higgs 1023/10 quant : q 0.83202 lv 931.3 2.04s | 3.11s
compute-sanitizer memcheck clean. nsys (higgs): per-split partition kernels
eliminated; histogram construct now dominates (43%), then the fused
gen+update (19%), split-inner (16%) and region copy (13%) batched kernels;
cudaDeviceSynchronize is down to the 2 per LEVEL readbacks.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ers, trim per-level/per-tree host overhead
Four independent optimizations of the hybrid level-batched growth, driven by
nsys profiles of the two cells still behind XGBoost (fraud deep: tiny-level
setup overhead; higgs shallow: leaf-wise tail + big-data kernel throughput):
1. Depth-capped final partial level (cuda_single_gpu_tree_learner.cpp).
When the leaf budget binds AND every splittable frontier leaf sits at
max_depth - 1, no child can ever become a leaf-wise candidate (leaves at
max_depth are marked invalid by the split finder) and splitting one frontier
leaf leaves the cached best splits of the others untouched. The leaf-wise
tail would therefore apply exactly the top-(remaining budget) candidates in
descending gain order, paying one full search+partition+3-sync round trip
per split -- including a wasted histogram construct for each depth-capped
child. Batch precisely that as one final partial level: sort candidates by
(gain desc, leaf index asc -- matching FindBestFromAllSplitsKernel's
strict-greater selection), truncate to the budget, apply in gain order (so
right-child indices match the sequential tail), then stop. On higgs 63/6
this removes ~31 tail iterations per tree and yields the IDENTICAL model
(md5-equal to the tail's) at 1.17s -> 0.95s per 100 trees.
2. Main index array double buffering (cuda_data_partition.{hpp,cpp,cu},
CUDAVector::Swap in cuda_utils.hu). The batched split-inner kernel already
writes each split leaf's partitioned indices at the leaf's main-layout
offsets of the out scratch; instead of copying every region back
(HybridCopyDataIndicesBatchKernel was 13% of GPU time on higgs), promote the
out buffer to the new main index array with an O(1) pointer swap and carry
only the complement -- terminal leaves' regions, computed host-side as the
gaps in the split regions' coverage of [0, root_num_data) -- via a sparse
batched copy. The batched tree-structure kernel points the child leaf-splits
structs at the new buffer; the classic per-split path keeps working off the
swapped members (its scratch is the old main). Verified bit-exact against
EXABOOST_HYBRID_BATCH_APPLY=0 (no swap) on quant models, including a
budget-bound tail-after-swap config (higgs 63/10) and bagging 0.7.
3. Batched construct grid sizing (CalcConstructHistogramBatchedKernelDim).
The per-leaf sizing forces min_grid_dim_y_=160 y-blocks even for tiny
leaves; every active block pays a fixed shared-hist zero + global atomic
merge (~2*num_bins_per_partition entries), which dominates small levels
(fraud deep levels spent 38us constructing ~20K-row pairs). Cap the y-grid
at 64 rows per thread but share the 160-block saturation floor across the
level's pairs (grid_y >= 160/num_pairs), so single-pair levels keep the full
grid (covtype, which is latency- not merge-bound, measured 0.98s vs 1.42s
under a hard cap) while many-pair levels shed the merge flood. Identical to
the per-leaf sizing whenever the cap is inactive; integer (quantized) hists
are bit-exact under any grouping, float quality gates re-verified exact.
EXABOOST_BATCH_CONSTRUCT_MINROWS overrides (0 = per-leaf sizing).
4. Host-overhead trims on hot paths. (a) The per-level descriptor uploads
(pair descs, apply descs, tree batch splits) become cudaMemcpyAsync on the
consuming stream -- ordered before their kernels by stream order, host
buffers rewritten only after each level's full sync, pageable async H2D
returns after staging. (b) FinishSplitBatch / SyncAllLeafBestSplitsToHost
drop their explicit device syncs: the synchronous legacy-stream D2H already
waits on all (blocking) streams. (c) Per-tree SetCUDAMemory sites
(CUDATree::InitCUDAMemory, CUDADataPartition::BeforeTrain,
CUDAHistogramConstructor::BeforeTrain, CUDALeafSplits::InitValues) become
plain async cudaMemset ordered by the legacy stream, and the leaf-splits
launchers' intra-stream device syncs are kept only on the NCCL path,
removing ~10 device syncs per tree.
Battery (100 trees, RTX 5090, best of 3; before = this branch's parent):
covtype 1023/10 quant : 1.08s (1.24) q 0.91952 lv 291.6 md5 5f4e7bdfff1e BIT-EXACT
covtype 1023/10 : 1.04s (1.06 measured, 1.11 ref) q 0.919-0.920
higgs 1023/10 : 2.20s (2.45) q 0.83208 exact, lv 931.9
higgs 1023/10 quant : 1.79s (1.76-2.04) q 0.83202, lv 931.3
year 1023/10 : 0.62s (0.94) rmse 8.994, lv 670.1
fraud 1023/10 : 0.18s (0.20) [target 0.12 not reached, see below]
fraud 63/6 : 0.12s (0.13) q 0.97469 exact, lv 37.2 [target 0.10]
higgs 63/6 quant : 0.95s (1.16) [target 0.70 not reached, see below]
higgs 63/6 : 1.05s (1.30) [target 0.85]
Fallbacks: EXABOOST_HYBRID_BATCH_APPLY=0 and =BATCH_KERNELS=0 reproduce the
covtype quant lock md5 bit-exactly; EXABOOST_HYBRID_GROWTH=0 (classic) clean.
compute-sanitizer memcheck clean (hybrid quant/nonquant, classic, partial-level
path). Remaining gaps are structural: higgs 63/6 wall is ~0.41s booster init
(row-data transpose + 750MB pageable H2D + first-launch module load, amortized
away at 500 trees) plus ~5.3ms/tree of bandwidth-bound kernels; fraud deep is
~115us/level of which two mandatory D2H sync points + 11 kernel launches make
~60us. Attempted but not landed: speculative next-level search enqueue before
the FinishSplitBatch readback (device-evaluated validity) -- would overlap the
apply readback with the next search but requires find/sync kernels to derive
leaf indices/validity/quant bit widths from device structs; deferred as the
risk outweighed the remaining ~30us/level.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ice arrays, deferred root sums
Collapses the hybrid level-batched growth from TWO host sync points per level
(best-split readback + FinishSplitBatch) to ONE, and removes the per-tree
cudaMalloc/cudaFree/D2H churn that an nsys trace of fraud/deep showed dominating
the remaining wall time. Four parts:
1. Single-sync (speculative) level pipeline (EXABOOST_HYBRID_ONE_SYNC, default
on; non-quantized batched-kernels+batched-apply configs only -- quantized
training keeps the classic two-sync flow because it selects histogram kernels
host-side from per-leaf bit widths). TrainLevelWisePrefixOneSync enqueues each
level's APPLY and the children's SEARCH back to back and reads back the
level's deferred split info AND the children's best splits with one device
synchronization (the second, deferred-info copy finds the device idle). The
search kernels no longer need any host value from the apply:
- The batched find/sync kernels derive leaf index, size and hessian sum from
the child leaf-splits structs the apply kernels write (the smaller/larger
designation is decided on-device by HybridSplitTreeStructureBatchKernel);
the descriptor validity flags now carry only the host-known max_depth gate,
ANDed on-device with the min_data/min_sum_hessian gates. In the classic
flow the host flags already include identical conditions, so the device
check is a no-op -- the quantized find kernel is untouched and the shared
sync kernel is value-identical (covtype quant md5 lock verified bit-exact).
- The batched construct kernel evaluates the construct gating (mirror of
ConstructHistogramForLeaf's early return) on-device from the structs.
- The speculative construct grid is sized by the parents' sizes (smaller
child <= floor(parent/2), host-known pre-level); the exact row-grouping
extent is recomputed on-device from the level's ACTUAL smaller-child sizes
(written per split by HybridAggregateBlockOffsetBatchKernel) with the exact
host sizing formula (HybridBatchedConstructGridDimY, now the single source
for host and device), so float histograms are bit-identical to the classic
sizing; excess blocks early-exit. The reduction+formula runs as one tiny
single-block kernel per level (ComputeBatchedConstructDimYKernel) -- an
earlier per-block variant regressed higgs 2.20 -> 2.40s.
- Ordering: the histogram stream waits on the partition's apply-done event;
level boundaries stay fully serialized by the (synchronous) readback.
- The final partial (depth-capped) level keeps its immediate readback, and
budget-bound levels defer to the leaf-wise tail exactly as before. Root
level uses device-derived gating too, with an exact host-known grid.
The classic loop and the one-sync loop share the extracted helpers
(CollectSplittableLeaves / ArbitrateLevelBudget / ApplyLevelBatched /
FinishLevelBookkeeping); all previous env fallbacks route to the classic loop.
2. Pooled per-tree device arrays (CUDATree + CUDAVector::SetView/Materialize).
Each tree paid ~17 cudaMallocs (ctor), ~16 cudaFrees + 16 synchronous D2H
copies (ToHost) -- ~53us of alloc API + ~90us of copies per tree, 30ms of the
fraud run's 170ms. CUDAVector gains a non-owning view mode (mutating
reallocation copies-then-owns; views are never freed), and CUDATree carves
all 16 max_leaves-sized arrays + the batch-split descriptors from a
learner-owned slab passed to the ctor. ToHost reads the whole slab back with
ONE fused D2H copy and scatters host-side; the retained leaf values are
materialized into owned memory (Resize on a view copies; Materialize covers
the full-tree case), everything else detaches without cudaFree, so the OOM
guard semantics are preserved. Fresh cudaMalloc'd memory was never zeroed
either, so slab reuse across trees is equivalent (write-first everywhere,
plus the same three first-element memsets).
3. Deferred root-sum readback. The 8-byte synchronous root-sum D2H in
CUDALeafSplits::InitValues blocked ~70us per tree on the (legacy-stream
ordered) full histogram-buffer zeroing. The single-sync flow derives every
root gate on-device, so BeforeTrain defers the readback (and Train defers the
root leaf-output init, which any split overwrites on device anyway);
EnsureRootSumsReadBack performs both lazily on the paths that still need host
root sums (classic prefix start, stump/no-valid-root-split case).
4. Trims: the dead construct_done event record in ConstructHistogramsForLevel
(the batched find only waits on the subtract event) is gone;
CUDADataPartition::ResetConfig now resizes cuda_split_info_buffer_ (was
missed) and the new per-level smaller-count array.
fraud/deep per-tree window (nsys, tree 50): 1328us -> ~1205us; per tree now 29
sync memcpys (was 47), 1 cudaMalloc (was 18), 0 frees on the hot path (was 16).
Remaining profile: ~690us/tree GPU pipeline (construct 287us -- saturation-floor
merge cost on tiny levels, find 138us, apply family ~150us) + 130 launches
(~234us API) + 11 level syncs absorbing the pipeline; further gains need
construct-kernel work shared with the (regression-locked) classic sizing.
Battery (100 trees, RTX 5090, best of 3; before = parent commit):
covtype 1023/10 quant : 1.00s (1.08) q 0.91952 lv 291.6 md5 5f4e7bdfff1e BIT-EXACT
covtype 1023/10 : 0.95s (1.02-1.04) q 0.918-0.921
higgs 1023/10 : 2.21s (2.20) q 0.83208 exact, lv 931.9
higgs 1023/10 quant : 1.78s (1.79) q 0.83202, lv 931.3
year 1023/10 : 0.63s (0.62) rmse 8.9944, lv 670.1
fraud 1023/10 : 0.16s (0.17-0.18) [target 0.13 not reached, see above]
fraud 63/6 : 0.10s (0.12) q 0.97469 exact, lv 37.2 [TARGET reached]
higgs 63/6 quant : 0.94s (0.94-0.95) q 0.81249 md5-stable
higgs 63/6 : 1.03s (1.05)
Fallbacks: EXABOOST_HYBRID_ONE_SYNC=0, =BATCH_KERNELS=0 and =BATCH_APPLY=0 all
reproduce the covtype quant lock md5 bit-exactly and fraud quality/leaves
exactly; EXABOOST_HYBRID_GROWTH=0 (classic) clean. Bagging 0.7 quant md5
identical with ONE_SYNC on/off; nonquant differences are within run-to-run
float-atomic noise. compute-sanitizer memcheck clean on the one-sync nonquant,
quantized classic and GROWTH=0 paths.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EYUXNtCAXDd3zYT71m8YSJ
… (fraud/deep)
Measure-first pass over the fraud/deep cell (285K x 29, trees ~83 leaves across
~11 levels; nsys tree window 1180us: construct 287us dominated by the
shared->global merge on small levels, 75us full-histogram-buffer memset, 120
launches, ~19us host turnaround per level). Four levers, three landed
default-on (all BIT-IDENTICAL to the previous outputs), one opt-in:
1. Partial histogram-buffer zeroing (BeforeTrain). Leaf k's histogram lives at
slot k and every tree assigns slots [0, num_leaves) consecutively, so after
a tree only that prefix is dirty and the buffer is all-clean before every
tree by induction. Zero just the previous tree's leaf count of slots
(SetNumDirtyLeaves from the learner, single-GPU only; -1 = full zeroing on
init/resets/NCCL): with a 1023-leaf buffer and ~83-leaf trees this cuts
~110MB (~65us) of memset per tree. Helps every dataset; value-identical.
2. Fused fix+subtract launch (non-quantized batched levels). One kernel does
both the most-frequent-bin fix of the smaller leaf and the parent-minus-
smaller subtraction: subtract blocks skip the need-fix features' mfb
gradient/hessian entries (precomputed cuda_fix_mfb_mask_); one fix block per
need-fix feature runs the exact FixHistogramInner reduction (same
FIX_HISTOGRAM_BLOCK_SIZE) and applies the subtraction at precisely those
skipped entries from the fixed values. Every entry is written by exactly one
block with the sequential launches' arithmetic => bit-identical, one launch
less per level. Quantized levels keep the sequential discretized kernels
untouched (covtype quant md5 lock byte-identical by construction).
3. Pinned staging for the two per-level synchronous readbacks
(SyncAllLeafBestSplitsToHost, FinishSplitBatch): a pageable sync D2H pays an
extra driver staging round trip on the level critical path; the host-side
memcpy of a few KB afterwards is negligible.
4. Folded dim_y computation: few-pair (<= 32) speculative levels evaluate the
HybridBatchedConstructGridDimY row-grouping formula inside the construct
kernel (block-uniform, <= 32 loads) instead of launching
ComputeBatchedConstructDimYKernel first -- identical integer result, one
launch less per level; many-pair levels keep the reduction kernel.
OPT-IN (default off, EXABOOST_SMALL_LEAF_ROWS=8192 enables): direct small-leaf
construct body. Pairs whose ACTUAL smaller leaf (on-device check, so the
speculative flow's parent/2 bounds cannot mask a tiny pair behind a skewed
sibling) is at most the threshold add each row straight to the global
histogram with device-scope double atomicAdd, skipping the shared-histogram
zero+merge entirely; excess blocks exit after one comparison. Measured only
~1-2% on fraud/deep (the expensive levels are the mid-sized ones that
legitimately need the shared path) and it changes the float accumulation
(per-row double adds vs per-block float partial sums), which flips the
fraud 63/6 exact-quality reproduction (0.97469 -> 0.97370) -- not worth it as
a default. Also fixed en route: the direct body must add the partition's
column_hist_offsets_full start (the per-column offsets are partition-relative).
Not landed, and why:
- Batched find+sync fusion (single kernel per level): SyncBestSplitForLevel's
1024-thread ReduceBestGain order is part of the bit-locks' tie-breaking; a
fused 256-thread tail could flip equal-gain winners, and the quantized flow
shares the kernel. Measured upside ~4-5% of the tree window; risk > reward.
- Construct saturation-floor tuning: exposed as EXABOOST_BATCH_CONSTRUCT_FLOOR
(default 160 = previous behavior); F in {0,16,64} moved fraud wall by ~0
(mid-level construct is row-bound, not merge-bound), so the default stays.
Env: EXABOOST_SMALL_LEAF_CONSTRUCT=0 kills the fused fix+subtract AND the
direct body (restores the previous launch sequence exactly);
EXABOOST_SMALL_LEAF_ROWS=<n> opts into the direct body.
Battery (100 trees unless noted, RTX 5090; before -> after, best of 3-5):
fraud 1023/10 : 0.16-0.17 -> 0.15s q 0.97548 lv 83.0 (exact-equal to parent)
fraud 1023/10 500t : 0.62-0.64 -> 0.54s q 0.97044 (gate 0.9704+-0.002 OK; XGB ref 0.52)
fraud 63/6 : 0.10s q 0.97469 lv 37.2 EXACT
covtype 1023/10 quant : 1.00-1.08 -> 0.99s q 0.91952 lv 291.6 md5 5f4e7bdfff1e BIT-EXACT
covtype 1023/10 : 0.95-1.04 -> 0.91-0.92s q 0.920-0.922
higgs 1023/10 : 2.21 -> 2.15-2.17s q 0.83208 EXACT, lv 931.9
higgs 63/6 quant : 0.94 -> 0.93-0.98s q 0.81249 md5-stable
year 1023/10 : 0.62-0.63 -> 0.58s rmse 8.9944 lv 670.1
Fraud/deep target 0.11s not reached in wall: the steady-state tree window is
1031us median (was ~1180us) = 0.103s/100 trees AT target; the remaining wall
gap is ~50ms one-time process warmup (first-launch module load + allocator
warm-up) inside the train() timer. Remaining steady profile per tree:
construct 240us (row-bound mid/large levels), find 136us (FP64 gain math,
leaf-size independent), ~230us of 2-4us apply/sync kernels, ~344us GPU idle
across ~11 level turnarounds (readback -> host bookkeeping -> relaunch).
Fallbacks verified bit-exact: EXABOOST_SMALL_LEAF_CONSTRUCT=0 reproduces the
fraud quality/leaves and the covtype quant md5 exactly; ONE_SYNC=0,
BATCH_KERNELS=0, BATCH_APPLY=0 all keep md5 5f4e7bdfff1e; HYBRID_GROWTH=0
clean. compute-sanitizer memcheck clean on fraud (default and
SMALL_LEAF_ROWS=8192) and covtype quant.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The benchmark harness times lgb.train() in a fresh process; Booster creation (learner Init) and the first Train() call sit inside that timer while Dataset construction does not. Profiling fraud/deep (285K x 29, 1023 leaves, depth 10, 500 trees) attributed ~44 ms of one-time cost to that window: - 37.4 ms: the first cudaLaunchKernel of the process pays the lazy CUDA module load of the device-linked fatbin (sm_120 SASS is present, so this is a plain module load, not PTX JIT). The first launch used to happen inside the learner's Init -> first Train() path. Fix: launch a trivial warm-up kernel from CUDAColumnData::Init (cuda_column_data.cu), which runs during Dataset construction -- the CUDA context is already created there, and the booster/learner does not exist yet. After the warm-up, every later kernel's first launch costs microseconds. - 3.7 ms: TrainLevelWisePrefix allocated num_leaves + 2 CUDALeafSplits(1) objects on the first tree, ~3 cudaMallocs each (3076 calls observed). The hybrid pair slots only ever expose their bare 1-element CUDALeafSplitsStruct (InitValues and the sum-reduction buffers are never used on slots), so replace the object vector with one device slab CUDAVector<CUDALeafSplitsStruct>, sized once: 1 cudaMalloc. Kernel-visible behavior is unchanged -- the split kernels initialize every slot they use before any read, same as before. Remaining Booster-create cost (~10 ms) is real one-time work on the critical path (row-wise MultiValBin build feeds CUDARowData's H2D upload, objective/score-updater init) and is left alone. fraud/deep 500 trees, fresh process, median of 3: 0.54s -> 0.51s (booster create 40 ms -> 13 ms, tree0 5.4 ms -> 2.0 ms, steady-state per-tree unchanged at ~945 us). Battery: covtype 1023/10 quant 0.91952 / 291.6 / md5 5f4e7bdfff1e (bit-exact); fraud 63/6 0.97469 / 37.2; fraud/deep quality 0.97044; higgs 1023/10 0.83208 @ 2.14s; covtype nonquant 0.88s; year 0.55s. All EXABOOST_HYBRID_* env fallbacks verified equivalent. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…gime Level batching is exactly leaf-wise-equivalent only when the budget can bind solely at max_depth (2^max_depth <= num_leaves + 1), where the final-level top-K-by-gain selection is optimal because no deeper candidates can exist. In budget-limited configs (num_leaves << 2^max_depth), applying whole levels can spend budget on shallow low-gain splits that best-first growth would invest deeper in richer subtrees: measured -1.1pp accuracy on covtype at num_leaves=64, max_depth=12 (0.9077 vs 0.9185), while homogeneous-gain data (numerai 4096/13) showed no measurable difference. Gate the prefix accordingly; budget-limited configs fall back to the classic loop until exact grow-then-prune (ancestry-closed top-k) selection is implemented. All benchmark regimes are depth-limited and keep their hybrid timings; the covtype quant reference stays bit-exact (md5 5f4e7bdfff1e). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Owner
Author
|
Semantics correction (credit: Felix's review question): the "leaf-wise equivalent" claim only holds in the depth-limited regime (2^max_depth <= num_leaves+1). In budget-limited configs, whole-level batching can spend budget on shallow low-gain splits that best-first growth would invest deeper — measured −1.1pp accuracy on covtype at 64 leaves/depth 12. Latest commit gates the hybrid prefix to the provably-exact regime (budget-limited configs fall back to the classic loop); all benchmark configs are depth-limited, so the reported numbers are unaffected. Follow-up for recovering hybrid speed in budget-limited configs: exact grow-then-prune with ancestry-closed top-k selection. |
…mited configs Opts into level batching where it is only approximately leaf-wise (num_leaves << 2^max_depth), trading the measured breadth bias (covtype 64/12: 0.9067 vs 0.9185 classic) for hybrid speed (2.06s vs 5.79s). Default remains the exact depth-limited gate. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ted configs
Budget-limited configs (num_leaves << 2^max_depth) previously fell back to the
classic one-split-at-a-time loop because plain level batching is breadth-biased
there. This adds a SELECTIVE level-batched mode that is exactly leaf-wise
equivalent and becomes the default for budget-limited configs (kill switch
EXABOOST_HYBRID_SELECTIVE=0; EXABOOST_HYBRID_AGGRESSIVE=1 still selects the
approximate plain batching; depth-limited configs keep the bit-exact-locked
exact-fit prefix unchanged).
Selection algorithm (per level, host side, O(C log C)):
- Leaf-wise growth is greedy selection: repeatedly apply the max-gain candidate
among AVAILABLE candidates (available = parent split already applied) until
num_leaves-1 splits or no positive gain. Monotonicity: a newly discovered
(deeper) candidate inserts into the greedy sequence right after its parent's
selection and consumes budget earlier, so it can only DISPLACE currently
selected candidates, never promote unselected ones. Hence per level:
1. run the greedy simulation (max-heap) over live applied splits + the newly
searched frontier candidates, budget num_leaves-1;
2. apply ONLY selected frontier candidates (batched apply + search as in the
exact-fit prefix; unselected candidates are permanently dormant);
3. eagerly COLLAPSE applied splits displaced from the selection (their whole
subtree is displaced by availability closure);
4. stop when no frontier candidate is selected: the live applied splits are
then exactly the classic leaf-wise selection.
- Classic leaf numbering is simulated along the way (the t-th selected split
keeps its leaf index for the left child, the right child takes t+1) and
equal-gain ties are resolved by an exact host replica of
FindBestFromAllSplitsKernel's strided-scan + shuffle-down reductions
(keep-current-on-tie) -- the device tie-break is NOT "smallest leaf index",
which an md5 bisect on covtype 128 exposed (two equal-gain splits applied in
reduction order, not leaf order).
Pruning = eager collapse with index recycling (no extra device memory):
- A collapsed subtree occupies one contiguous window of the main data index
array (child regions nest), so collapse = revert the topmost displaced node
to a leaf at its left-child index, rewrite the window's
data-index-to-leaf-index entries to it (CollapseLeafWindowsKernel, one
batched launch per level on the apply stream), and recycle every right-child
leaf index of the subtree. Recycled histogram slots (slot == right-child
index) are re-zeroed (ZeroHistSlots) before reuse. Live leaves never exceed
num_leaves (selection holds <= num_leaves-1 splits; fresh indices are only
allocated when every allocated index is live), so ALL existing device
buffers keep their sizes; the quantized parent-bit lookup gets an explicit
parent field in HybridPendingPair because recycled right-child indices break
the old min(smaller, larger) == left assumption.
- Tree rebuild: the CUDA tree object is not touched during growth. At tree end
the final greedy sequence is replayed host-side (CUDATree::
RebuildFromHostSplits) from the per-split CUDASplitInfo host copies captured
at apply time -- bit-identical arithmetic to SplitKernel on the same device-
computed values, in classic order with classic node/leaf numbering, so
deterministic (quantized) configs produce byte-identical model files to the
classic loop. Device leaf values are uploaded for scoring/shrinkage/renewal;
ToHost() is skipped (RebuildFromHostSplits shares its shrink+release tail).
- Score/partition consistency: one RemapDataIndexToLeafIndexKernel pass maps
every used row's leaf index from hybrid to final numbering (pruned rows were
already rewritten to their reverted dormant leaf at collapse), and the final
per-leaf data layout is uploaded so AddPredictionToScore, RenewTreeOutput
and quantized leaf renewal all see classic numbering. The next BeforeTrain
zeroes the hybrid allocation peak of histogram slots.
- Both level pipelines are supported: the non-quantized single-sync
(speculative) flow and the two-sync flow (quantized, or per-pair search
fallback for shapes the batched kernels do not cover, e.g. numerai's compact
view). Linear trees keep the classic loop.
EXACTNESS (identical model md5, GROWTH=1 vs GROWTH=0, quantized =
deterministic; 100 trees):
covtype 64/12 quant: 01a30b7ff02f == classic (0.88s vs 2.79s)
covtype 128/12 quant: d7ff4fcac6d3 == classic (1.02s vs 5.34s)
fraud 32/12 quant: 97cbd25fb86e == classic (0.12s vs 0.23s)
Non-quant parity (float atomics are order-sensitive, so quality gate):
covtype 64/12 200t: 0.91937 vs classic 0.91850 (+0.0009, within +-0.002)
fraud 32/12 200t: 0.97162 == classic exactly; numerai rmse 0.22368 == classic
SPEED (best of 3, same HEAD, RTX 5090; classic GROWTH=0 -> selective):
covtype 64/12 200t : 5.89s -> 1.80s (3.3x, target <=3.5 met)
covtype 128/12 200t : 11.10s -> 2.17s (5.1x, target <=6.5 met)
fraud 32/12 200t : 0.50s -> 0.19s (2.6x, target <=0.35 met)
numerai 4096/13 100t: 56.31s -> 37.47s (1.5x, target <=40 met; runs 40.22/37.47/37.63)
CHURN (displaced/applied, cumulative over the run):
covtype 64/12 quant: 47982/92082 (52%; applied = 2.09x final splits)
covtype 128/12 quant: 72305/161194 (45%)
fraud 32/12 quant: 1683/4783 (35%)
numerai 4096/13 : 120301/529801 (23%; applied = 1.29x final, ~14 levels/tree)
EXISTING LOCKS unchanged (depth-limited exact-fit path untouched):
covtype 1023/10 quant: 0.91952 / 291.6 / md5 5f4e7bdfff1e BIT-EXACT (0.94s)
fraud 63/6: 0.97469 / 37.2 (0.08s); higgs 1023/10: 0.83208 exact @ 2.12s
year 1023/10: rmse 8.9945 @ 0.54s; covtype 1023/10 nonquant 0.88s
Fallbacks: GROWTH=0 classic; SELECTIVE=0, BATCH_KERNELS=0 and BATCH_APPLY=0
route budget-limited configs to the classic loop (prior behavior, all md5
01a30b7ff02f == classic on covtype 64/12 quant); ONE_SYNC=0 runs selective
through the two-sync flow (covtype 64 quant md5 unchanged, nonquant quality
verified); AGGRESSIVE=1 still selects approximate level batching (0.90668,
matching the historical breadth-biased number). feature_fraction 0.5 quant,
bagging 0.7 quant and the per-pair-search selective flow (quant +
lambda_l1=0.5, which disables the batched find kernel and exercises the
recycled-index parent-bits fix) are all md5-identical to classic; numerai runs
the compact-view per-pair fallback end to end with exact quality parity (rmse
0.22368 both). compute-sanitizer memcheck clean on the selective one-sync,
two-sync quant and per-pair (L1) paths.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Owner
Author
|
Update: exact grow-then-prune (selective level growth) landed in cc5c00a — the hybrid is now exactly leaf-wise-equivalent in ALL regimes, including budget-limited configs, with md5-proven equivalence in deterministic quant mode. Speed on budget-limited configs vs classic leaf-wise: covtype 64/12 3.3x, covtype 128/12 5.1x, fraud 32/12 2.6x, numerai 4096/13 1.5x — also 2.5-5.4x faster than XGBoost lossguide on the same configs. Churn (speculative splits later pruned): 23-52%. The earlier depth-limited gate and the AGGRESSIVE flag remain as documented alternatives. |
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…shape specialization, persisted tuning cache) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…atasets (numerai shape)
The batched per-level SEARCH path fell back to per-pair kernels on
numerai-shaped data (5.4M rows x 2748 features, ~7 bins/feature,
feature_fraction=0.1), capping the hybrid gains. Diagnosis (verified on the
real dataset with the new env-gated EXABOOST_HYBRID_DIAG dump):
1. CUDABestSplitFinder::SupportsBatchedLevel: num_tasks_=2748 > 1024
(one sync block) -- gated BOTH quant and non-quant runs.
2. CUDAHistogramConstructor::SupportsBatchedLevel: use_compact_view_=1
(the per-tree compact column gather that feature_fraction triggers)
-- gated non-quant runs. use_global_memory_ was 0 (that concern was
spuriously conservative: ~7 bins/feature).
Batched support for both shapes:
- Multi-block batched sync (finder): SyncBestSplitForLevelKernel gains a
task-block grid dimension (blockIdx.z), each block reducing its
1024-task slice with the identical slicing/ReduceBestGain order as the
per-pair SyncBestSplitForLeafKernel and writing slot
leaf + block * num_leaves; a new SyncBestSplitForLevelKernelAllBlocks
replicates the per-pair AllBlocks sequential strictly-greater merge
(lowest task block wins ties) per (leaf role, pair). Bit-identical to the
per-pair reduction: numerai quant md5s equal the per-pair path exactly.
- Compact-view batched construct (histogram constructor): the batched dense
construct kernel is fed the per-tree compact data/metadata and compact
block dims (mirroring the per-pair compact launch); fix/subtract are
layout-independent. This also unlocks the one-sync speculative prefix and
the selective one-sync flow for non-quant numerai.
Wide-shape efficiency (all masks nullptr -- and hence byte-identical
behavior -- without feature sampling):
- Find-task compaction: the batched find launches one block per USED task
only (per-tree list built in CUDABestSplitFinder::BeforeTrain); output
slots keep the original task index and the batched sync masks unused
lanes not-found, which is exactly what the full-grid find would have
written (is_valid=false lanes never propagate through ReduceBestGain).
10x fewer find blocks at feature_fraction=0.1.
- Per-tree bin-used mask (built in SetFeatureUsedBytree): the batched
fix/subtract kernels (float + discretized + the fused small-leaf kernel)
and the batched construct's shared-hist zero/global-merge skip histogram
entries of features outside the tree's sample -- nothing in the tree ever
reads them (the find is feature-masked), so they are dead storage.
- The tree learner's per-tree column-view gather sources from the histogram
constructor's compact matrix instead of re-reading the full bin matrix
(10.1ms -> 4.1ms per tree; the full-matrix gather effectively streamed
all 14.8GB through L2 once per tree). A fused col-major second output in
the fill kernel itself was measured SLOWER (uncoalesced slot-major
writes) and is not used.
- Register-accumulation construct body (kRegHistMaxBins=8) on the batched
compact path when every feature has <= 8 bins: per-thread register sums
with one flush instead of two same-address shared atomics per row
(numerai construct turned out latency-bound, so the measured gain is
small; kept for its contention-free flush). EXABOOST_BATCH_REGHIST=0
kill switch; non-quantized only (float order changes).
Kill switch for the whole wide-shape support: EXABOOST_BATCH_WIDE=0
restores the previous per-pair fallback for num_tasks > 1024 and for the
compact view. All existing env fallbacks unchanged.
Acceptance (RTX 5090):
- numerai 4096 13 quant 20 trees: GROWTH=1 md5 e64e9e021cc7 == GROWTH=0
(selective classic-numbering equivalence holds through the multi-block
sync + host tie-break).
- numerai 32 5 quant 50 trees: batched md5 23fac3b87259 == per-pair hybrid
(BATCH_WIDE=0) md5, i.e. the new kernels are bit-exact vs what they
replace. (GROWTH=1 vs GROWTH=0 md5s differ here as they already did at
HEAD: the depth-limited exact-fit prefix keeps level-order leaf
numbering by design -- same tree, same predictions, different node ids;
pre-existing, orthogonal to this change.)
- numerai 32 5 non-quant 100 trees: rmse 0.22367 both GROWTH=1/0.
- Locks unchanged: covtype 1023/10/1 -> 0.91952/291.6/5f4e7bdfff1e;
covtype 64/12/1 md5 == GROWTH=0 (01a30b7ff02f); fraud 63/6/0 ->
0.97469/37.2; higgs 1023/10/0 -> 0.83208 @ 2.13s; timings covtype
0.88s / year 0.44s / higgs 2.13s / fraud 0.07s per 100 trees (all
within the 3% budget of 0.92/0.58/2.12/0.08).
- Speed: numerai 4096 13 0 100: 39.4s -> 27.2s (1.45x, target <= 30s met).
numerai 32 5 0 200 (example-config shape): 22.4s -> 21.1s total; the
wall clock is dominated by ~14s of one-time first-train setup, marginal
per-tree cost 41.6ms -> 34.8ms (1.20x). Per the fresh nsys profile the
remaining per-tree time is histogram construct (55%, latency-bound
scattered reads) and the per-tree compact fills (41%), not the search
path (find+sync now < 1%).
- compute-sanitizer memcheck clean on numerai 32/5 non-quant (one-sync
compact path) and numerai 4096/13 quant (wide selective path).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…construct-latency follow-ups Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…+ config at init Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… items Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…l bin
Booster creation on wide dense datasets (numerai: 5.4M x 2748, uint8 bins)
spent ~13.9s building CUDARowData through the host MultiValDenseBin: a 15GB
zero-filled multi-val bin, a per-element iterator transpose to fill it, a
second 15GB zero-filled partition buffer, a repack, and the H2D copy. On the
CUDA path the host multi-val bin exists only to feed
CUDARowData::GetDenseDataPartitioned, so:
- CUDARowData::Init now builds the partitioned row-major buffer directly from
the Dataset's per-group column bins (4/8/16/32-bit, incl. 4-bit packed) with
a cache-tiled parallel transpose into an uninitialized staging buffer. For
non-multi-val dense groups the multi-val bin stores the raw group bin value
unchanged, so the transpose is byte-identical to the old path.
- Dataset::GetShareStates skips the host multi-val bin build entirely when the
caller is the CUDA tree learner and the dataset qualifies (dense row-wise,
no multi-val or sparse groups); CUDARowData recovers bit_type from
column_hist_offsets. Everything else the CUDA learner reads from the share
state (feature_hist_offsets, num_hist_total_bin, column_hist_offsets) comes
from CalcBinOffsets, and all TrainingShareStates methods already no-op on a
null multi-val bin wrapper, so bagging / quantized / feature_fraction /
compact-view paths are unaffected (verified functionally).
- Where the old repack remains as a fallback, the redundant zero-fill of the
partition buffer is skipped too.
EXABOOST_FAST_ROWDATA=0 restores the old path bit-for-bit;
EXABOOST_FAST_ROWDATA_VERIFY=1 builds both buffers and Log::Fatal's on any
byte mismatch (passes on numerai, covtype, fraud, year, higgs, epsilon,
incl. quantized, bagging and feature_fraction configs).
numerai (RTX 5090): Booster create 13.8s -> 2.3s, peak host RSS -14.6GB;
construct and per-tree times unchanged. Model md5 regression locks vs the
previous build hold for covtype 1023/10 quant, covtype 64/12 quant and
year 63/6 quant under EXABOOST_HYBRID_GROWTH={0,1} x EXABOOST_FAST_ROWDATA={0,1}
(fraud non-quant is run-to-run nondeterministic upstream; quality and tree
stats match). CPU path untouched: deterministic covtype run md5-identical
with the flag on/off.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EYUXNtCAXDd3zYT71m8YSJ
…ion item Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Accept int8/int16 numpy matrices in lgb.Dataset zero-copy (new C_API_DTYPE_INT16; INT8 previously only bypassed the float copy) and bin them through per-column value->bin lookup tables instead of a per-value double conversion + bin-boundary binary search: - Bin::PushBlock: block push of pre-encoded bins with a kSkipBin sentinel; DenseBin overrides it with a devirtualized inlined loop. - FeatureGroup::EncodeBinForPush/PushTargetBin expose PushData's encoding (most-freq skip, offset, multi-val adjust) for LUT folding. - Dataset::PushDenseSmallIntRows: tiled (512-row) column-blocked push through fully encoded LUTs; row-major and col-major, raw_data kept for linear trees. - c_api: SampleDenseSmallInt column-parallel bin-mapper sampler for the int paths (ints pass the kZeroThreshold check iff nonzero). - basic.py: int16 passes through like int8 instead of astype(float32), removing the 4x input copy. Bit-identical models vs float32 input (md5 gates): numerai int8 slice (quant CUDA 20 trees + CPU deterministic), covtype as int16 (CPU+CUDA), synthetic int8 with negatives/constant/sparse cols in C and F order. Float paths untouched: covtype 1023/10/quant hybrid=1 md5 5f4e7bdfff1e and hybrid=0 fcb9f6c2ab87 unchanged. Full-shape numerai (5.43M x 2748, RTX 5090, medians of 3): f32 memmap fed: construct 38.9s, booster create 2.34s, peak RSS 86.4G int8 memmap fed: construct 15.8s, booster create 2.34s, peak RSS 43.9G Remaining int8 construct time is dominated by EFB FindGroups (~9s), which is out of scope here. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EYUXNtCAXDd3zYT71m8YSJ
…n/hist elevated (speed over determinism) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
When every column fits in <=16 bins (numerai: 6 max), store the dense CUDA bin structures at two cells per byte instead of one: - CUDARowData: the partitioned row-major matrix is packed on the host (fast column-transpose path and multi-val fallback) and uploaded packed. Alignment invariant: per-partition packed row width = ceil(cols/2) bytes (column count padded to even PER PARTITION), column j in byte j>>1, nibble j&1 (low nibble = even column, mirroring DenseBin 4-bit). - Per-tree compact view: the compact matrix is packed the same way; a new fill kernel owns one DESTINATION byte (a compact column pair) per thread, reading source cells by precomputed nibble index. - Consumers unpack in-kernel: dense construct bodies (non-quant fp64/fp32 incl. the register-accumulation <=8-bin body and the small-leaf direct body; quantized int16/int32, per-pair and batched) and the row-to-column compact gather. Env: EXABOOST_ROWDATA_4BIT (default on when eligible, =0 restores 8-bit bit-for-bit) and EXABOOST_ROWDATA_4BIT_VERIFY=1 (builds both representations and element-wise compares the unpacked values, Log::Fatal on mismatch). Sparse, 16/32-bit, >16-bin and large-bin-partition datasets are untouched. numerai scoreboard shape (5.43M x 2748, RTX 5090), measured: - 2000-tree train_s 87.58 -> 66.27 (-24%); 300-tree median 35.47 -> 32.64 ms - gather source fill kernel 11.64 -> 5.10 ms/tree (packed source at 1.47 TB/s), row-to-col gather 4.46 -> 4.01 ms/tree; construct unchanged (~21 ms/tree, latency-bound on scattered gradient/hessian reads, not bin bytes) - peak device memory 19.26 -> 11.47 GB; Booster create 2.25 -> 1.44 s - construct() 37.6s and RMSE/mean_leaves/predictions identical to pre-change (rmse 0.22367504 at 2000 trees, byte-equal pred head) Gates: EXABOOST_ROWDATA_4BIT_VERIFY passes on the full numerai shape and on a synthetic 400,001 x 1001 13-bin dataset (odd rows/columns/partition widths, feature_fraction 0.37); quant md5 identical for 4BIT=1/4BIT=0/ pre-change on the 1.5M-row int8-fed numerai slice (763c75c0d9cb...), also with EXABOOST_HYBRID_GROWTH=0 and EXABOOST_FAST_ROWDATA=0; covtype 1023/10 quant locks unchanged (GROWTH=1 5f4e7bdfff1e, GROWTH=0 fcb9f6c2ab87; covtype >16 bins proves the fallback path untouched). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EYUXNtCAXDd3zYT71m8YSJ
….3s) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…rai -10.6%/tree) Two follow-ups on the 4-bit packed compact path (numerai shape: 5.4M x 2748, ~5.5 bins/feature, feature_fraction 0.1, non-quant, hybrid growth): 1. EXABOOST_GH_INTERLEAVE (default on): BeforeTrain interleaves the per-row gradient/hessian arrays into a float2 buffer (0.05ms/tree at 5.4M rows); the batched dense construct kernels then touch ONE 32B sector per scattered row instead of two. Bit-identical values. Compile-time USE_GH2 template dispatch so the disabled instantiation stays identical to the historical kernel (a runtime branch on the pointer measurably slowed the kernel for both modes). Construct kernel 22.9 -> 22.7 ms/tree vs old tip (1.0ms vs the branchless interleave-off build). 2. EXABOOST_SPLIT_PACKED_READ (default on): the batched apply kernels read the split columns' bins straight from the 4-bit packed compact matrix (bit_type=4 descriptors carrying base byte, per-row byte stride, nibble shift), dropping the per-tree 1.5GB column-major gather (CUDARowToColCompactKernel, 4.2ms/tree). The scattered packed reads cost one 32B sector per row: gen-bit-vector 0.41 -> 0.78, split inner 0.21 -> 0.28 ms/tree; net -3.8ms/tree. Classic per-split Split() consumers (leaf-wise tail, forced splits, batched-apply env fallback) lazily run the skipped gather once per tree via EnsureClassicColumnView before the first classic split. Same-session A/B (RTX 5090, numerai 300-tree median): 34.36 (old tip) -> 30.72 ms/tree (-10.6%); kill-switches off: 34.91 (parity). Kernels/tree (60-tree nsys): construct 22.94 -> 22.70, row-to-col gather 4.18 -> 0, gen bitvec 0.41 -> 0.78, interleave fill +0.05. Full 2000-tree bench flow: 59.96s (official baseline 66.24s; wall clocks drift up to ~15% between sessions, same-session ratio 0.88). Gates: covtype 1023/10 quant GROWTH=1 md5 5f4e7bdfff1e, GROWTH=0 fcb9f6c2ab87 (unchanged); numerai quant int8-fed 1.5M-row slice md5 763c75c0d9cb with the new envs both on and off (quant hybrid takes the packed read against the full 4-bit matrix); numerai non-quant 100-tree RMSE 0.22366795, mean_leaves 32.000, predictions bit-identical envs on vs off; covtype non-quant on/off within its inherent run-to-run CUDA-atomic nondeterminism. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EYUXNtCAXDd3zYT71m8YSJ
…onstruction + EFB precheck Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds hybrid tree growth to the CUDA learner: whole levels are grown depth-wise-batched while the leaf budget cannot bind (provably leaf-wise-equivalent, since split decisions are node-local), with an exact leaf-wise tail once it might. All search-phase and apply-phase kernels are batched per level, and non-quantized training runs a single-synchronization-per-level pipeline.
Benchmark: 10 wins, 1 tie, 0 losses vs XGBoost 3.3.0 (RTX 5090, 500-tree aligned regimes — gbm-bench suite + Numerai v5 full data, median of 3; full report in
benchmarks/workspace):(airline pending: kt.ijs.si has been down)
Commits
Correctness
🤖 Generated with Claude Code
https://claude.ai/code/session_01EYUXNtCAXDd3zYT71m8YSJ