Performance
pixi run -e cuda bench-cuda produces the figures.
GPU
NVIDIA A100 40 GB (PCIe), driver 570, AdaptiveCpp 25.10.0 for CUDA
12.9 (pixi run -e cuda bench-cuda): 2^28 elements into USM
device memory, by tandem-cuda's method. Each row first runs its own fill
for two seconds, then times 21 fills, each to the queue's wait.
build/cuda/bench prints the median and the fastest of the
21, and the table gives the median. 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 (build/cuda/bench <part of its label>)
gives the figure it gives in the table, within 3 %. Two runs agreed
within 3 %. The GPU had no other process during the runs. The
tandem-cuda column is the figure of its speed
page on the same card by the same method. The cuRAND column is
Philox4x32-10 of cuRAND 10.3.10 in the same run, by the same method,
through the CUDA runtime that AdaptiveCpp's CUDA backend uses. Where
cuRAND has no such output the column gives the nearest call, marked
"nearest": it has no 8-, 16- or 64-bit integers for Philox, so
curandGenerate writes the same bytes as 32-bit words, and
no half floats, bools, bounded integers or exponentials. A normal row's
start draw is a generator offset in cuRAND. Narrow types write fewer
bytes for the same element count.
| GiB/s written | tandem-sycl | tandem-cuda | cuRAND Philox4x32-10 | cuRAND call |
|---|---|---|---|---|
fill uint32_t, tile kernel (default on
GPUs for K >= 8) |
1361 | 1377 | 1255 | curandGenerate |
fill uint64_t, tile kernel |
1381 | 1389 | 1295 | curandGenerate, nearest |
fill float, tile kernel |
1359 | 1379 | 1245 | curandGenerateUniform |
fill double, tile kernel |
1378 | 1388 | 778 | curandGenerateUniformDouble |
fill uint32_t, chunk kernel |
1280 | 1316 | 1255 | curandGenerate |
fill uint64_t, chunk kernel |
1295 | 1320 | 1295 | curandGenerate, nearest |
fill float, chunk kernel |
1279 | 1314 | 1245 | curandGenerateUniform |
fill double, chunk kernel |
1291 | 1318 | 778 | curandGenerateUniformDouble |
fill uint16_t |
1325 | 1356 | 1222 | curandGenerate, nearest |
fill sycl::half (tandem-cuda:
fill_f16_bits) |
1236 | 1341 | 1222 | curandGenerate, nearest |
fill uint8_t |
1125 | 1320 | 1155 | curandGenerate, nearest |
fill bool (one byte per bit) |
1187 | 1221 | 1155 | curandGenerate, nearest |
fill std::complex<double> |
1387 | 788 | curandGenerateUniformDouble |
|
fill_below uint32_t, range 1000 |
1359 | 1332 | 1255 | curandGenerate, nearest |
fill_below uint64_t, range 1000, 32-bit
draws (tandem-cuda: fill_u32_below into
int64_t) |
1375 | 1363 | 1285 | curandGenerate, nearest |
fill_below uint32_t, range 1000, chunk
kernel |
1317 | 1255 | curandGenerate, nearest |
|
fill_normal float |
1160 | 1219 | 858 | curandGenerateNormal |
fill_normal float, start at an odd Float32
draw |
1063 | 737 | curandGenerateNormal, offset 1 |
|
fill_normal float, start at word 4 (first
block at lane 1) |
1093 | 806 | curandGenerateNormal, offset 4 |
|
fill_normal double |
1066 | 1064 | 566 | curandGenerateNormalDouble |
fill_normal double, start at an odd
Float64 draw |
993 | 985 | 547 | curandGenerateNormalDouble, offset 1 |
fill_normal double, start at word 6 |
996 | 946 | 548 | curandGenerateNormalDouble, offset 3 |
fill_exponential float |
829 | 926 | 1232 | curandGenerateUniform, nearest |
fill_exponential double |
874 | 955 | 778 | curandGenerateUniformDouble, nearest |
The uniform and bounded fills run near the card's memory bandwidth, as in tandem-cuda, except the 8- and 16-bit fills, 2 % to 15 % below tandem-cuda's. The float normal fills from an odd draw and from word 4 cost 6 % to 9 % against a start at a multiple of 32 words. The double normal rows are the ziggurat with tandem-cuda's two kernels. Their table pass stores in octets, which keeps a start at word 6 at the rate of an odd start. When each lane stored its own block, word 6 put every group's 128 bytes 16 bytes off the 128-byte lines, touched five 32-byte sectors per group and step instead of four, and ran at 708 GiB/s by the earlier method.
The normal and exponential fills run below the uniforms because the
card holds 250 W: their arithmetic lowers its clock until it bounds
them, see tandem-cuda's design.
The f64 exponential takes tandem-cuda's folded logarithm, the same bits
with two products fewer. The f32 exponential takes tandem-cuda e98daee's
two-float logarithm, 10 f32 operations per draw more than the
single-float one before it. tandem-cuda's device f32 division and its
f32 normal's square root are PTX, which AdaptiveCpp's generic target
does not compile, so the f32 exponential keeps the IEEE division and
runs at 829 GiB/s against tandem-cuda's 926. A build for the
cuda:sm_80 target took the PTX division but ran the f32
normal from an odd draw at 889, so the generic target stays. The f32
exponential and uint64_t bounded rows come from a later run
than the rest, on the same card by the same method.
fill_below into uint64_t with a range up to
2^32 takes 32-bit draws. The tile kernel's write phase hands each work
item half a block, two elements in one 16-byte store, as tandem-cuda
does. With a whole block per work item, two stores 32 bytes apart across
work items, the fill ran at 558 GiB/s, and with the half blocks but the
rejection retry inlined beside the store, at 1120. cuRAND leads the
uint8_t and f32 exponential rows, where its nearest call
does less work: 32-bit words where the fill stores single bytes, and
uniforms without the logarithm.