SYprox

A SYCL header-only library for approximate computing on heterogeneous architectures.

SYprox extends the SYCL buffer and accessor classes with approximate computing techniques — data perforation, signal reconstruction, and mixed precision — and lets them be composed to explore a wide space of performance/accuracy trade-offs. Because it is a plain header-only library, it needs no custom compiler, no language extension, and no source-to-source tooling: it works with a standard SYCL compiler and is portable across AMD, Intel, and NVIDIA GPUs.

SYprox is the artifact of the ICS '25 paper SYprox: Combining Host and Device Perforation with Mixed Precision Approximation on Heterogeneous Architectures (see Citation).

Key ideas

Device perforation. The classic approach: all data are transferred to the device, and the kernel accesses them in a perforated shape. PerforatedAccessor overloads operator[] so that an access to a[i] is translated to a[i * SKIP_RATE], leaving the original data untouched in memory.

Host perforation. Introduced by SYprox. Data are perforated on the host, before the transfer, so only the perforated subset ever crosses the PCIe link. PerforatedBuffer intercepts the sycl::buffer constructor and applies the perforation scheme to the host pointer and to the range. This has two advantages over device perforation: less data movement, and better cache behaviour, since the surviving elements are contiguous rather than strided.

Reconstruction. Perforated elements can be filled back in, either before the computation (input reconstruction, using local memory) or after it (output reconstruction). Two strategies are provided: nearest neighbour (nn) and linear interpolation (lerp). Input reconstruction is generally more accurate; output reconstruction is generally faster.

Mixed precision. Since PerforatedBuffer controls buffer construction, host data can be narrowed to a lower-precision type (sycl::half, bfloat16, …) on the way to the device and combined freely with any of the above.

Requirements

• A SYCL 2020 compiler — either Intel DPC++/oneAPI or hipSYCL/AdaptiveCpp
• CMake ≥ 3.13
• A C++17 compiler

Building

SYprox is header-only, so there is nothing to compile for the library itself — it is exported as a CMake INTERFACE target. The build system exists to configure the SYCL backend and to build the samples.

Two cache variables control the configuration:

Variable	Values	Meaning
AX_SYCL_BACKEND	dpcpp, hipSYCL	SYCL implementation to use (required)
AX_HW_SUPPORT	nvidia	Target architecture; adds -fsycl-targets=nvptx64-nvidia-cuda under DPC++
AX_CUDA_ARCH	e.g. sm_70	CUDA architecture, when AX_HW_SUPPORT=nvidia
AX_BUILD_SAMPLES	ON/OFF (default OFF)	Build the programs under samples/

Intel DPC++

cmake -B build -DAX_SYCL_BACKEND=dpcpp -DCMAKE_CXX_COMPILER=clang++ -DAX_BUILD_SAMPLES=ON
cmake --build build

DPC++ targeting an NVIDIA GPU

cmake -B build \
  -DAX_SYCL_BACKEND=dpcpp \
  -DCMAKE_CXX_COMPILER=clang++ \
  -DAX_HW_SUPPORT=nvidia \
  -DAX_CUDA_ARCH=sm_70 \
  -DAX_BUILD_SAMPLES=ON
cmake --build build

hipSYCL / AdaptiveCpp

cmake -B build -DAX_SYCL_BACKEND=hipSYCL -DAX_BUILD_SAMPLES=ON
cmake --build build

CMAKE_BUILD_TYPE defaults to Release if not set.

Using SYprox in your own project

Add the repository as a subdirectory and link against the SYprox interface target:

add_subdirectory(SYprox)
target_link_libraries(my_app PRIVATE SYprox)

Or simply put SYprox/include on your include path — the library has no dependencies beyond SYCL itself.

Usage

Everything lives in the syprox namespace and is pulled in by a single header:

#include "syprox/syprox.hpp"

namespace ax = syprox;

Host perforation

PerforatedBuffer<T, Dim, Schema> takes the same type and dimensionality as sycl::buffer, plus a perforation schema. The host data are perforated in place and the resulting buffer has a reduced range, which you can query with get_range():

constexpr int SKIP_RATE = 2;

std::vector<int> vec_a(N), vec_b(N);
// ... fill inputs ...

// After construction the underlying sycl::buffer has range N / SKIP_RATE
ax::PerforatedBuffer<int, 1, ax::PerforationSchema_1D<int, SKIP_RATE>>
    buff_a{vec_a.data(), sycl::range<1>{N}};
ax::PerforatedBuffer<int, 1, ax::PerforationSchema_1D<int, SKIP_RATE>>
    buff_b{vec_b.data(), sycl::range<1>{N}};

sycl::buffer<int, 1> buff_c{vec_c.data(), sycl::range<1>{N}};

q.submit([&](sycl::handler &cgh) {
    auto acc_a = buff_a.get_access<sycl::access_mode::read>(cgh);
    auto acc_b = buff_b.get_access<sycl::access_mode::read>(cgh);
    auto acc_c = buff_c.get_access<sycl::access_mode::read_write>(cgh);

    // Launch over the perforated range, not the original one
    cgh.parallel_for(sycl::range<1>{buff_a.get_range()},
                     vec_add(acc_a, acc_b, acc_c));
}).wait();

Note: PerforatedBuffer perforates the host array in place. The first range / skip_rate elements of the vector you pass in are overwritten with the surviving elements. Keep a copy if you need the original data afterwards.

Device perforation

PerforatedAccessor keeps the full data on the device and strides the accesses instead:

q.submit([&](sycl::handler &cgh) {
    ax::PerforatedAccessor<float, 2, SKIP_RATE> in_acc{in_buf, cgh};
    ax::PerforatedAccessor<float, 2, SKIP_RATE> out_acc{out_buf, cgh};

    cgh.parallel_for(sycl::range<2>{N / SKIP_RATE, N}, [=](sycl::id<2> id) {
        // in_acc[i][j] resolves to the underlying element [i * SKIP_RATE][j]
        out_acc[id[0]][id[1]] = in_acc[id[0]][id[1]] * 2.0f;
    });
}).wait();

The template parameters are <DataT, Dimensions, SKIP_RATE, AccessMode, AccessTarget, IsPlaceholder>; the last three default to the same values as sycl::accessor. Subscripting is supported for 1D, 2D, and 3D accessors.

Reconstruction

Reconstruction is exposed as static member functions on the schema classes, called from inside the kernel. Output reconstruction takes the value the work-item computed and writes both it and the reconstructed neighbours:

void operator()(sycl::id<1> id) const {
    int sum = acc_a[id] + acc_b[id];
    ax::PerforationSchema_1D<int, SKIP_RATE>::nn_out_reconstruction(id, acc_c, sum);
}

Input reconstruction runs before the computation and works through local memory, so it needs an nd_item and a sycl::local_accessor:

void operator()(sycl::nd_item<2> it) const {
    value_type val_a = acc_a[it.get_global_id()];
    value_type val_b = acc_b[it.get_global_id()];

    ax::RowSchema<value_type>::local_memory_load_input<0>(it, local_acc_a, val_a);
    ax::RowSchema<value_type>::local_memory_load_input<0>(it, local_acc_b, val_b);

    ax::RowSchema<value_type>::nn_input_reconstruction(it, local_acc_a, val_a);
    ax::RowSchema<value_type>::nn_input_reconstruction(it, local_acc_b, val_b);

    // ... compute on the reconstructed local data ...
}

Available reconstruction entry points on each schema:

Function	Kind	Strategy
nn_out_reconstruction	output	nearest neighbour
linear_interpolation_output_reconstruction	output	linear interpolation
nn_input_reconstruction	input	nearest neighbour
linear_interpolation_input_reconstruction	input	linear interpolation
local_memory_load_input / local_memory_load_output	—	helper that stages data in local memory

The input helpers take a SKIP_SYNCH template parameter controlling barrier placement.

Perforation schemes

Class	Dimensions	Behaviour
PerforationSchema_1D<T, skip_rate>	1D	strided — keeps every skip_rate-th element
RowSchema<T, skip_rate>	2D	keeps every skip_rate-th row
ColSchema<T, skip_rate>	2D	keeps every skip_rate-th column

All three derive from ApproxSchema<T, Dim>, which defines the interface a scheme must implement:

template <typename T, int dim>
class ApproxSchema {
public:
    ApproxSchema(T *host_data, sycl::range<dim> range_data);
    virtual T *perforate_data();                 // rewrite the host array
    virtual sycl::range<dim> perforate_range();  // shrink the range accordingly
};

To write your own scheme, derive from ApproxSchema and implement perforate_data() and perforate_range() (plus whichever reconstruction statics your kernels call). Any class satisfying that interface can be passed as the Schema parameter of PerforatedBuffer.

Mixed precision

Declare the buffer with a narrower type than the host vector; the conversion happens during buffer construction:

std::vector<float> in(N * N);
sycl::buffer<sycl::half, 2> in_buf{/* ... */};   // half on the device
sycl::buffer<float, 2>      out_buf{out.data(), sycl::range<2>{N, N}};

Any precision defined by the SYCL standard is supported, as are vendor extensions such as bfloat16 under DPC++.

Samples

Build with -DAX_BUILD_SAMPLES=ON. Each sample runs the accurate version alongside several approximated variants so the outputs can be compared directly.

Sample	What it shows
samples/vector_add	1D strided perforation; nn/lerp, input and output reconstruction
samples/matrix_add	2D row and column schemes with all four reconstruction variants
samples/matrix_mul	Perforation applied to a tiled matrix multiplication

./build/samples/vector_add
./build/samples/matrix_add
./build/samples/matrix_mul

Repository layout

SYprox/
├── SYprox/
│   ├── CMakeLists.txt              # INTERFACE target
│   └── include/syprox/
│       ├── syprox.hpp              # umbrella header
│       ├── perforated_buffer.hpp   # PerforatedBuffer — host perforation
│       ├── perforated_accessor.hpp # PerforatedAccessor — device perforation
│       └── perforation_schema.hpp  # ApproxSchema and the built-in schemes
├── samples/
│   ├── vector_add/
│   ├── matrix_add/
│   └── matrix_mul/
└── CMakeLists.txt

Naming: paper vs. code

The ICS '25 paper presents the interface with short names for readability. The mapping to the identifiers in this repository is:

Paper	Code
pbuffer	syprox::PerforatedBuffer
paccessor	syprox::PerforatedAccessor
prow	syprox::RowSchema
pcol	syprox::ColSchema
strided scheme	syprox::PerforationSchema_1D
prange	not implemented — pass buffer.get_range() to parallel_for

Results

On eight benchmarks (median, sobel, blur, tv, gaussian, hotspot, lavaMD, leukocytes) over 100 input datasets, evaluated on AMD MI100, Intel Max 1100, and NVIDIA V100S:

• Host perforation reaches 1.3×–2× data-transfer speedup for float and 2.5×–4.2× for half precision; device perforation gives none, since all data are still transferred.
• Kernel speedup is 1.7×–4.2× for host perforation against 0.5×–3.5× for device perforation. The gap comes from cache behaviour — an L2 hit rate of 76% versus 49% under the column scheme, where device perforation loads lines it never uses.
• Combining host and device perforation with mixed precision reaches 3.5× speedup at 7% MAPE.
• The SYprox Pareto front dominates those of HPAC and Maier et al. on all benchmarks but tv, and error is consistent across all three vendors, confirming that the approximations are data-dependent rather than hardware-dependent.

See the paper for the full evaluation.

Citation

@inproceedings{carpentieri2025syprox,
  author    = {Carpentieri, Lorenzo and Cosenza, Biagio},
  title     = {{SYprox}: Combining Host and Device Perforation with Mixed Precision
               Approximation on Heterogeneous Architectures},
  booktitle = {Proceedings of the 2025 International Conference on Supercomputing
               (ICS '25)},
  year      = {2025},
  address   = {Salt Lake City, UT, USA},
  publisher = {ACM},
  doi       = {10.1145/3721145.3725741}
}

Paper: https://doi.org/10.1145/3721145.3725741 · preprint

Acknowledgments

This work was supported by the National Recovery and Resilience Plan (NRRP), call for tender No. 104 published on 02/02/2022 by the Italian Ministry of University and Research (MUR), funded by the European Union – NextGenerationEU, Mission 4, Component 1, CUP D53D23008590001, project LibreRT.

We thank CINECA for providing access to the Intel Max 1100 GPU.