Add SYCL backend

README.rst

@@ -1,6 +1,12 @@
 GiMMiK
 ======
-Generator of Matrix Multiplication Kernels - GiMMiK - is a tool for generation of high performance matrix multiplication kernel code for various accelerator platforms. Currently C, CUDA, HIP, ISPC, Metal, and OpenCL are supported.
+Generator of Matrix Multiplication Kernels - GiMMiK - is a tool for generation of high performance matrix multiplication kernel code for various accelerator platforms. Currently C, CUDA, HIP, ISPC, Metal, OpenCL, and SYCL are supported.
+
+The SYCL backend (``platform='sycl'``) emits self-contained launcher functions of
+the form ``sycl::event kname(sycl::queue& q, ...)`` that submit the generated
+kernel to a queue, and can be compiled with any SYCL 2020 compiler (e.g. Intel
+oneAPI DPC++). See ``bench/`` for an OpenCL-vs-SYCL benchmark harness targeting
+Intel GPUs.
 
 What does GiMMiK do?
 --------------------

bench/.gitignore

@@ -0,0 +1,4 @@
+build/
+vbuild/
+*.o
+*.bin

bench/bench_gen.py

@@ -0,0 +1,137 @@
+#!/usr/bin/env python3
+"""Generate OpenCL and SYCL GiMMiK kernels + shared input data for benchmarking.
+
+Produces, under bench/build/:
+  ocl/<name>.cl      - one OpenCL kernel per variant (unique entry name)
+  sycl/<name>.cpp    - one SYCL launcher per variant (unique entry name)
+  B.bin, Cref.bin    - shared fp64 input / reference output (row-major, k x n / m x n)
+  manifest.json      - description of every kernel + problem dims
+"""
+import json
+import os
+import sys
+
+import numpy as np
+
+sys.path.insert(0, os.path.dirname(os.path.dirname(os.path.abspath(__file__))))
+from gimmik import OpenCLMatMul, SYCLMatMul
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+BUILD = os.path.join(HERE, 'build')
+
+
+def make_operator(m, k, sparsity, seed=42):
+    rng = np.random.default_rng(seed)
+    A = rng.standard_normal((m, k))
+    A[rng.random((m, k)) < sparsity] = 0.0
+    # guarantee no fully-zero row so every output column is exercised
+    for j in range(m):
+        if not A[j].any():
+            A[j, rng.integers(k)] = rng.standard_normal()
+    return A
+
+
+def main():
+    m = int(os.environ.get('GMK_M', 32))
+    k = int(os.environ.get('GMK_K', 48))
+    n = int(os.environ.get('GMK_N', 200000))
+    sparsity = float(os.environ.get('GMK_SPARSITY', 0.5))
+
+    ldb = ldc = n
+    A = make_operator(m, k, sparsity)
+    nnz = int((A != 0).sum())
+
+    rng = np.random.default_rng(7)
+    B = rng.standard_normal((k, n))
+    Cref = A @ B
+
+    os.makedirs(os.path.join(BUILD, 'ocl'), exist_ok=True)
+    os.makedirs(os.path.join(BUILD, 'sycl'), exist_ok=True)
+
+    B.astype('<f8').tofile(os.path.join(BUILD, 'B.bin'))
+    Cref.astype('<f8').tofile(os.path.join(BUILD, 'Cref.bin'))
+
+    manifest = {
+        'm': m, 'k': k, 'n': n, 'ldb': ldb, 'ldc': ldc,
+        'nnz': nnz, 'sparsity': sparsity, 'dtype': 'double', 'dsize': 8,
+        'nbix': int(np.count_nonzero(np.any(A != 0, axis=0))),
+        'kernels': [],
+    }
+
+    backends = [('opencl', OpenCLMatMul, 'ocl', 'cl'),
+                ('sycl', SYCLMatMul, 'sycl', 'cpp')]
+
+    for plat, cls, subdir, ext in backends:
+        mm = cls(A, beta=0.0, n=n, ldb=ldb, ldc=ldc)
+        for idx, (src, meta) in enumerate(mm.kernels(np.float64, kname='gimmik_mm')):
+            tpl = meta['tplname']
+            width = meta['width']
+            entry = f'gmk_{plat}_{idx}_{tpl.replace("-", "_")}_w{width}'
+            src = src.replace('gimmik_mm', entry)
+            fname = f'{entry}.{ext}'
+            with open(os.path.join(BUILD, subdir, fname), 'w') as f:
+                f.write(src)
+
+            gws = meta.get('global_work_size')
+            lws = meta.get('local_work_size')
+            manifest['kernels'].append({
+                'platform': plat, 'entry': entry, 'file': f'{subdir}/{fname}',
+                'tpl': tpl, 'width': width,
+                'gws': list(gws) if gws else None,
+                'lws': list(lws) if lws else None,
+                'local_mem_size': meta.get('local_mem_size', 0),
+            })
+
+    with open(os.path.join(BUILD, 'manifest.json'), 'w') as f:
+        json.dump(manifest, f, indent=2)
+
+    # Shared problem dimensions
+    with open(os.path.join(BUILD, 'dims.h'), 'w') as f:
+        f.write('#pragma once\n')
+        f.write(f'#define GMK_M {m}\n#define GMK_K {k}\n#define GMK_N {n}\n')
+        f.write(f'#define GMK_LDB {ldb}\n#define GMK_LDC {ldc}\n')
+        f.write(f'#define GMK_NNZ {nnz}\n#define GMK_NBIX {manifest["nbix"]}\n')
+
+    ocl = [x for x in manifest['kernels'] if x['platform'] == 'opencl']
+    scl = [x for x in manifest['kernels'] if x['platform'] == 'sycl']
+
+    # OpenCL registry (entry name, source file, work sizes)
+    with open(os.path.join(BUILD, 'ocl_registry.h'), 'w') as f:
+        f.write('#pragma once\n')
+        f.write('typedef struct { const char* entry; const char* file; '
+                'const char* tpl; int gdim; size_t g0,g1; size_t l0,l1; } '
+                'OclKernel;\n')
+        f.write('static const OclKernel g_ocl[] = {\n')
+        for x in ocl:
+            g = x['gws']
+            l = x['lws']
+            gdim = len(g)
+            g0 = g[0]
+            g1 = g[1] if gdim > 1 else 1
+            l0 = l[0] if l else 0
+            l1 = l[1] if l else 0
+            f.write(f'  {{"{x["entry"]}", "{x["file"]}", "{x["tpl"]}", '
+                    f'{gdim}, {g0}, {g1}, {l0}, {l1}}},\n')
+        f.write('};\n')
+        f.write(f'static const int g_ocl_n = {len(ocl)};\n')
+
+    # SYCL registry (declarations + function pointer table)
+    with open(os.path.join(BUILD, 'sycl_registry.cpp'), 'w') as f:
+        f.write('#include "sycl_common.hpp"\n')
+        for x in scl:
+            f.write(f'sycl::event {x["entry"]}'
+                    f'(sycl::queue&, const double*, double*);\n')
+        f.write('const SyclKernel g_sycl[] = {\n')
+        for x in scl:
+            f.write(f'  {{"{x["entry"]}", "{x["tpl"]}", &{x["entry"]}}},\n')
+        f.write('};\n')
+        f.write(f'const int g_sycl_n = {len(scl)};\n')
+
+    print(f'A: {m}x{k}, nnz={nnz} ({100*nnz/(m*k):.1f}% dense), '
+          f'used B rows={manifest["nbix"]}')
+    print(f'n={n}  B={B.nbytes/1e6:.1f} MB  C={Cref.nbytes/1e6:.1f} MB')
+    print(f'Generated {len(manifest["kernels"])} kernels into {BUILD}')
+
+
+if __name__ == '__main__':
+    main()

bench/ocl_bench.cpp

@@ -0,0 +1,183 @@
+// OpenCL host benchmark for GiMMiK-generated kernels on the Intel GPU.
+#define CL_TARGET_OPENCL_VERSION 300
+#include <CL/cl.h>
+
+#include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <algorithm>
+#include <string>
+#include <vector>
+
+#include "build/dims.h"
+#include "build/ocl_registry.h"
+
+#define CK(x) do { cl_int _e = (x); if (_e != CL_SUCCESS) { \
+    std::fprintf(stderr, "OpenCL error %d at %s:%d (%s)\n", _e, __FILE__, \
+                 __LINE__, #x); std::exit(1); } } while (0)
+
+static std::vector<double> read_bin(const char* path, size_t count) {
+    std::vector<double> v(count);
+    FILE* f = std::fopen(path, "rb");
+    if (!f) { std::fprintf(stderr, "cannot open %s\n", path); std::exit(1); }
+    if (std::fread(v.data(), sizeof(double), count, f) != count) {
+        std::fprintf(stderr, "short read %s\n", path); std::exit(1);
+    }
+    std::fclose(f);
+    return v;
+}
+
+static std::string read_text(const std::string& path) {
+    FILE* f = std::fopen(path.c_str(), "rb");
+    if (!f) { std::fprintf(stderr, "cannot open %s\n", path.c_str()); std::exit(1); }
+    std::fseek(f, 0, SEEK_END);
+    long sz = std::ftell(f);
+    std::fseek(f, 0, SEEK_SET);
+    std::string s(sz, '\0');
+    if (std::fread(&s[0], 1, sz, f) != (size_t)sz) { std::exit(1); }
+    std::fclose(f);
+    return s;
+}
+
+static cl_device_id pick_gpu(cl_platform_id* out_plat) {
+    cl_uint np = 0;
+    CK(clGetPlatformIDs(0, nullptr, &np));
+    std::vector<cl_platform_id> plats(np);
+    CK(clGetPlatformIDs(np, plats.data(), nullptr));
+    for (auto p : plats) {
+        cl_uint nd = 0;
+        if (clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, 0, nullptr, &nd) != CL_SUCCESS
+            || nd == 0)
+            continue;
+        std::vector<cl_device_id> devs(nd);
+        CK(clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, nd, devs.data(), nullptr));
+        char vendor[256] = {0};
+        clGetDeviceInfo(devs[0], CL_DEVICE_VENDOR, sizeof(vendor), vendor, nullptr);
+        if (std::strstr(vendor, "Intel") || std::strstr(vendor, "INTEL")) {
+            *out_plat = p;
+            return devs[0];
+        }
+    }
+    std::fprintf(stderr, "No Intel GPU found\n");
+    std::exit(1);
+}
+
+int main(int argc, char** argv) {
+    const int reps = argc > 1 ? std::atoi(argv[1]) : 50;
+    const size_t nB = (size_t)GMK_K * GMK_N;
+    const size_t nC = (size_t)GMK_M * GMK_N;
+
+    auto B = read_bin("bench/build/B.bin", nB);
+    auto Cref = read_bin("bench/build/Cref.bin", nC);
+
+    cl_platform_id plat;
+    cl_device_id dev = pick_gpu(&plat);
+    char name[256] = {0};
+    clGetDeviceInfo(dev, CL_DEVICE_NAME, sizeof(name), name, nullptr);
+    std::printf("# OpenCL device: %s\n", name);
+
+    cl_int err;
+    cl_context ctx = clCreateContext(nullptr, 1, &dev, nullptr, nullptr, &err);
+    CK(err);
+    cl_command_queue_properties qp[] = {CL_QUEUE_PROPERTIES,
+                                        CL_QUEUE_PROFILING_ENABLE, 0};
+    cl_command_queue q = clCreateCommandQueueWithProperties(ctx, dev, qp, &err);
+    CK(err);
+
+    cl_mem dB = clCreateBuffer(ctx, CL_MEM_READ_ONLY, nB * sizeof(double),
+                               nullptr, &err); CK(err);
+    cl_mem dC = clCreateBuffer(ctx, CL_MEM_READ_WRITE, nC * sizeof(double),
+                               nullptr, &err); CK(err);
+    CK(clEnqueueWriteBuffer(q, dB, CL_TRUE, 0, nB * sizeof(double), B.data(),
+                            0, nullptr, nullptr));
+
+    std::vector<double> Cout(nC);
+    const double gflop = 2.0 * GMK_NNZ * (double)GMK_N / 1e9;
+    const double gbyte = (double)(GMK_NBIX + GMK_M) * GMK_N * 8.0 / 1e9;
+
+    std::printf("# %-26s %10s %10s %10s %12s\n",
+                "kernel", "ms", "GFLOP/s", "GB/s", "max_relerr");
+
+    for (int ki = 0; ki < g_ocl_n; ki++) {
+        const OclKernel& K = g_ocl[ki];
+        std::string src = read_text(std::string("bench/build/") + K.file);
+        const char* csrc = src.c_str();
+        size_t slen = src.size();
+
+        cl_program prog = clCreateProgramWithSource(ctx, 1, &csrc, &slen, &err);
+        CK(err);
+        cl_int be = clBuildProgram(prog, 1, &dev, "-cl-std=CL2.0", nullptr, nullptr);
+        if (be != CL_SUCCESS) {
+            size_t ls = 0;
+            clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 0, nullptr, &ls);
+            std::string log(ls, '\0');
+            clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, ls, &log[0], nullptr);
+            std::fprintf(stderr, "build failed for %s:\n%s\n", K.entry, log.c_str());
+            std::exit(1);
+        }
+        cl_kernel kern = clCreateKernel(prog, K.entry, &err); CK(err);
+        CK(clSetKernelArg(kern, 0, sizeof(cl_mem), &dB));
+        CK(clSetKernelArg(kern, 1, sizeof(cl_mem), &dC));
+
+        // rounded global work size
+        size_t gws[2] = {(size_t)K.g0, (size_t)K.g1};
+        size_t lws[2] = {(size_t)K.l0, (size_t)K.l1};
+        const size_t* lp = nullptr;
+        if (K.l0 > 0) {
+            lp = lws;
+            gws[0] = ((gws[0] + lws[0] - 1) / lws[0]) * lws[0];
+        }
+
+        // correctness
+        double zero = 0.0;
+        CK(clEnqueueFillBuffer(q, dC, &zero, sizeof(double), 0,
+                               nC * sizeof(double), 0, nullptr, nullptr));
+        CK(clEnqueueNDRangeKernel(q, kern, K.gdim, nullptr, gws, lp, 0,
+                                  nullptr, nullptr));
+        CK(clFinish(q));
+        CK(clEnqueueReadBuffer(q, dC, CL_TRUE, 0, nC * sizeof(double),
+                               Cout.data(), 0, nullptr, nullptr));
+        double maxrel = 0.0;
+        for (size_t i = 0; i < nC; i++) {
+            double d = std::fabs(Cout[i] - Cref[i]);
+            double r = d / (std::fabs(Cref[i]) + 1e-30);
+            maxrel = std::max(maxrel, r);
+        }
+
+        // warmup
+        for (int w = 0; w < 5; w++)
+            CK(clEnqueueNDRangeKernel(q, kern, K.gdim, nullptr, gws, lp, 0,
+                                      nullptr, nullptr));
+        CK(clFinish(q));
+
+        double best_ms = 1e30;
+        for (int r = 0; r < reps; r++) {
+            cl_event ev;
+            CK(clEnqueueNDRangeKernel(q, kern, K.gdim, nullptr, gws, lp, 0,
+                                      nullptr, &ev));
+            CK(clWaitForEvents(1, &ev));
+            cl_ulong t0, t1;
+            clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START,
+                                    sizeof(t0), &t0, nullptr);
+            clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END,
+                                    sizeof(t1), &t1, nullptr);
+            best_ms = std::min(best_ms, (t1 - t0) * 1e-6);
+            clReleaseEvent(ev);
+        }
+
+        std::printf("%-28s %10.4f %10.1f %10.1f %12.2e  %s\n",
+                    K.entry, best_ms, gflop / (best_ms * 1e-3),
+                    gbyte / (best_ms * 1e-3), maxrel,
+                    maxrel < 1e-9 ? "OK" : "FAIL");
+
+        clReleaseKernel(kern);
+        clReleaseProgram(prog);
+    }
+
+    clReleaseMemObject(dB);
+    clReleaseMemObject(dC);
+    clReleaseCommandQueue(q);
+    clReleaseContext(ctx);
+    return 0;
+}

bench/run.sh

@@ -0,0 +1,31 @@
+#!/usr/bin/env bash
+# Build + run the OpenCL vs SYCL GiMMiK benchmark on the Intel GPU.
+# Run from the GiMMiK repository root.
+set -e
+
+REPS=${1:-50}
+ROOT=$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)
+cd "$ROOT"
+
+source /opt/intel/oneapi/compiler/latest/env/vars.sh >/dev/null 2>&1
+ONEAPI=/opt/intel/oneapi/compiler/latest
+
+echo "== Generating kernels =="
+python3 bench/bench_gen.py
+
+echo "== Building OpenCL host =="
+icpx -O3 -std=c++17 -Ibench -I"$ONEAPI/include" \
+     bench/ocl_bench.cpp -o bench/build/ocl_bench \
+     -L"$ONEAPI/lib" -lOpenCL
+
+echo "== Building SYCL host =="
+icpx -fsycl -O3 -std=c++17 -Ibench \
+     bench/sycl_bench.cpp bench/build/sycl_registry.cpp bench/build/sycl/*.cpp \
+     -o bench/build/sycl_bench
+
+echo
+echo "########## OpenCL ##########"
+bench/build/ocl_bench "$REPS"
+echo
+echo "########## SYCL ##########"
+ONEAPI_DEVICE_SELECTOR='level_zero:gpu' bench/build/sycl_bench "$REPS"