I built and benchmarked Intel oneAPI 2026.1 DPC++ GEMM kernels against NumPy/OpenBLAS on Intel Iris Xe (Gen11 Xe-LP). The results are counterintuitive: the SYCL kernel itself is 1.5–1.9x faster than a naive kernel, but the end-to-end performance is 64x slower than NumPy due to a subprocess initialization overhead I didn't expect.
Here's everything: the code, the methodology, the real numbers, and the trap I hit.
Test Environment
| Component | Specification |
|---|---|
| CPU | Intel Core i7-1195G7 @ 2.90 GHz (4C/8T) |
| GPU | Intel Iris Xe (96 EU @ 1.40 GHz, Gen11 Xe-LP) |
| RAM | 32 GB DDR4 |
| oneAPI | 2026.1 DPC++ |
| Python | 3.12.9, NumPy 1.26.4 (OpenBLAS) |
| OS | Windows 10, WSL2 (Ubuntu) |
The SYCL Tiled GEMM Kernel
Here's the tiled SYCL kernel I benchmarked:
// tiled_gemm.cl - Intel oneAPI SYCL 2020
#include <sycl/sycl.hpp>
template<typename dtype, int TILE>
void tiled_gemm(sycl::queue& q, int n, const dtype* a, const dtype* b, dtype* c) {
q.submit([&](sycl::handler& h) {
h.parallel_for(sycl::nd_range<2>(
sycl::range<2>(n, n),
sycl::range<2>(TILE, TILE)),
[&](sycl::nd_item<2> item) {
int row = item.get_global_id(0);
int col = item.get_global_id(1);
dtype sum = 0;
for (int k = 0; k < n; k++) {
sum += a[row * n + k] * b[k * n + col];
}
c[row * n + col] = sum;
});
});
}
The naive kernel is identical but without tiling.
Benchmark Methodology
- Naive SYCL vs Tiled SYCL — measures kernel-level speedup from tiling
- Tiled SYCL vs NumPy/OpenBLAS — measures end-to-end performance
- Subprocess overhead decomposition — isolates initialization cost
Results
Tiled SYCL vs Naive SYCL (kernel only)
| Matrix | Naive (ms) | Tiled (ms) | Speedup | GFLOP/s (tiled) |
|---|---|---|---|---|
| 128³ | 0.46 | 0.28 | 1.63x | 14.74 |
| 256³ | 0.92 | 0.52 | 1.76x | 64.24 |
| 512³ | 4.10 | 2.19 | 1.87x | 122.43 |
| 1024³ | 23.62 | 15.63 | 1.51x | 137.39 |
Tiled SYCL vs NumPy/OpenBLAS (end-to-end)
| Matrix | NumPy (ms) | NumPy GFLOP/s | SYCL (ms) | Ratio |
|---|---|---|---|---|
| 512³ | 1.70 | 206.12 | 5.50* | 0.308x |
| 1024³ | 8.41 | 286.98 | 19.85* | 0.424x |
- Includes subprocess overhead (155ms initialization)
Subprocess Overhead Decomposition (512³)
| Component | Time (ms) | Percentage |
|---|---|---|
| SYCL kernel compute | 8.59 | 5.2% |
| oneAPI runtime init | 154.99 | 94.8% |
| Total | 165.33 | 100% |
The Subprocess Trap
Every SYCL invocation spawns a new process. The oneAPI runtime initialization takes 155ms per call. To beat NumPy (2.59ms), you'd need negative kernel time — impossible.
Code: The Benchmark Harness
import numpy as np
import time
import subprocess
def benchmark_numpy(n, iterations=10):
times = []
a = np.random.rand(n, n).astype(np.float32)
b = np.random.rand(n, n).astype(np.float32)
for _ in range(iterations):
start = time.perf_counter()
c = np.dot(a, b)
elapsed = time.perf_counter() - start
times.append(elapsed)
median_time = np.median(times)
gflops = 2 * n**3 / median_time / 1e9
return median_time, gflops
def benchmark_sycl(n, iterations=10):
times = []
for _ in range(iterations):
start = time.perf_counter()
result = subprocess.run(
['sycl_runner', f'{n}', f'{n}', f'{n}'],
capture_output=True, text=True
)
elapsed = time.perf_counter() - start
times.append(elapsed)
median_time = np.median(times)
gflops = 2 * n**3 / median_time / 1e9
return median_time, gflops
if __name__ == '__main__':
for n in [64, 128, 256, 512, 1024]:
t_np, g_np = benchmark_numpy(n)
t_sycl, g_sycl = benchmark_sycl(n)
print(f"n={n}: NumPy={t_np*1000:.4f}ms ({g_np:.2f} GF/s), SYCL={t_sycl*1000:.4f}ms ({g_sycl:.2f} GF/s)")
Actionable Findings
What Works:
- Tiled SYCL kernel: 1.5–1.9x faster than naive
- Branchless math: 2.09x speedup in Python
- NumPy/OpenBLAS baseline: 287 GFLOP/s
What Doesn't:
- Subprocess SYCL: 155ms init makes it 64x slower
- FP16 via Python: type conversion overhead (0.02x)
- Small TILE (16): should use 32-64
What To Do:
- Replace subprocess with pybind11
- Increase TILE to 32-64
- Use SYCL USM for zero-copy
- Enable sub_groups for SIMD shuffle
Reproducibility
syclcc -O3 -fsycl tiled_gemm.cl -o tiled_gemm
python benchmark_gemm.py
Tested on Intel Iris Xe (Gen11 Xe-LP) with oneAPI 2026.1.
Top comments (0)
Some comments may only be visible to logged-in visitors. Sign in to view all comments.