silvic-gpu

The Metal and CUDA backends of silvic-core. The kernels are written once, against a small set of macros that one prelude defines for Metal Shading Language and another for CUDA C++, and compiled when a backend starts: Metal with fast math off, NVRTC with --fmad=false. GpuBackend implements the core’s Backend once over the device primitives, so the two backends differ only in how they allocate, copy, compile and launch.

let mut backend = silvic_gpu::metal::backend(None)?;
let model = silvic_core::train(&params, &train_set, &[], &mut backend, &mut LocalCollective, &mut |_| Ok(()))?;

Public API

metal::backend(index) feature metal, macOS

A compiled backend on Metal device index, or the system default device. Buffers use shared storage; each sync commits one command buffer.

cuda::backend(ordinal) feature cuda

A backend on CUDA device ordinal, default 0, through cudarc with the driver and NVRTC loaded at run time, on one stream. It builds without a CUDA toolkit.

metal::is_available(), cuda::is_available()

Whether the device is there. Neither panics: a missing or mismatched CUDA driver is a false here and a device error from cuda::backend, never an abort.

GpuBackend::new(device), trait GpuDevice

The backend over any device that can allocate, upload, download, compile, launch and sync.

Also public MetalDevice, CudaDevice, Arg, and kernel_source(prelude), the source a device compiles.

On the device

The training rows, their gradients, scores and the row order of every leaf stay on the device; a split brings back only its histogram and the count of rows going left.

Gradients

L2 and binary per row; lambdarank with one threadgroup per query of up to 1024 rows, larger queries on the host.

Quantization and bagging

Bit for bit what the CPU computes, from the same counter-based hash.

Splits

A stable partition of the leaf and the histogram of the child with fewer rows, in one submission and one wait.

Histograms

Quantized gradient and hessian packed into one 32-bit integer per row, added with one atomic per feature into threadgroup memory, then flushed to the device histogram. Float histograms as well, without the exactness.

Leaves and scores

Exact fixed-point gradient sums per leaf for renewal, and each new tree added to every dataset’s scores.

Split search stays on the host, where every rank searches the same reduced histogram. On the Metal backend, quantization, histograms, partitions, leaf sums and score updates equal the CPU backend’s bit for bit, and quantized L2 models are byte-identical; binary and lambdarank gradients agree to 1e-5, Metal computing them in f32.

CUDA

The CUDA backend compiles, its kernel source passes clang’s CUDA front end, and it never panics on a missing driver. Training has been benchmarked on an NVIDIA Tesla T4; see the training benchmark for the workload and timings.