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.

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.

SyclReproSum 1.2.6

SyclReproSum 1.2.6 fixes a reproducibility issue in generic SPIR float
conversion. On CPU OpenCL, some high-fold sums and cumulative sums could differ
by one ULP between work-group sizes.

The library remains header-only, C++17, and source-compatible with version
1.2.5.

Highlights

• Guard generic SPIR float conversion against the observed CPU OpenCL JIT
  rounding difference.
• Preserve the existing arithmetic order and native NVIDIA and AMD paths.
• Add regression coverage for sums and cumulative sums across work-group sizes,
  fold counts, input permutations, scaling, in-place scans, and devices.
• Update the encoded library version to 1.2.6.
• Refresh the README benchmark results.

Testing and validation

• All 768 tests passed on the four-device NVIDIA, AMD, Intel GPU, and Intel CPU
  system.
• All 567 correctness tests passed on the Intel Arc Pro B70, Intel iGPU, and
  Intel CPU system.
• Hostile and unsafe floating-point environment validation passed.

Compatibility and upgrade notes

• Existing adn::sum, adn::cumsum, and adn::validate_environment calls
  require no changes.
• Recompile applications after updating repro_sum.hpp.
• Native NVIDIA and AMD conversion paths are unchanged.

Full comparison:
v1.2.5...v1.2.6

Attribution and license

SyclReproSum is licensed under the BSD 3-Clause License. Portions of the binned
accumulator primitives are adapted from ReproBLAS. See THIRD_PARTY_NOTICES
for details.

SyclReproSum 1.2.5

SyclReproSum 1.2.5 fixes two correctness hazards: an observed CPU OpenCL JIT
miscompilation in high-fold double conversion, and an environment-probe cache
collision between operation specializations compiled in different translation
units.

The library remains header-only, source-compatible with version 1.2.4, C++17,
and based on the Ahrens-Demmel-Nguyen (ADN) binned floating-point format. The
binned primitives remain adaptations of the ReproBLAS reference implementation.

Highlights

• Guard generic SPIR double conversion against an observed CPU OpenCL JIT
  miscompilation.
• Keep native NVPTX and AMDGCN sum and cumulative-sum operation kernels
  unchanged.
• Isolate device arithmetic probes and their caches by operation, floating-point
  type, fold count, and work-group size.
• Add a mixed-translation-unit regression test that combines strict and unsafe
  device floating-point options.
• Add high-fold conversion-order coverage for scalar and cumulative sums across
  input permutations and small work-group sizes.
• Update the encoded library version to 1.2.5.

SPIR double-conversion workaround

Some CPU OpenCL runtimes can miscompile the double accumulator conversion loop
when consuming generic SPIR device code. The failure was observed with higher
fold counts and small work-groups and can produce an incorrect final value after
large-magnitude cancellation.

Version 1.2.5 makes the running double conversion value volatile only in generic
SPIR device images. The arithmetic sequence, bin contents, scaling rules, and
rounding order are unchanged. Native NVIDIA and AMD device images retain the
ordinary non-volatile conversion loop, as does every float conversion.

The new regression coverage exercises:

• Double fold counts 16, 32, and 52.
• Float fold count 21.
• Work-group sizes 2, 4, and 64.
• All permutations of a large cancelling pair and a retained signal.
• Inputs spanning work-item, tile, and recursive-scan boundaries.
• Near-overflow scaled conversion terms.
• Bit-identical scalar sums and cumulative prefixes.

Both adn::sum and adn::cumsum use the corrected conversion path.

Translation-unit-safe environment checks

Previous releases named the device arithmetic probe and cached its successful
result only by backend, device, and floating-point type. In a program containing
translation units built with different device floating-point options, a strict
translation unit could populate that cache before an operation from an unsafe
translation unit ran. The later operation could then skip the arithmetic probe.

Version 1.2.5 gives each sum and cumsum specialization its own probe identity
and cache. The identity includes the operation, floating-point type, K, and
WG_SIZE. An explicit adn::validate_environment<T>(q) call retains its own
cache and cannot suppress the check required by a later operation
specialization.

A new validation target compiles one translation unit with strict defaults and
another with -fdenormal-fp-math=positive-zero, links them into one application,
and verifies that every otherwise supported device rejects the unsafe
operation. The test fails against version 1.2.4 and passes with this release on
Intel, NVIDIA, and AMD devices.

Performance impact

The accumulator update, deposit, renormalization, merge, reduction, scan, and
operation-kernel conversion sequences are unchanged by the validation-cache
fix. Comparison of the common operation kernels found:

• Identical NVIDIA SASS.
• Identical NVIDIA PTX instruction sequences, apart from compiler-generated
  local label numbering.
• Identical AMDGPU machine-code bytes.

Each new operation specialization performs its arithmetic probe once per device
and then uses the cached result. The normal throughput benchmark performs a
warmup call before timing, so this first-use check is outside the reported
measurements. Applications that instantiate many K and WG_SIZE combinations
may emit more small probe kernels and perform more one-time checks; steady-state
sum and cumulative-sum throughput is unaffected.

The README clarifies the three-invocation benchmark method and refreshes the
SPIR measurements and toolchain metadata. Published native NVIDIA and AMD
throughput values remain applicable because their timed operation kernels are
unchanged.

Testing and validation

The complete suite now contains:

• 179 correctness cases and four throughput benchmarks per selected device.
• Ten cross-device reproducibility cases.
• One version test.
• 743 total Google Tests on the four-device validation system.
• One standalone mixed-translation-unit validation test.

Release validation covered:

• NVIDIA RTX PRO 4500 Blackwell via CUDA (sm_120).
• AMD Radeon Pro W7500 via HIP (gfx1102).
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.
• Additional randomized sum and cumulative-sum checks on NVIDIA GB10 and Intel
  Arc Pro B70 systems.

Validation results:

• CUDA, SPIR-V, and AMDGCN example and test fat binaries built successfully.
• All 743 tests passed, including sixteen 100-million-element benchmarks.
• All 727 non-benchmark tests passed under hostile host floating-point modes.
• Denormal and signed-zero unsafe-device variants passed all expected rejection
  tests.
• The compile-time unsafe evaluation-mode rejection check passed.
• The mixed-translation-unit regression passed on all enabled backends.
• Clang-format 18, ASCII, and release-diff whitespace checks passed.

Compatibility and upgrade notes

• Version 1.2.5 is source-compatible with version 1.2.4.
• Existing adn::sum, adn::cumsum, and adn::validate_environment calls
  require no changes.
• Recompile applications after updating the header.
• Supported K, WG_SIZE, pointer, dependency-event, and backend interfaces are
  unchanged.
• Native NVIDIA and AMD steady-state operation performance is unchanged.
• Explicit environment preflight remains available, while every operation
  specialization performs and caches its own first-use check.

Changes since 1.2.4

• Add the generic-SPIR double-conversion JIT workaround.
• Preserve the native NVIDIA and AMD conversion paths.
• Add high-fold and small-work-group conversion-order regression tests.
• Cover scalar sums, cumulative sums, permutations, and scaled terms.
• Isolate environment probes and caches by operation specialization.
• Add mixed-translation-unit unsafe-device validation.
• Clarify benchmark methodology and refresh SPIR benchmark results.
• Update the encoded library version from 1.2.4 to 1.2.5.

Full comparison:
v1.2.4...v1.2.5

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.

SyclReproSum 1.2.4

SyclReproSum 1.2.4 works around an Intel IGC miscompilation affecting
two-fold (K=2) binned-accumulator merges. It also moves AMD libspirv
compatibility handling into a project-local Clang resource overlay and adds
Intel Arc Pro B70 benchmark results.

The library remains header-only, source-compatible with version 1.2.3,
C++17, and based on the Ahrens-Demmel-Nguyen (ADN) binned floating-point
format. The binned primitives remain adaptations of the ReproBLAS reference
implementation.

Highlights

• Work around an Intel IGC K=2 merge miscompilation with a loop-free
  specialization.
• Add double- and single-precision regression coverage across shuffled
  inputs and multiple work-group sizes.
• Create the AMD libspirv compatibility layout under the project build
  directory instead of modifying the compiler installation.
• Re-emit the AMD signed-character libspirv bitcode with the expected
  AMDGCN target triple.
• Add Intel Arc Pro B70 throughput and cross-device reproducibility results.
• Update the encoded library version to 1.2.4.

Intel IGC K=2 merge workaround

Intel IGC's PropagateCmpUniformity pass could miscompile the generic
two-iteration merge loop used by K=2 accumulators, merging bin zero twice
and skipping bin one.

Version 1.2.4 provides an explicit, loop-free K=2 specialization. It
preserves the original bin additions, carry updates, evaluation order, and
final renormalization. The generic merge path for K > 2 is unchanged.

New double- and single-precision regression tests cover:

• The minimal 1 + 1 reduction.
• Cancellation-heavy inputs with a small residual signal.
• Multiple shuffled input orders.
• Work-group sizes 64, 128, and 256.
• Bit-identical results across all tested configurations.

The workaround does not introduce FMA, reassociation, or order-dependent
rounding.

Project-local AMD resource overlay

Previous releases created an AMD libspirv compatibility alias inside the
selected compiler's resource directory.

Version 1.2.4 instead creates a configuration-specific overlay under:

build/clang-resource-<configuration-key>

The overlay reuses the compiler's normal resources while re-emitting
libspirv.l64.signed_char.bc with the expected amdgcn-amd-amdhsa target
triple. It is selected automatically for AMD-enabled builds and removed by
make clean.

This avoids modifying the DPC++ installation and requires write access only
to the project build directory. Existing AMD_GPU_ARCH and AMD_LIBSPIRV
overrides remain available.

Intel Arc Pro B70 results

The README now includes Intel Arc Pro B70 Level Zero measurements using
100-million-element inputs.

Operation	Double ADN / baseline	Float ADN / baseline
adn::sum	21.1 / 52.6 (2.5x slower)	47.4 / 85.5 (1.8x slower)
adn::cumsum	2.945 / 18.047 (6.1x slower)	5.806 / 40.862 (7.0x slower)

The documented wide-range cross-device sum produces
0x430FC878C605717F on the Intel Arc Pro B70, matching the NVIDIA, AMD,
Intel iGPU, and Intel CPU results, including after input shuffling.

Reproducibility and numerical behavior

Version 1.2.4 restores the intended K=2 behavior on affected Intel IGC
versions without changing the public API.

The accumulator bin layout, deposit, renormalization, and conversion paths
are unchanged. Within the documented capacity:

• Scalar sums remain independent of input permutation, work-group size,
  scheduling, and validated device.
• Cumulative prefixes remain bit-identical across runs, supported
  work-group sizes, and validated devices for a fixed input sequence.
• The last cumulative output remains bit-identical to the corresponding
  scalar sum.
• NaN, infinity, signed-zero, subnormal, and capacity behavior is unchanged.

Testing and validation

The complete suite now contains:

• 176 correctness cases and four throughput benchmarks per selected device.
• Ten cross-device reproducibility cases.
• One version test.
• 731 total tests on the four-device validation system.

Validation covered:

• NVIDIA RTX PRO 4500 Blackwell via CUDA (sm_120).
• AMD Radeon Pro W7500 via HIP (gfx1102).
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.

Release validation results:

• CUDA, SPIR-V, and AMDGCN example and test fat binaries built successfully.
• The example ran successfully.
• All 715 non-benchmark tests passed.
• The same 715 tests passed under hostile host floating-point modes.
• Both unsafe-device variants passed all expected rejection tests.
• The compile-time unsafe evaluation-mode rejection check passed.
• All 731 tests passed, including sixteen 100-million-element benchmarks.
• The GitHub clang-format workflow and release-diff whitespace checks passed.

Compatibility and upgrade notes

• Version 1.2.4 is source-compatible with version 1.2.3.
• Existing adn::sum, adn::cumsum, and adn::validate_environment calls
  require no changes.
• Users running K=2 reductions on Intel GPUs should upgrade.
• K > 2 merge behavior is unchanged.
• AMD-enabled builds no longer modify the compiler resource directory.
• Backend switches, architecture overrides, and runtime device selection are
  unchanged.
• No compiled library is required; update repro_sum.hpp and continue
  compiling as C++17 with SYCL.

Changes since 1.2.3

• Add a loop-free K=2 merge specialization.
• Work around the Intel IGC two-iteration merge-loop miscompilation.
• Add double- and single-precision K=2 regression tests.
• Move AMD libspirv handling into a project-local resource overlay.
• Re-emit the signed-character libspirv bitcode with the AMDGCN target triple.
• Add Intel Arc Pro B70 benchmark and reproducibility results.
• Update the encoded library version from 1.2.3 to 1.2.4.

Full comparison:
v1.2.3...v1.2.4

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.

SyclReproSum 1.2.3

SyclReproSum 1.2.3 automatically detects the installed NVIDIA GPU
architecture and passes the corresponding CUDA target to DPC++. CUDA builds
now target the local GPU instead of silently using the compiler's default
architecture.

The release remains header-only, source-compatible with version 1.2.2,
C++17, and based on the Ahrens-Demmel-Nguyen (ADN) binned floating-point
format. The binned primitives remain adaptations of the ReproBLAS reference
implementation.

Highlights

• Detect the NVIDIA compute capability with nvidia-smi when a CUDA target
  is enabled.
• Convert the reported capability to the DPC++ CUDA architecture form, such
  as 12.0 to sm_120.
• Pass the detected architecture through
  --cuda-gpu-arch=<architecture> to the NVIDIA SYCL target backend.
• Allow cross-builds and explicit selection with CUDA_ARCH, for example
  CUDA_ARCH=sm_86.
• Fail before compilation with a descriptive message when a CUDA target is
  enabled but no architecture can be detected or supplied.
• Document automatic detection, cross-compilation, and multi-GPU selection.
• Update the encoded library version to 1.2.3.

NVIDIA architecture detection

Previous releases enabled the nvptx64-nvidia-cuda SYCL target but did not
select a CUDA GPU architecture. The architecture was therefore left to the
compiler default, which may be older than the installed GPU. On the release
validation system, the compiler default was sm_75 while the installed
NVIDIA RTX PRO 4500 Blackwell reports compute capability 12.0.

Version 1.2.3 queries the first capability reported by:

nvidia-smi --query-gpu=compute_cap --format=csv,noheader

The Makefile removes the decimal point, prefixes the value with sm_, and
passes the resulting architecture to the CUDA backend. On the validation
system, the normal fat-binary build now includes:

-Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_120

The detected architecture is also displayed by make check-config before
compilation:

NVIDIA target: nvptx64-nvidia-cuda (sm_120)

This change affects build-time target selection only. It does not change
runtime device selection, which continues to use ONEAPI_DEVICE_SELECTOR.

Overrides, cross-compilation, and multi-GPU systems

Set CUDA_ARCH explicitly when cross-compiling, when nvidia-smi is not
available, or when targeting a different installed architecture:

make CUDA_ARCH=sm_86
make test CUDA_ARCH=sm_86
make test-validation CUDA_ARCH=sm_86

On a multi-GPU system, automatic detection uses the architecture reported for
the first GPU. Use CUDA_ARCH to select another architecture.

When the active SYCL_TARGETS list contains nvptx64-nvidia-cuda and neither
automatic detection nor CUDA_ARCH provides an architecture, the Makefile
stops with a descriptive error and an example override. Builds without the
CUDA target do not require nvidia-smi or CUDA_ARCH.

The architecture flag is part of the recorded build configuration, so
changing CUDA_ARCH causes the affected binaries to be rebuilt.

Reproducibility and numerical behavior

This release does not change the ADN accumulator arithmetic, bin layout,
deposit, renormalization, merge, or conversion paths. It also does not change
the public adn::sum, adn::cumsum, or adn::validate_environment APIs.

Version 1.2.3 retains the reproducibility and numerical behavior of version
1.2.2:

• Scalar sums remain independent of input permutation, work-group size,
  work-item scheduling, and validated device.
• Cumulative prefixes remain bit-identical across runs, supported work-group
  sizes, and validated devices for a fixed input sequence.
• The last cumulative output remains independent of input permutation and
  bit-identical to the corresponding scalar sum.
• Canonical NaN handling, infinity propagation, signed zero, subnormal
  support, work-group capacity validation, and accumulator-capacity limits
  are unchanged.

The CUDA architecture selection changes code generation for the selected
NVIDIA target but does not alter the library's bit-reproducibility contract.

Testing and validation

The complete suite remains:

• 174 correctness cases and four throughput benchmarks per selected device.
• Ten cross-device reproducibility cases.
• One version test.
• 723 total tests on the four-device validation system.

The CUDA, SPIR-V, and AMDGCN fat binaries were validated on:

• NVIDIA RTX PRO 4500 Blackwell via CUDA, automatically targeted as
  sm_120.
• AMD Radeon Pro W7500 (gfx1102) via HIP.
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.

Release validation results:

• The CUDA, SPIR-V, and AMDGCN example and test fat binaries built
  successfully.
• The example fat binary ran successfully.
• All 707 non-benchmark tests passed.
• The same 707 tests passed under hostile host floating-point modes.
• Both unsafe-device variants passed all four expected rejection tests.
• The compile-time unsafe evaluation-mode rejection check passed.
• All 723 tests passed with the sixteen 100-million-element scalar and
  cumulative throughput benchmarks.
• All three AMD-only benchmark invocations passed all four benchmarks.
• Automatic sm_120 detection was verified on the NVIDIA validation GPU.
• A CUDA_ARCH=sm_86 override build compiled and ran successfully on the
  NVIDIA validation GPU.
• The missing-architecture path produced the expected descriptive failure.
• clang-format 18.1.8, project-authored ASCII content, ASCII filenames,
  trailing-whitespace, and diff-whitespace checks passed.

The cross-device, shuffled-input, work-group-size, repeated-run,
special-value, cumulative-prefix, and work-group-capacity tests continue to
pass with the explicitly targeted NVIDIA device image.

Compatibility and upgrade notes

• Version 1.2.3 is source-compatible with version 1.2.2.
• Existing calls to adn::sum, adn::cumsum, and
  adn::validate_environment require no changes.
• Runtime device selection and the default CUDA, SPIR-V, and AMDGCN fat-binary
  target list are unchanged.
• Native CUDA builds now require working nvidia-smi architecture detection
  or an explicit CUDA_ARCH value.
• Cross-build and headless environments should set CUDA_ARCH explicitly.
• Multi-GPU systems should set CUDA_ARCH when the first reported GPU is not
  the intended build target.
• ENABLE_NVIDIA=0 or a SYCL_TARGETS list without the CUDA target continues
  to build without NVIDIA architecture detection.
• No compiled library or link step is required. Replace or update
  repro_sum.hpp and continue compiling as C++17 with SYCL.

Build requirements

• A C++17 compiler with SYCL support.
• Intel DPC++/LLVM with the required CUDA, SPIR-V, and/or AMDGCN targets.
• CUDA toolkit support when ENABLE_NVIDIA=1.
• nvidia-smi for automatic CUDA architecture detection, or an explicit
  CUDA_ARCH override.
• ROCm and a supported AMD GPU architecture when ENABLE_AMD=1.
• An Intel GPU or CPU runtime when ENABLE_SPIRV=1.
• At least one fp64-capable SYCL CPU or GPU device that passes runtime
  environment validation.
• Shared USM for every valid nonempty operation.
• Device USM for cumulative sums and sums of plain host arrays.
• Sufficient work-group and local-memory capacity for the selected K and
  WG_SIZE.
• oneDPL headers for the cumulative-sum throughput baseline.
• The Google Test submodule for the test suite.

Typical commands remain:

make all
make run
make test
make test-validation

For a faster correctness-only run:

./repro_test "--gtest_filter=-*Bench*"

Changes since 1.2.2

• Detect NVIDIA compute capability with nvidia-smi.
• Convert the detected capability to a DPC++ sm_XX architecture.
• Pass --cuda-gpu-arch to the NVIDIA SYCL target backend.
• Add the CUDA_ARCH override for cross-builds and explicit target
  selection.
• Reject CUDA builds with a descriptive error when no architecture is
  available.
• Document NVIDIA auto-detection and multi-GPU override behavior.
• Update the encoded library version from 1.2.2 to 1.2.3.

Full comparison:
v1.2.2...v1.2.3

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.

SyclReproSum 1.2.2

SyclReproSum 1.2.2 validates requested work-group and local-memory capacity
before submitting reduction or scan kernels. Unsupported K and WG_SIZE
combinations now fail with descriptive library exceptions instead of opaque
backend out-of-resources errors.

The release also strengthens strict floating-point enforcement inside the
runtime device environment probe. The library remains header-only,
source-compatible with version 1.2.1, C++17, and based on the
Ahrens-Demmel-Nguyen (ADN) binned floating-point format. The binned primitives
remain adaptations of the ReproBLAS reference implementation.

Highlights

• Validate WG_SIZE against the device's maximum supported work-group size.
• Validate the local-memory footprint of the per-work-item binned
  accumulators before kernel submission.
• Report requested and available capacity in descriptive
  std::runtime_error exceptions.
• Apply the strict floating-point pragma directly inside the runtime device
  environment probe kernel.
• Add parameterized correctness coverage for supported and rejected
  work-group configurations.
• Update the encoded library version to 1.2.2.

Work-group capacity validation

The reduction and cumulative-scan kernels keep one Binned<T, K>
accumulator per work-item in local memory. Although the public template
parameters already require WG_SIZE to be a power of two from 2 through
1024, an otherwise valid compile-time configuration may still exceed the
limits of a particular device.

Version 1.2.2 checks two device-specific limits for every nonempty
adn::sum and adn::cumsum operation:

• WG_SIZE must not exceed
  sycl::info::device::max_work_group_size.
• WG_SIZE * sizeof(Binned<T, K>) must not exceed
  sycl::info::device::local_mem_size.

For example, a configuration with additional folds and a large work-group:

double result = adn::sum<9, 512>(q, input, count);

may fit on one device but exceed the work-group or local-memory capacity of
another. The library now detects that condition before submitting a
reduction or scan kernel.

If the requested work-group is too large, the exception reports the
requested and maximum work-group sizes. If the binned accumulators require
too much local memory, the exception reports both the required and
available byte counts and suggests reducing WG_SIZE or K.

This replaces backend-dependent failures such as
UR_RESULT_ERROR_OUT_OF_RESOURCES with a descriptive API-level
std::runtime_error.

The reported local-memory size is an upper bound. A backend may reserve part
of that memory internally, so passing this validation cannot guarantee that
every backend will accept a configuration at the exact advertised limit.

Stronger device floating-point probe

The runtime environment probe verifies effective device behavior for
round-to-nearest arithmetic, gradual underflow, and signed-zero
preservation.

Version 1.2.2 applies the strict floating-point pragma directly inside the
probe kernel body. Propagation of a pragma from the enclosing host function
into a SYCL kernel lambda is not guaranteed, so applying it at the device
expression site makes the intended floating-point requirements explicit.

The advertised capability checks and the arithmetic values tested by the
runtime probe are unchanged.

Reproducibility and numerical behavior

This release does not change the ADN accumulator arithmetic, bin layout,
deposit, renormalization, merge, or conversion paths. It also does not change
the public adn::sum, adn::cumsum, or adn::validate_environment APIs.

Valid configurations retain the same bit-reproducibility and numerical
behavior as version 1.2.1. Within the documented accumulator capacity:

• Scalar sums remain independent of input permutation, work-group size,
  work-item scheduling, and validated device.
• Cumulative prefixes remain bit-identical across runs, supported work-group
  sizes, and validated devices for a fixed input sequence.
• The last cumulative output remains independent of input permutation and
  bit-identical to the corresponding scalar sum.
• Canonical NaN handling, infinity propagation, signed zero, subnormal
  support, and accumulator-capacity limits are unchanged.

The new validation only rejects configurations that exceed the advertised
work-group or local-memory capacity of the selected device.

Testing and validation

One parameterized correctness case was added per selected device:

• WorkGroupCapacityValidated evaluates a K=9, WG_SIZE=512
  configuration against the device's advertised limits.
• When the configuration fits, the test verifies successful validation and
  a correct scalar result.
• When it does not fit, the test verifies that validation, adn::sum, and
  adn::cumsum throw std::runtime_error before the backend can surface an
  opaque resource failure.

The complete suite now contains:

• 174 correctness cases and four throughput benchmarks per selected device.
• Ten cross-device reproducibility cases.
• One version and capacity test.
• 723 total tests on the four-device validation system.

The CUDA, SPIR-V, and AMDGCN fat binaries were validated on:

• NVIDIA RTX PRO 4500 Blackwell via CUDA.
• AMD Radeon Pro W7500 (gfx1102) via HIP.
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.

Release validation results:

• All 707 non-benchmark tests passed.
• The same 707 tests passed under hostile host floating-point modes.
• Both unsafe-device variants passed all four expected rejection tests.
• The compile-time unsafe evaluation-mode rejection check passed.
• All 723 tests passed with the sixteen 100-million-element scalar and
  cumulative throughput benchmarks.
• All three AMD-only benchmark invocations passed.
• The CUDA, SPIR-V, and AMDGCN example fat binary built and ran successfully.
• clang-format 18, ASCII-content, and whitespace checks passed.

The cross-device, shuffled-input, work-group-size, repeated-run,
special-value, and cumulative-prefix reproducibility tests continue to pass.

Compatibility and upgrade notes

• Version 1.2.2 is source-compatible with version 1.2.1.
• Existing calls to adn::sum, adn::cumsum, and
  adn::validate_environment require no changes.
• Supported K and WG_SIZE configurations retain their existing behavior.
• Configurations that exceed a device's advertised work-group or local-memory
  limits now fail earlier with std::runtime_error.
• Applications that catch backend-specific out-of-resources failures should
  also handle the new library-level exception.
• Build targets, backend switches, compiler requirements, and public template
  defaults are unchanged.
• No compiled library or link step is required. Replace or update
  repro_sum.hpp and continue compiling as C++17 with SYCL.

Build requirements

• A C++17 compiler with SYCL support.
• Intel DPC++/LLVM with the required CUDA, SPIR-V, and/or AMDGCN targets.
• CUDA toolkit support when ENABLE_NVIDIA=1.
• ROCm and a supported AMD GPU architecture when ENABLE_AMD=1.
• An Intel GPU or CPU runtime when ENABLE_SPIRV=1.
• At least one fp64-capable SYCL CPU or GPU device that passes runtime
  environment validation.
• Shared USM for every valid nonempty operation.
• Device USM for cumulative sums and sums of plain host arrays.
• Sufficient work-group and local-memory capacity for the selected K and
  WG_SIZE.
• oneDPL headers for the cumulative-sum throughput baseline.
• The Google Test submodule for the test suite.

Typical commands remain:

make all
make run
make test
make test-validation

For a faster correctness-only run:

./repro_test "--gtest_filter=-*Bench*"

Changes since 1.2.1

• Validate WG_SIZE against the device maximum work-group size.
• Validate the per-work-group binned-accumulator local-memory requirement.
• Produce descriptive exceptions with requested and available capacity
  information.
• Apply strict floating-point controls directly inside the environment probe
  kernel.
• Add parameterized work-group capacity coverage for scalar and cumulative
  sums.
• Update the encoded library version from 1.2.1 to 1.2.2.

Full comparison:
v1.2.1...v1.2.2

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.

SyclReproSum 1.2.1

SyclReproSum 1.2.1 adds validated AMD HIP support alongside the existing
CUDA and SPIR-V targets. The standard Make targets now build a single
CUDA + SPIR-V + AMDGCN fat binary by default, while backend switches let
users disable toolchains that are not installed.

The library remains header-only, source-compatible with version 1.2.0,
C++17, and based on the Ahrens-Demmel-Nguyen (ADN) binned floating-point
format. The binned primitives remain adaptations of the ReproBLAS reference
implementation.

Highlights

• Add validated AMD HIP support for the AMD Radeon Pro W7500 (gfx1102).
• Build CUDA, SPIR-V, and AMDGCN device images from the normal make,
  make all, make run, make test, and make test-validation targets.
• Add ENABLE_AMD, ENABLE_NVIDIA, and ENABLE_SPIRV switches, all enabled
  by default.
• Detect the AMD GPU architecture with rocm_agent_enumerator, with
  AMD_GPU_ARCH available as an explicit override.
• Select only runtime backends represented in the compiled fat binary.
• Add the AMD libspirv lookup compatibility alias and the verified
  GlobalOffset LTO pipeline required by the current compiler.
• Document AMD scalar-sum and cumulative-sum throughput.
• Expand release validation to three GPUs and one CPU across HIP, CUDA,
  Level Zero, and OpenCL.
• Update the encoded library version to 1.2.1.

AMD HIP support

With all default backends enabled, the offload target list is:

nvptx64-nvidia-cuda,spir64,amdgcn-amd-amdhsa

The AMD architecture is detected automatically when
rocm_agent_enumerator is available. It can also be specified directly:

make test AMD_GPU_ARCH=gfx1102

Users can disable unavailable backends without introducing separate build or
test workflows:

# CUDA + SPIR-V only
make test ENABLE_AMD=0

# SPIR-V only
make test ENABLE_AMD=0 ENABLE_NVIDIA=0

# HIP only
make test ENABLE_NVIDIA=0 ENABLE_SPIRV=0

The selected target list also controls ONEAPI_DEVICE_SELECTOR, so a disabled
backend is not enumerated by the test suite when its device image is absent.
An explicitly supplied ONEAPI_DEVICE_SELECTOR still takes precedence.

The current Intel DPC++ AMD toolchain requires two compatibility measures:

• The AMD libspirv device library must contain the group non-uniform shuffle
  builtins supplied by
  intel/llvm#22665.
• The compiler driver looks for AMD libspirv in amdgcn-amd-amdhsa, while the
  build installs it in amdgcn-amd-amdhsa-llvm. The Makefile creates the
  resource-directory alias used as a local workaround for
  intel/llvm#19339.

The Makefile also supplies the verified AMD offload linker pipeline:

--lto-newpm-passes=globaloffset,lto<O3>

These measures affect compiler and device-library integration only. They do
not patch or alter the SyclReproSum accumulator implementation.

Reproducibility and numerical behavior

This release does not change the ADN accumulator arithmetic, bin layout,
deposit, renormalization, merge, or conversion paths. It also does not change
the public adn::sum, adn::cumsum, or adn::validate_environment APIs.

The one-million-element double wide-range cross-device test produces:

0x430FC878C605717F

The same bits, including after independent input shuffles, were verified on:

• NVIDIA RTX PRO 4500 Blackwell via CUDA.
• AMD Radeon Pro W7500 (gfx1102) via HIP.
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.

The established guarantees for canonical NaNs, infinities, signed zero,
subnormal values, accumulator capacity, input order, work-group size, and
validated devices are unchanged from version 1.2.0.

Performance

The AMD benchmarks process 100 million elements. Each value is the mean of
three benchmark invocations, and each invocation averages ten timed runs
after one warmup. Throughput is measured in billions of elements per second.

Scalar-sum throughput

Device	Type	adn::sum	Plain sycl::reduction	Slowdown
AMD Radeon Pro W7500 via HIP	double	9.2 GElements/s	14.8 GElements/s	1.6x
AMD Radeon Pro W7500 via HIP	float	15.7 GElements/s	20.0 GElements/s	1.3x

Cumulative-sum throughput

Device	Type	adn::cumsum	oneDPL inclusive_scan	Slowdown
AMD Radeon Pro W7500 via HIP	double	0.751 GElements/s	4.876 GElements/s	6.5x
AMD Radeon Pro W7500 via HIP	float	1.089 GElements/s	9.126 GElements/s	8.4x

The baselines are not reproducible and are provided only as throughput
comparison points.

Testing and validation

The release fat binary was validated on:

• NVIDIA RTX PRO 4500 Blackwell via CUDA.
• AMD Radeon Pro W7500 (gfx1102) via HIP.
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.

The complete suite contains:

• 173 correctness cases and four throughput benchmarks per selected device.
• Ten cross-device reproducibility cases.
• One version test.
• 719 total tests on the four-device validation system.

Release validation results:

• All 703 non-benchmark tests passed.
• The same 703 tests passed under hostile host floating-point modes.
• Both unsafe-device variants passed all four expected rejection tests.
• The compile-time unsafe evaluation-mode rejection check passed.
• All 719 tests passed with the sixteen 100-million-element benchmarks.
• All three AMD-only benchmark invocations passed.
• clang-format 18, ASCII-content, and whitespace checks passed.

The cross-device, shuffled-input, work-group-size, repeated-run,
special-value, and cumulative-prefix reproducibility tests continue to pass.

Compatibility and upgrade notes

• Version 1.2.1 is source-compatible with version 1.2.0.
• Existing calls to adn::sum, adn::cumsum, and
  adn::validate_environment require no changes.
• No compiled library or link step is required. Replace or update
  repro_sum.hpp and continue compiling as C++17 with SYCL.
• CUDA, SPIR-V, and AMD support are now enabled by default in the Makefile.
  Disable any backend whose toolkit is unavailable with its corresponding
  ENABLE_* switch.
• AMD builds require ROCm, an AMD-capable DPC++ compiler, the patched libspirv
  shuffle builtins, and an fp64-capable device that passes environment
  validation.
• Switching backend options automatically updates the build configuration and
  rebuilds stale binaries; a manual make clean is not required.

Build requirements

• A C++17 compiler with SYCL support.
• Intel DPC++/LLVM with the required CUDA, SPIR-V, and/or AMDGCN targets.
• CUDA toolkit support when ENABLE_NVIDIA=1.
• ROCm and a supported AMD GPU architecture when ENABLE_AMD=1.
• An Intel GPU or CPU runtime when ENABLE_SPIRV=1.
• At least one fp64-capable SYCL CPU or GPU device that passes runtime
  environment validation.
• Shared USM for every valid nonempty operation.
• Device USM for cumulative sums and sums of plain host arrays.
• oneDPL headers for the cumulative-sum throughput baseline.
• The Google Test submodule for the test suite.

Typical commands remain:

make all
make run
make test
make test-validation

For a faster correctness-only run:

./repro_test "--gtest_filter=-*Bench*"

Changes since 1.2.0

• Add AMDGCN compilation and AMD HIP runtime selection.
• Add configurable CUDA, SPIR-V, and AMD backend switches.
• Add AMD architecture detection and explicit override support.
• Add AMD libspirv and offload linker compatibility handling.
• Validate reproducibility and hostile-mode behavior on the AMD Radeon Pro
  W7500.
• Document AMD scalar and cumulative throughput.
• Expand the default validation system from three devices to four.
• Update the encoded library version from 1.2.0 to 1.2.1.

Full comparison:
v1.2.0...v1.2.1

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.

SyclReproSum 1.2.0

SyclReproSum 1.2.0 adds explicit SYCL event dependencies to the reproducible
sum and cumulative-sum APIs, clarifies the synchronous execution contract, and
strengthens input and USM capability validation. The release remains
header-only, source-compatible with existing calls, C++17, and based on the
Ahrens-Demmel-Nguyen (ADN) binned floating-point format, with binned primitives
adapted from the ReproBLAS reference implementation.

Highlights

• New dependency-aware overloads for both adn::sum and adn::cumsum.
• Producer events can be passed directly to the library without a separate
  host-side wait.
• Dependencies are attached to the first library operation that accesses user
  memory, including plain-host-pointer copy paths.
• Existing overloads remain available and synchronous, with their caller-side
  USM synchronization responsibility now documented explicitly.
• adn::sum now rejects null pointers for nonempty inputs with
  std::invalid_argument.
• adn::sum now rejects element counts whose byte size would overflow
  size_t with std::length_error.
• Shared and device USM requirements are now checked explicitly and documented
  instead of being left to allocation failure behavior.
• Missing USM capabilities produce descriptive std::runtime_error
  exceptions.
• Cross-queue producer tests verify dependency ordering for scalar and
  cumulative sums.
• The test suite now contains 173 correctness cases and four 100-million-
  element benchmarks per selected device, plus ten cross-device cases and one
  version test.
• Version macros now report 1.2.0.

Dependency-aware APIs

The new overloads accept a list of prerequisite SYCL events:

template <int K = 3, int WG_SIZE = 256, typename T>
T sum(sycl::queue &q, const T *arr, size_t N,
   const std::vector<sycl::event> &dependencies);

template <int K = 3, int WG_SIZE = 256, typename T>
void cumsum(sycl::queue &q, const T *input, T *output, size_t N,
   const std::vector<sycl::event> &dependencies);

For example, an input transfer can feed directly into a reproducible sum:

double *device_input = sycl::malloc_device<double>(count, q);
sycl::event input_ready =
   q.memcpy(device_input, host_input, count * sizeof(double));

double result = adn::sum(q, device_input, count, {input_ready});

The caller does not need to wait for input_ready separately. The library
orders its first input access after that event and waits for its own submitted
work before returning.

The same model applies to cumulative sums:

sycl::event input_ready = producer_queue.submit([&](sycl::handler &h) {
   h.parallel_for(sycl::range<1>(count), [=](sycl::id<1> i) {
      input[i] = static_cast<double>(i[0] + 1);
   });
});

adn::cumsum(q, input, output, count, {input_ready});

The producer and consumer queues may be distinct when the supplied event is a
valid dependency for the target context.

Ordering and synchronous behavior

Both the original and dependency-aware overloads are synchronous. A return
from adn::sum or adn::cumsum means all work submitted by that library call
has completed. These APIs intentionally do not return a SYCL event.

The dependency list controls when the library may begin accessing the user
ranges:

Operation and pointer path	First ordered library access
sum with USM input	First reduction kernel
sum with plain host input	Host-to-device copy
cumsum with USM input and output	Tile-total kernel
cumsum with any plain host range	Required copy or the first dependent kernel, with later accesses ordered transitively

For overloads without a dependency list, callers must complete earlier
operations that write or conflict with USM input or output ranges before the
call. Callers must also prevent conflicting concurrent access while the
synchronous operation is running.

This makes the ownership boundary explicit: the dependency overloads express
known producer ordering to the library, while the original overloads preserve
their existing simple synchronous interface.

Stronger input validation

adn::sum now matches the defensive behavior already provided by
adn::cumsum for invalid nonempty input.

Null pointers

For N > 0, a null sum input is rejected deterministically:

adn::sum(q, static_cast<const double *>(nullptr), 1);
// Throws std::invalid_argument.

For N == 0, the pointer is not inspected and the scalar sum still returns
positive zero. Empty cumulative sums likewise return without touching either
range.

Byte-size overflow

adn::sum now verifies that N * sizeof(T) can be represented by size_t.
An overflowing count throws std::length_error before pointer classification,
allocation, copying, or kernel submission. The existing conservative
adn::max_reproducible_count<T> check remains in place and is evaluated first.

These checks provide stable API-level errors instead of backend-dependent
copy, allocation, or kernel failures.

Explicit USM requirements

Version 1.2.0 makes the implementation's USM requirements part of the
documented and validated device contract:

• Every valid, nonempty adn::sum and adn::cumsum operation requires shared
  USM support for environment probes. Scalar sums also store their partial
  binned accumulators in shared USM.
• adn::cumsum additionally requires device USM for tile totals, tile
  prefixes, and other internal scan storage.
• adn::sum requires device USM only when a plain host input must be copied to
  a temporary device allocation.
• A direct sum from device, shared, or host USM does not add a device-USM
  requirement beyond the capabilities needed by that pointer itself.

adn::validate_environment<T>(q) now validates shared USM in addition to the
existing floating-point requirements. Internal shared and device allocation
helpers also check their corresponding SYCL aspects before attempting an
allocation. Missing support produces a descriptive std::runtime_error
instead of an incidental std::bad_alloc or backend allocation exception.

The established floating-point requirements remain unchanged: supported
devices must provide fp64, round-to-nearest behavior, gradual underflow,
Inf/NaN handling, and signed-zero preservation as verified by the advertised
capabilities and runtime device probes.

Reproducibility and numerical behavior

This release does not change the ADN accumulator arithmetic, bin layout,
deposit, renormalization, merge, or conversion paths. Valid inputs retain the
same bit-reproducibility guarantees and numerical behavior as version 1.1.0.

Within adn::max_reproducible_count<T>, scalar sums remain independent of
input permutation, work-group size, work-item scheduling, and validated
device. Cumulative prefixes remain bit-identical across runs, supported
work-group sizes, and validated devices for a fixed input sequence, with the
last prefix retaining full input-order independence.

Canonical NaN handling, infinity propagation, signed zero, subnormal support,
capacity limits, fold-count behavior, and work-group-size configuration are
unchanged.

Testing and validation

Four parameterized correctness cases were added per selected device:

• Sum_USMDependencies holds an unfinished cross-queue producer and verifies
  that adn::sum cannot access the shared-USM input or return early.
• Cumsum_USMDependencies performs the same ordering check for both cumulative
  input and output, then validates every generated prefix.
• Sum_ByteSizeOverflowRejected verifies the new deterministic
  std::length_error path.
• USMCapabilitiesValidated covers shared and device USM aspect checks.

The empty-input sum test now also covers a null pointer for both zero and
nonzero counts, and environment validation tests account for shared USM.

The complete suite now contains:

• 173 correctness cases and four throughput benchmarks per selected device.
• Ten cross-device reproducibility cases.
• One version and capacity test.
• 542 total tests on the validation system with two GPUs and one CPU.

The release was built as one CUDA plus SPIR-V fat binary and validated on:

• NVIDIA RTX PRO 4500 Blackwell via CUDA.
• Intel Graphics [0x7d67] via Level Zero.
• Intel Core Ultra 7 265 CPU via OpenCL.

Release validation results:

• All 530 non-benchmark tests passed.
• The same 530 tests passed under hostile host floating-point modes.
• Both unsafe-device variants passed all three expected rejection tests.
• The compile-time unsafe evaluation-mode rejection check passed.
• All 542 tests passed when the twelve 100-million-element scalar and
  cumulative throughput benchmarks were included.
• clang-format 18, ASCII-content, and whitespace checks passed.

The dependency and validation changes are outside the binned arithmetic, and
the existing shuffled-input, work-group matrix, repeated-run, special-value,
and cross-device bit-identity tests continue to pass.

Compatibility and upgrade notes

• Version 1.2.0 is an additive, source-compatible release for valid existing
  calls.
• Existing three-argument adn::sum and four-argument adn::cumsum calls do
  not require source changes.
• Existing overloads remain synchronous and now document that callers are
  responsible for completing conflicting prior USM operations.
• Applications with producer events can adopt the new overloads incrementally.
• A nonempty null adn::sum input now throws std::invalid_argument instead
  of reaching backend-dependent behavior.
• A sum count with an unrepresentable byte size now throws
  std::length_error.
• Devices without the required USM aspects now fail early with
  std::runtime_error and a descriptive requirement.
• No compiled library or link step is required. Replace or update
  repro_sum.hpp and continue compiling as C++17 with SYCL.

Build requirements

• A C++17 compiler with SYCL support.
• Intel DPC++/LLVM with the required offload target support.
• An fp64-capable SY...

SyclReproSum 1.1.0

SyclReproSum 1.1.0 adds bit-reproducible cumulative sums and strengthens the
floating-point, capacity, device-reduction, and resource-safety guarantees of
the scalar summation API. The library remains header-only, C++17, and based on
the Ahrens-Demmel-Nguyen (ADN) binned floating-point format, with binned
primitives adapted from the ReproBLAS reference implementation.

Highlights

• New adn::cumsum API for reproducible inclusive cumulative sums of float
  and double values on SYCL CPU and GPU devices.
• Bit-identical cumulative results across runs, thread schedules, supported
  work-group sizes, and validated devices for a fixed input sequence.
• Exact in-place cumulative sums with input == output.
• Independent support for plain host pointers and device, shared, or host USM
  pointers for both cumulative-sum input and output.
• A hierarchical binned scan that stores binned state only for tile totals and
  tile prefixes, rather than one full accumulator per output element.
• Stronger fail-closed environment validation for both adn::sum and
  adn::cumsum, including effective device arithmetic probes.
• Explicit accumulator-capacity checks through the public
  adn::max_reproducible_count<T> constant.
• Final scalar reduction and conversion now remain on the device with explicit
  event ordering, making the result independent of hostile host floating-point
  modes.
• Exception-safe ownership and cleanup for all internal USM allocations.
• Version macros now report 1.1.0.

New reproducible cumulative-sum API

adn::cumsum computes an inclusive reproducible prefix sum:

#include <sycl/sycl.hpp>
#include <vector>
#include "repro_sum.hpp"

sycl::queue q{sycl::gpu_selector_v};

std::vector<double> input = {1.0, 2.0, 3.0, 4.0};
std::vector<double> output(input.size());

adn::cumsum(q, input.data(), output.data(), input.size());
// output is {1.0, 3.0, 6.0, 10.0}

The template parameters match adn::sum:

adn::cumsum<6>(q, input, output, count);       // More folds
adn::cumsum<3, 512>(q, input, output, count); // Custom work-group size

For every valid index i, output[i] is the ADN reproducible sum of
input[0] through input[i]. The call is synchronous: when it returns, a
plain host output buffer has been copied back and is ready to use.

Exact in-place operation is supported:

adn::cumsum(q, data, data, count);

Partially overlapping ranges other than exact input == output are not
supported.

Cumulative-sum reproducibility model

Input order defines which values belong to each prefix. For a fixed input
sequence, every prefix result is bit-identical across runs, supported
work-group sizes, thread schedules, and validated CPU or GPU devices.

Shuffling an input generally changes the intermediate prefixes because it
changes prefix membership. The last cumulative output is special: it contains
all input values and retains the full order-independent guarantee of
adn::sum. In other words:

• Every output is reproducible for a fixed sequence.
• The last output is also independent of input permutation.
• output[count - 1] is bit-identical to adn::sum(q, input, count).

NaNs and infinities follow the same deterministic rules as scalar sums. Any
NaN, or a prefix containing both positive and negative infinity, produces the
fixed positive canonical quiet NaN for that and all following prefixes:

• float: 0x7FC00000
• double: 0x7FF8000000000000

Subnormal values, signed zero, overflow, and gradual underflow are handled by
the same binned primitives used by adn::sum.

How the cumulative scan works

adn::cumsum is a hierarchical scan that keeps intermediate totals in the ADN
binned format:

1. Contiguous input tiles are reduced to reproducible binned totals.
2. Tile totals undergo a recursive exclusive scan while remaining in binned
   form.
3. A final pass replays each tile from its incoming binned accumulator and
   converts every inclusive prefix to float or double.

Only tile totals and tile prefixes require temporary global storage. The
implementation does not materialize a complete K-fold binned accumulator for
every output element.

Pointer handling

The cumulative-sum input and output pointer types are detected independently
at runtime with sycl::get_pointer_type:

Pointer source	Input behavior	Output behavior
Device USM	Used directly	Used directly
Shared USM	Used directly	Used directly
Host USM	Used directly	Used directly
Plain host pointer	Copied to temporary device storage	Computed in temporary device storage and copied back

This supports combinations such as a plain host input with device output or a
device input with plain host output. "Used directly" means no library-side
copy; the SYCL runtime may still migrate or otherwise manage USM storage.

Stronger scalar-sum guarantees

This release also hardens adn::sum without changing its public call syntax.

Device-only final reduction

Per-work-group accumulators are now merged and converted by a final device
work-group with explicit SYCL event dependencies. Host-side floating-point
rounding, denormal, and evaluation modes therefore cannot affect the final
result.

Explicit capacity enforcement

The exact binned carry representation has a finite conservative capacity
inherited from ReproBLAS:

Type	Maximum reproducible input count
float	4,294,966,784
double	9,223,372,036,854,773,760

The limits are available as adn::max_reproducible_count<float> and
adn::max_reproducible_count<double>. Both adn::sum and adn::cumsum reject
larger inputs with std::length_error before pointer inspection or allocation.
adn::cumsum also rejects a byte-size overflow and requires non-null pointers
when the count is nonzero.

Splitting a larger input into scalar chunks and adding the rounded chunk
results does not preserve the library's reproducibility guarantee.

Strict floating-point requirements

The header continues to reject fast-math and equivalent unsafe compilation
modes. Version 1.1.0 additionally:

• Verifies the required IEEE 754 binary32 and binary64 layouts at compile time.
• Requires device expressions to evaluate in their declared type.
• Explicitly disables reassociation, reciprocal transformations, and FP
  contraction in accumulator arithmetic where supported.
• Checks the device's advertised round-to-nearest, denormal, Inf/NaN, and fp64
  capabilities.
• Executes runtime-loaded device probes for effective rounding behavior,
  gradual underflow, and signed-zero preservation.

Both float and double operations require device fp64 support because final
float conversion accumulates in double precision.

Successful validation is cached per backend, device, and floating-point type
and shared safely across host threads. Root-enumerated devices use the fixed
cache; dynamically created sub-devices are validated on each use. Applications
can preflight explicitly with:

adn::validate_environment<float>(q);
adn::validate_environment<double>(q);

Validation throws std::runtime_error when the library cannot establish the
required arithmetic guarantees.

Exception-safe USM cleanup

Internal shared and device USM allocations now use context-aware RAII owners.
If allocation, submission, copying, or synchronization throws, owned storage
is released without throwing from cleanup. When a later queue submission
fails after earlier work was submitted, cleanup waits for pending work before
releasing the associated USM allocation.

Accuracy and configuration

The default remains K=3 folds and a work-group size of 256. Supported
work-group sizes are powers of two from 2 through 1024, subject to the target
device's work-group and local-memory limits. Increasing K trades performance
and storage for additional low-order coverage.

Type	Bin width	Maximum K	Dominant absolute-error scale at K=3
double	40 bits	52	approximately N * 2^-80 * max abs input
float	13 bits	21	approximately N * 2^-26 * max abs input

The error scale is absolute rather than relative. Severe cancellation can make
the error large relative to the exact final sum. The result is reproducible,
but it is not an arbitrary-precision exact sum.

Performance

The included benchmarks process 100 million elements. Each value is the mean
of three benchmark invocations, and each invocation averages ten timed runs
after one warmup. Throughput is measured in billions of elements per second.

Cumulative-sum throughput

The baseline is oneDPL inclusive_scan using a standard device policy and
std::plus on the same device-USM input and output. The baseline is not
reproducible and is provided only as a throughput comparison.

Device	Type	adn::cumsum	oneDPL baseline	Slowdown
NVIDIA GB10 via CUDA	double	2.364 GElements/s	7.606 GElements/s	3.2x
NVIDIA GB10 via CUDA	float	3.413 GElements/s	14.980 GElements/s	4.4x
NVIDIA RTX PRO 4500 Blackwell via CUDA	double	4.855 GElements/s	21.383 GElements/s	4.4x
NVIDIA RTX PRO 4500 Blackwell via CUDA	float	6.696 GElements/s	40.480 GElements/s	6.0x
Intel Graphics [0x7d67] via Level Zero	double	0.167 GElements/s	0.610 GElements/s	3.7x
Intel Graphics [0x7d67] via Level Zero	float	0.331 GElements/s	2.445 GElements/s	7.4x
Intel Core Ultra 7 265 CPU via OpenCL	double	0.118 GElements/s	1.490 GElements/s	12.6x
Intel Core Ultra 7 265 CPU via OpenCL	float	0.138 GElements/s	2.536 GElements/s	18.4x

Scalar-sum throughput

For completeness, scalar-sum results remain:

Device	Type	adn::sum	Plain sycl::reduction	Slowdown
NVIDIA GB10 via CUDA	double	26.3 GElements/s	31.7 GElements/s	1.2x
NVIDIA GB10 via CUDA	float	55.5 ...