tandem rng tandem-cuda GitHub

Performance

pixi run -e cuda12 make bench produces the figures.

GPU

NVIDIA A100 40 GB (PCIe), CUDA 12.8 built by clang 19 as the nvcc host compiler with -std=c++20 (pixi run -e cuda12 make bench): 2^28 elements into device memory. Each row first runs its own fill for two seconds, then times 21 fills by cudaEvent. tools/bench prints the median and the fastest of the 21, and the table gives the median. In the last run the fastest lay at most 4 % above the median on every row. The fastest once hid an allocation that slowed most fills, see design. After the two seconds every fill runs at the card's 250 W power cap, so the rows give the capped rate. A row run alone (tools/bench <part of its name>) gives the figure it gives in the table. Rows that name a length write fewer elements. The GPU had no other process during the runs, and two consecutive runs agreed within 3 %. The elements are 2^28 of each type, so the rows for narrow types write fewer bytes.

The cuRAND column is Philox4x32-10 of cuRAND 10.3.9 from the same run, by the same method, through the host API with the call named. It writes the row's output type where cuRAND has one. Where it has none, the column gives the nearest call, marked "nearest": cuRAND has no 8-, 16- or 64-bit integer output for Philox, so curandGenerate writes the same bytes as 32-bit words, and it has no half floats, bools, bounded integers or exponentials. A normal row's start draw is a generator offset in cuRAND. The Thrust rows of cuRAND use its documented device API idiom, curand_init(seed, i, 0) per element.

GiB/s written cuRAND Philox4x32-10 cuRAND call
tandem::fill_u32, tile kernel (default for K >= 8) 1377 1285 curandGenerate
tandem::fill_u64, tile kernel 1389 1300 curandGenerate, nearest
tandem::fill_f32, tile kernel 1379 1262 curandGenerateUniform
tandem::fill_f64, tile kernel 1388 778 curandGenerateUniformDouble
tandem::fill_u32, direct kernel (K < 8) 1316 1285 curandGenerate
tandem::fill_u64, direct kernel 1320 1300 curandGenerate, nearest
tandem::fill_f32, direct kernel 1314 1262 curandGenerateUniform
tandem::fill_f64, direct kernel 1318 778 curandGenerateUniformDouble
tandem::fill_u16, tile kernel 1356 1282 curandGenerate, nearest
tandem::fill_f16_bits, tile kernel 1341 1282 curandGenerate, nearest
tandem::fill_u8, tile kernel 1320 1252 curandGenerate, nearest
tandem::fill_bool (one byte per bit) 1221 1252 curandGenerate, nearest
tandem::fill_u32_below(1000) 1332 1285 curandGenerate, nearest
tandem::fill_u64_below(1000) 1354 1300 curandGenerate, nearest
tandem::fill_u32_below(2^32 - 2) (rejects about every draw's threshold check) 1329 1285 curandGenerate, nearest
tandem::fill_u64_below(2^64 - 2) 1299 1300 curandGenerate, nearest
tandem::fill_u32_below with a low bound into int32_t 1329 1285 curandGenerate, nearest
tandem::fill_u32_below with a low bound into int64_t (8-byte elements) 1363 1300 curandGenerate, nearest
tandem::fill_u64_below with a low bound into int64_t 1302 1300 curandGenerate, nearest
tandem::fill_normal_f64 1064 567 curandGenerateNormalDouble
tandem::fill_normal_f64, start at an odd Float64 draw 985 554 curandGenerateNormalDouble, offset 1
tandem::fill_normal_f64, start at word 6 (draw 3) 946 548 curandGenerateNormalDouble, offset 3
tandem::fill_normal_f64, 2^24 elements, even and odd start 798, 788 528, 515 curandGenerateNormalDouble, offset 0 and 1
tandem::fill_normal_f64, 2^20 elements, even and odd start 238, 218 218, 201 curandGenerateNormalDouble, offset 0 and 1
tandem::fill_normal_f32 1219 873 curandGenerateNormal
tandem::fill_exponential_f64 955 778 curandGenerateUniformDouble, nearest
tandem::fill_exponential_f32 926 1262 curandGenerateUniform, nearest
thrust::transform of tandem::uniform<uint32_t> into a device vector 64 687 curand
thrust::transform of tandem::uniform<double> into a device vector 125 1245 curand_uniform_double

The tile kernel stages eight steps of 32 groups in 32 KiB of shared memory and writes them as 512 contiguous bytes per warp. It runs at the card's memory bandwidth, about 1.5 TB/s. The direct kernel stores each block straight from registers, so the eight threads of a group cover one 128-byte line per step. Both kernels pick a 16-byte vector store at compile time when the output's blocks are 16-byte aligned. The bool fill expands each bit to a byte, so it stages one step of 32 groups (32 KiB) and writes 1024 contiguous bytes per group. The bounded fills run at the fill speed, because a rejection is rare and its retry runs out of line. A 32-bit range into 8-byte elements gives each thread one 16-byte output slot from two draws, so a warp still writes 512 contiguous bytes. Two stores per 16-byte block of draws gave 955 GiB/s. The f64 ziggurat runs a table pass and a second kernel for the 0.43 % of misses, see design. At 2^24 elements it reaches three quarters of the 2^28 rate: the table pass alone runs at 1168 GiB/s, and the misses kernel takes 25 µs of the 150. The Box-Muller fill it replaced ran at 833 and 679 GiB/s. A start at word 6 takes the octet stores of tandem-sycl, which keep each group's 128 bytes on whole sectors, see design. Before them it ran at 689.

The normal and exponential fills run at 67 % to 88 % of the uniform rate. Every fill draws the card's 250 W, so the SM clock sets each fill's rate once its arithmetic per byte is high enough. The uniform fills run at 1170 to 1260 MHz and are bound by memory. The f32 normal, the exponentials and the f64 normal run at 1035 to 1110 MHz, where the arithmetic per element bounds them: a logarithm, a square root, a sine and a cosine per f32 normal pair, the specification's f64 logarithm, about 22 f64 operations per element, the f32 exponential's two-float logarithm, about 20 f32 operations and a reciprocal per element, or the ziggurat's table read and its misses kernel, which takes 16 % of the f64 normal fill. The f32 exponential spends 10 of its operations on keeping every draw within 0.571 ulp, so that 1 - exp(-x) maps back to the draw's own 2^-24 grid point. See design for the costs that were removed.

cuRAND is faster on five rows. On the bool, the 2^64 - 2 bounded and the f32 exponential rows its nearest call does less work than the row: no expansion of bits to bytes, no bounding, no logarithm. The Thrust rows are slower in these measurements: element i of a Tandem stream needs its chunk's seeding, eight rounds of the step, followed by 1 through K steps, while a Philox element is one ten-round block of its counter. Across complete groups in this contiguous benchmark, K = 32 averages 24.5 steps per requested element. Independent indexed calls repeat that work for each word or pair of words in a block. The fill kernels share it across outputs, so use a fill when materialising a contiguous array.