SyclBenchmark is a small C++17 program that measures SYCL device throughput without calling cuBLAS, CUTLASS, oneMKL, or another compute library.

• FP32 and, when supported, FP64 fused multiply-add throughput
• FP16, TF32, INT8, and FP64 joint-matrix throughput
• Device-memory read and write bandwidth
• Host-to-device, device-to-host, and concurrent USM transfer bandwidth

Matrix tests report both complete GEMM throughput and an instruction issue-rate ceiling. The issue rate is not a complete GEMM result. Unsupported tests are shown as unavailable.

• A C++17 SYCL compiler, runtime, and device with device USM support
• Host USM support for the transfer tests
• sycl::ext::oneapi::experimental::matrix, a 32-lane subgroup, and a matching matrix_combinations entry for matrix tests
• A CUDA toolkit/runtime setup and compatible NVIDIA driver for CUDA builds

This project is currently tested with lilohuang/llvm, a fork of intel/llvm. The fork includes custom patches that are used by some benchmarks in this project and have not yet been merged into intel/llvm. Until those patches are upstreamed, builds or benchmark availability may differ when using the upstream compiler.

The default compiler is:

~/sycl_workspace/llvm/build/bin/clang++

The default build contains CUDA NVPTX and generic SPIR-V images, targeting sm_120 for CUDA.

make

Optional overrides:

make -B SYCL_ROOT=/path/to/llvm/build
make -B OFFLOAD_TARGETS=nvptx64-nvidia-cuda CUDA_ARCH=sm_86
make -B OFFLOAD_TARGETS=nvptx64-nvidia-cuda CUDA_ARCH=sm_120
make -B OFFLOAD_TARGETS=spir64

An architecture-specific CUDA build must be run on a compatible NVIDIA device.

./sycl-bench --list
./sycl-bench 2
./sycl-bench 0 2
./sycl-bench
./sycl-bench --help

With no device ID, every visible device is benchmarked.

The program warms up each test and reports the best of multiple runs. Kernel event profiling is used when available, with host timing as a fallback. FMA is counted as two operations. Matrix initialization and result checks are outside the timed region, and several output positions are validated.

concurrent-both is the aggregate bandwidth of both directions. reported.clock comes from SYCL's max_clock_frequency property. It is static metadata and may not reflect the current boost clock. Results vary with compiler, backend, clocks, power state, and memory path, so a boosted result can exceed a nominal specification.

Measured on 2026-07-24 with DPC++ 7.1.0 pre-release (Clang 23, a57498b), CUDA_ARCH=sm_120, NVIDIA driver 575.57.08, and CUDA 12.9:

SYCL throughput report
  device.id       : 2
  device.name     : NVIDIA RTX PRO 4500 Blackwell
  device.vendor   : NVIDIA Corporation
  device.driver   : CUDA 12.9
  execution.units : 82 units
  reported.clock  : 2407 MHz (static, not live)
  memory.capacity : 32117 MiB global, 99 KiB local
  sample.bytes    : 268435456

Arithmetic throughput
  FP64          0.779 TFLOP/s
  FP32         55.017 TFLOP/s

Matrix throughput
  FP16->FP32        67.098 TFLOP/s  (2048 x 2048 x 2048)
    issue-rate     235.098 TFLOP/s
  TF32->FP32        35.802 TFLOP/s  (2048 x 2048 x 2048)
    issue-rate      58.830 TFLOP/s
  INT8->INT32       74.425 TIOP/s  (2048 x 2048 x 2048)
    issue-rate     235.225 TIOP/s
  FP64->FP64         0.780 TFLOP/s  (2048 x 2048 x 2048)
    issue-rate       0.812 TFLOP/s

Device-memory throughput
  read                   691.701 GB/s
  write                  797.609 GB/s

USM transfer throughput
  host-to-device         41.180 GB/s
  device-to-host         24.276 GB/s
  concurrent-both        35.626 GB/s

Released under the MIT License.

Copyright (c) 2026 Lilo Huang.