SyclReproSum 1.3.0

SyclReproSum 1.3.0

SyclReproSum 1.3.0 is a performance-focused release for reproducible scalar
and cumulative summation. It reduces synchronization, temporary-storage
traffic, and scan overhead while preserving the library's central guarantee:
results remain bit-identical across input permutations, work-group sizes,
execution schedules, and validated devices.

The library remains header-only, C++17, and source-compatible with version
1.2.6. The public adn::sum, adn::cumsum, and
adn::validate_environment APIs are unchanged.

Highlights

• Improve large default cumulative sums by processing wider contiguous chunks
  and reducing the number of tile totals and recursive scan operations.
• Add a single-tile cumulative-sum fast path for CUDA and Level Zero devices.
• Store CUDA scalar-reduction partials in device USM and copy back only the
  final scalar.
• Use a smaller persistent reduction grid for large default sums on Intel
  GPUs.
• Add built-in small-N latency benchmarks against native SYCL and oneDPL
  baselines.
• Expand correctness coverage for every optimized execution path.
• Refresh all documented benchmark results.
• Update the encoded library version to 1.3.0.

Large cumulative-sum optimization

The default cumulative-sum configuration previously processed eight
contiguous elements per work-item, giving a 2,048-element tile for
WG_SIZE=256.

For inputs of at least 67,108,864 elements, version 1.3.0 processes 16 elements
per work-item. This doubles the default tile size to 4,096 elements and
approximately halves:

• The number of tile totals.
• Temporary binned-prefix storage.
• Work performed by the recursive tile scan.
• Global traffic associated with tile-prefix propagation.

The selection depends only on N, K, and WG_SIZE. It does not inspect the
backend, vendor, GPU model, or device properties. Smaller inputs and custom
fold or work-group configurations retain the previous eight-item path.

Observed 100M-element improvements

The following figures compare the published 1.3.0 ADN throughput with the
published 1.2.6 results. Throughput is measured in GElements/s; percentages
are calculated from the rounded documented values.

Device	double cumsum	float cumsum
NVIDIA GB10	2.524, +6.8%	3.500, +2.5%
NVIDIA RTX PRO 4500	5.304, +9.2%	6.815, +1.8%
AMD Radeon Pro W7500	0.818, +8.9%	1.181, +8.4%
Intel Arc Pro B70	3.926, +26.6%	7.181, +20.2%
Intel Core Ultra 7 265 iGPU	0.191, +26.5%	0.357, +24.4%
Intel Core Ultra 7 265 CPU	0.144, +24.1%	0.173, +24.5%

Every documented device and precision recorded higher 100M-element
cumulative-sum throughput. These measurements are observations from the
tested systems, not performance guarantees for other hardware or runtime
versions.

Single-tile cumulative sums

CUDA and Level Zero devices now avoid the preliminary tile-total kernel and
recursive tile scan when the complete input fits in one tile.

With the defaults K=3 and WG_SIZE=256, this path applies through N=2048.
The incoming binned prefix is initialized directly to zero, after which the
final prefix kernel performs the same ordered binned accumulation and
conversion as before.

The N=2049 benchmark is intentionally retained as the first multi-tile
control.

Scalar-sum improvements

CUDA reductions now keep partial binned accumulators in device USM rather than
host-visible shared USM. Only the final floating-point scalar is copied back
to the host. This reduces host-visible memory traffic and can make adn::sum
faster than a plain sycl::reduction for several small double input sizes on
the tested CUDA systems.

Large default reductions on Intel GPUs use a maximum of 256 persistent
work-groups instead of 2,048. This applies only to the default K=3,
WG_SIZE=256 specialization once more than 2,048 groups would otherwise be
required. Other devices and custom specializations retain the original grid
limit.

The Intel selection uses the backend, device type, and Intel vendor
identifier. No individual GPU model IDs are hard-coded.

Scalar performance remains workload- and device-dependent; this release does
not claim that reproducible summation universally outperforms conventional
non-reproducible reduction.

Built-in small-N latency benchmarks

Version 1.3.0 adds latency benchmarks at:

• N=1
• N=64
• N=1,024
• N=2,048
• N=2,049

The baselines are the same algorithmic baselines used by the 100M-element
benchmarks:

• Plain non-reproducible sycl::reduction for adn::sum.
• oneDPL inclusive_scan for adn::cumsum.

Each test performs three warmup calls followed by 201 synchronous measurements
in alternating baseline/ADN order and reports the median latency. The README
tables use the median from five separate test-binary invocations.

Input allocation, initial transfer, and warmup are excluded. Internal ADN
allocation and synchronization remain included.

The baseline is intentionally not the previous SyclReproSum release.
Version-to-version A/B measurements remain useful for regression analysis,
but they do not describe the cost of reproducibility relative to a
conventional SYCL implementation.

Reproducibility and numerical behavior

The accumulator primitives and their scalar arithmetic order are unchanged:

• update
• deposit
• renorm
• merge
• conv

The optimized paths only change storage placement, launch geometry, tile
partitioning, or eliminate redundant passes. Intermediate values continue to
be represented as binned accumulators and combined through exact reproducible
merges.

No FMA contraction, reassociated scalar addition, relaxed rounding, or
order-dependent shortcut was introduced.

Within adn::max_reproducible_count<T>:

• Scalar sums remain independent of input order, work-group size, scheduling,
  and validated device.
• Every cumulative prefix remains bit-identical across supported work-group
  sizes and validated devices for a fixed input sequence.
• The final cumulative prefix remains bit-identical to the corresponding
  scalar sum.
• Canonical NaN handling, infinity propagation, signed zero, subnormal
  behavior, fold-count behavior, and capacity limits are unchanged.

Expanded correctness coverage

Four correctness cases were added per selected device:

• Verify default sum grid selection and specialization isolation.
• Verify shuffled large reductions remain bit-identical across tuned and
  untuned work-group configurations.
• Compare the single-tile cumulative path with a forced multi-tile path over
  complete output arrays.
• Compare eight-item and 16-item cumulative tiles for both float and
  double, including shuffled cancellation-heavy inputs.

The suite now contains:

• 189 correctness cases per selected device.
• Four small-N latency benchmarks per device.
• Four 100M-element throughput benchmarks per device.
• 11 cross-device reproducibility cases.
• One version test.

This produces 800 tests on the four-device validation system.

Release validation

All binaries were rebuilt from the same 43acfbf source snapshot before
testing.

Validation system	Selected devices	Result
CUDA + HIP + SPIR-V system	RTX PRO 4500, Radeon Pro W7500, Intel iGPU, Intel CPU	800/800 passed
NVIDIA GB10 CUDA system	NVIDIA GB10	197 passed, 12 expected skips
Intel Arc Pro B70 SPIR-V system	Arc Pro B70, Intel iGPU, Intel CPU	603/603 passed

The GB10 skips consist of cross-device cases on a single-device system and one
SSE-specific host denormal-mode test unavailable on ARM.

An additional B70-only run, including small-N and 100M benchmarks, completed
with 198 passed tests, 11 expected cross-device skips, and exit status zero.

Intel multi-device runtime note

On the B70 validation host, a process simultaneously exercising the Level Zero
GPUs and OpenCL CPU may print Intel's generic crash-report diagnostic during
ADNSumCrossDevice.Float_WideRange.

The test still completes successfully, all 11 cross-device tests pass, and the
process exits with status zero. Complete B70-only, iGPU-only, and CPU-only runs
do not reproduce the message. No incorrect result, SYCL exception, or Google
Test failure was observed; current evidence limits it to the Intel
multi-device compiler/runtime interaction.

Compatibility and upgrade notes

• Existing public API calls require no source changes.
• The library remains header-only and requires no additional link step.
• Custom K and WG_SIZE configurations retain their previous large-input
  tile and reduction-grid behavior.
• CUDA reductions now use device USM for internal partial accumulators; shared
  USM remains required for environment validation.
• Applications should replace repro_sum.hpp and recompile.
• Version macros now report 1.3.0, encoded as 10300.

Changes since 1.2.6

• Move CUDA sum partial accumulators to device USM.
• Copy only the final CUDA sum result back to the host.
• Add the CUDA and Level Zero single-tile cumulative-sum path.
• Tune the large default Intel GPU reduction grid.
• Add the generic 16-item large cumulative-sum tile.
• Add reproducibility tests for every optimized path.
• Add native SYCL and oneDPL small-N latency benchmarks.
• Expand the suite from four to eight performance benchmarks per device.
• Refresh performance documentation for all validated devices.
• Update the library version from 1.2.6 to 1.3.0.

Full comparison:

v1.2.6...v1.3.0

Attribution and license

SyclReproSum is licensed under the BSD 3-Clause License. Portions of the
binned accumulator primitives are adapted from ReproBLAS, Copyright 2016,
University of California. See THIRD_PARTY_NOTICES for the complete
attribution and license terms.