SyclReproSum 1.4.0

SyclReproSum 1.4.0

SyclReproSum 1.4.0 improves repeated scalar and cumulative summation through
reusable scratch storage. It reduces allocation and initialization overhead
while preserving the library's reproducibility guarantees.

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

Highlights

• Reuse scalar-reduction scratch buffers across calls on each host thread.
• Reuse the zero-prefix buffer for single-tile CUDA and Level Zero scans.
• Handle context changes, concurrent threads, reentrant calls, exceptions,
  and cache teardown.
• Add nine correctness cases per selected device.
• Add a CUDA memory initialization check target.
• Improve AMD libspirv discovery across compiler installation layouts.
• Refresh all README performance tables from complete test binaries.
• Update the library version to 1.4.0, encoded as 10400.

Reusable reduction storage

adn::sum now retains one reduction buffer per host thread and (T, K)
combination. Repeated calls on the same context and device reuse that
allocation.

The buffer grows when a larger reduction grid is needed and is replaced
when the context or device changes. Each buffer holds at most 2,048 binned
accumulators. Every call overwrites all partials it subsequently reads.

Concurrent host threads use separate buffers. Reentrant calls use an
independent temporary allocation while the outer call holds the cached
buffer. Exception handling releases the lease so later calls can reuse
the cache.

Single-tile cumulative sums

The existing CUDA and Level Zero single-tile path now retains a separate
zero-prefix buffer. Initialization occurs when the buffer is allocated,
removing repeated allocation and zero-initialization work from subsequent
calls on the same context and device.

With the defaults K=3 and WG_SIZE=256, this path applies through
N=2048. The N=2049 benchmark remains the first multi-tile control.

Single-tile and multi-tile execution share the prefix-writing kernel.
Input dependency events remain honored, and calls remain synchronous.

Performance measurements

All README performance tables were remeasured, including ADN timings,
the corresponding algorithm baselines, and their relative performance.

Before/after comparisons used complete test binaries with the same compiler,
oneDPL headers, and compilation options within each validation system.

For the measured single-tile cumsum sizes from N=1 through N=2048,
the observed latency reductions across float and double were:

Device	Approximate latency reduction
NVIDIA GB10	12% to 35%
NVIDIA RTX PRO 4500 Blackwell	11% to 35%
Intel Arc Pro B70	29% to 46%
Intel Core Ultra 7 265 iGPU	26% to 47%

The primary matched measurements also showed approximately 20% higher
double sum throughput and 67% higher float sum throughput on the
Intel Arc Pro B70 at 100M elements.

These comparisons use newly measured results for both implementations.
Large-input cumsum results varied in both directions across repeated
comparisons; the improvements above do not imply that every workload
became faster.

The README continues to compare adn::sum with plain sycl::reduction
and adn::cumsum with oneDPL inclusive_scan. Its tables contain the
complete primary measurements rather than selected faster follow-up runs.

Storage lifetime and shutdown behavior

Cached buffers remain allocated between calls and retain their associated
SYCL context. They are released when replaced or when their owning thread
exits.

Calls made after cache destruction use temporary storage that is released
before returning. This avoids accessing or recreating a destroyed cache.
The queue and SYCL runtime must still be usable during such calls.

Applications using many host threads or accumulator specializations should
account for the retained storage.

Reproducibility and numerical behavior

The update, deposit, renorm, merge, and conv primitives are
unchanged. This release preserves their scalar arithmetic order.

Within the documented accumulator capacity:

• Scalar sums remain bit-identical across input permutations, supported
  work-group sizes, and validated devices.
• For a fixed input sequence, cumulative prefixes remain bit-identical
  across supported work-group sizes and validated devices.
• The final cumulative prefix remains bit-identical to the scalar sum.
• Existing NaN, infinity, signed-zero, subnormal, and capacity semantics
  remain unchanged.

A shared 1M-element wide-range input was checked across seven selected
CPU and GPU devices. Both implementations produced
0x430FC878C605717F, including after input shuffling and a work-group-size
change. The result also matched an independent math.fsum reference.

Expanded correctness coverage

Nine additional correctness cases per selected device cover:

• Scratch reuse across growing and shrinking inputs, including shuffles.
• Context and queue changes.
• Concurrent calls from multiple host threads.
• Reentrancy and recovery after exceptions.
• Single-tile in-place scans and output bounds.
• Scratch access during thread and process teardown.

The suite now contains 198 correctness cases and eight benchmarks per
selected device, plus 11 cross-device cases and one version test.

This gives 836 tests on a system with three selected GPUs and one CPU.

Build and diagnostic tooling

CUDA initialization checks

The new target runs CUDA correctness tests under Compute Sanitizer:

make test-initcheck

To focus on scratch reuse and single-tile initialization:

make test-initcheck INITCHECK_FILTER='*Scratch*:*SingleTile*'

The default configuration checks global and shared memory initialization.
Older Compute Sanitizer versions without --initcheck-address-space
support can use INITCHECK_FLAGS= for global-memory checks.

AMD device-library discovery

The Makefile now searches these compiler resource directories for AMD
libspirv, in order:

• lib/amdgcn-amd-amdhsa
• lib/amdgpu-amd-amdhsa
• lib/amdgcn-amd-amdhsa-llvm

This supports both current and legacy compiler layouts. AMD_LIBSPIRV
remains available as an explicit override.

Validation

The optimized implementation completed correctness validation on three
systems:

Selected devices	Correctness results, excluding benchmarks
RTX PRO 4500, Radeon Pro W7500, Core Ultra 7 265 iGPU and CPU	804 passed
NVIDIA GB10	198 passed, 12 expected skips
Arc Pro B70 and Core Ultra 7 265K CPU	408 passed

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

All selected-device latency and throughput benchmark cases passed.
The final version-only update passed the version test separately.
Formatting and whitespace checks also passed.

Compatibility and upgrade notes

• Existing public API calls require no source changes.
• Replace repro_sum.hpp and recompile.
• Scratch storage may now persist between calls and retain its context
  until replacement or thread exit.
• Synchronous completion and dependency-event handling are preserved.
• Version macros report 1.4.0, encoded as 10400.

Changes since 1.3.0

Attribution and license

The project continues under the BSD 3-Clause License. Its binned primitives
include adaptations of ReproBLAS, copyright 2016, University of California.
The complete attribution and applicable terms remain in
THIRD_PARTY_NOTICES.