From 2e4c95e83934737f22e9cb6336dc53b42e5e46be Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 09:57:07 -0500
Subject: [PATCH 001/201] serve: a llama.cpp engine behind the same boundary
(--engine llama)
The Strata engine is CUDA. On a GPU without it - written for and measured on an Intel Arc Pro
B70 - the same GGUF weights run through llama.cpp's SYCL backend. serve/engine_llama.py puts
llama-server behind Engine.generate(ids, max_new, sampling, cancel), so the OpenAI and
Anthropic endpoints, streaming, tool calls, MCP and the web app are unchanged above it.
It sends the prompt as token ids (the server's own tokenizer, from the same GGUF) with
return_tokens, streams ids back with llama-server's timings mapped onto Strata's `last`,
heartbeats while a long prompt is read, and stops the server by closing the connection when a
client goes away. Attach to a running server ("llama": {"url"}) or spawn one ("llama": {"exe"}).
Not yet: images (Strata's Vision produces embeddings for the CUDA engine; llama-server takes
the image itself) - start without "vision" and image requests are refused cleanly.
Tests: a scripted llama-server (the wire format observed on b29c606), plus a live round trip
when STRATA_LLAMA_URL is set. Measured on the B70 with Coder IQ1_M: 23-25 tok/s decode,
~150 tok/s prefill, correct code, GPU-bound.
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
serve/engine_llama.py | 300 +++++++++++++++++++++++++++++++++++++
serve/server.py | 21 ++-
serve/test_engine_llama.py | 197 ++++++++++++++++++++++++
3 files changed, 516 insertions(+), 2 deletions(-)
create mode 100644 serve/engine_llama.py
create mode 100644 serve/test_engine_llama.py
diff --git a/serve/engine_llama.py b/serve/engine_llama.py
new file mode 100644
index 00000000000..948c8086349
--- /dev/null
+++ b/serve/engine_llama.py
@@ -0,0 +1,300 @@
+"""serve/engine_llama.py - Strata's engine boundary over llama.cpp's `llama-server`.
+
+The Strata engine (`strata --serve`) is CUDA: sm_80+ tensor-core kernels, cudaGraph, cuBLAS. On a GPU without
+CUDA - the Intel Arc Pro B70 this was written for - the same GGUF weights run through llama.cpp's SYCL backend.
+This class puts that engine behind the same boundary the server already speaks:
+
+ Engine.generate(prompt_ids, max_new, sampling, cancel) -> iterator of token ids (None = heartbeat)
+
+so everything above it - the OpenAI and Anthropic endpoints, streaming, tool calls, MCP, the web app - is
+unchanged. It talks to llama-server's `/completion` with the prompt AS TOKEN IDS (the server's own tokenizer
+produced them, from the same GGUF, so the ids agree) and `return_tokens`, which streams ids back one chunk
+at a time with the timings Strata reports.
+
+Two ways to get an engine:
+
+ LlamaEngine.attach("http://127.0.0.1:8081") a llama-server someone else runs (a container, say)
+ LlamaEngine.spawn([exe, *args], log=..., cwd=...) start `llama-server` ourselves and own its lifetime
+
+Both settle `max_context` from `/props` and expose the same `info` the Monitor tab reads.
+
+What it does not do (yet): images. Strata's `Vision` encodes images with `strata-vision` into embeddings the
+CUDA engine reads (`GENI`); llama-server has its own multimodal path (`--mmproj`) that takes the image itself.
+Until that is wired, run the server with no `vision` in the config and it refuses image requests cleanly.
+"""
+from __future__ import annotations
+
+import http.client
+import json
+import os
+import subprocess
+import threading
+import time
+import urllib.error
+import urllib.parse
+import urllib.request
+from pathlib import Path
+from typing import Iterator
+
+HEARTBEAT_S = 10 # yield None this often while the server is quiet (reading a long prompt)
+READY_TIMEOUT_S = 20 * 60 # a 30 GB model from a cold SSD can take this long to load
+STOP_TYPES = {"eos": "stop", "word": "stop", "limit": "length"} # llama-server's stop_type -> Strata's finish
+
+
+class EngineDied(RuntimeError):
+ """The engine went away in the middle of a request."""
+
+
+class LlamaEngine:
+ def __init__(self, url: str, spawn: list[str] | None = None, cwd: str | None = None, log: str | None = None,
+ env: dict | None = None, model_path: str = "", name: str = "llama.cpp"):
+ self.url = url.rstrip("/")
+ self.spawn_cmd, self.cwd, self.log_path, self.env = spawn, cwd, log, env
+ self.model_path = model_path
+ self.proc: subprocess.Popen | None = None
+ self.log = None
+ self.ended = False
+ self.max_context = 0
+ self.can_stop = True # closing the HTTP response stops llama-server's generation
+ self.last: dict = {}
+ self.info: dict = {"engine": name}
+ self.progress = None # llama-server reports no prompt-reading progress on this path
+ self._lock = threading.Lock()
+ if spawn:
+ self._start()
+ self._wait_ready()
+
+ # ------------------------------------------------------------------------------------ construction
+ @classmethod
+ def attach(cls, url: str, model_path: str = "") -> "LlamaEngine":
+ """Use a llama-server that is already running (started by hand, by a container, by the render front end)."""
+ return cls(url, model_path=model_path)
+
+ @classmethod
+ def spawn(cls, cmd: list[str], url: str, cwd: str | None = None, log: str | None = None,
+ env: dict | None = None, model_path: str = "") -> "LlamaEngine":
+ """Start `llama-server` ourselves. `cmd` is the whole command line, `url` where it will listen."""
+ return cls(url, spawn=cmd, cwd=cwd, log=log, env=env, model_path=model_path)
+
+ def _start(self):
+ self.log = open(self.log_path, "a", encoding="utf-8") if self.log_path else subprocess.DEVNULL
+ self.proc = subprocess.Popen(self.spawn_cmd, cwd=self.cwd, stdout=self.log, stderr=subprocess.STDOUT,
+ env=self.env)
+ self.ended = False
+
+ def _wait_ready(self):
+ """`/health` says 503 until the model is loaded; `/props` answers once it is and carries n_ctx."""
+ t0 = time.time()
+ while True:
+ if self.proc is not None and self.proc.poll() is not None:
+ raise EngineDied(f"llama-server exited during load (code {self.proc.returncode})"
+ + (f"; see {self.log_path}" if self.log_path else ""))
+ try:
+ props = self._get("/props", timeout=5)
+ gen = props.get("default_generation_settings") or {}
+ self.max_context = int(gen.get("n_ctx") or 0)
+ if self.max_context:
+ self.info.update(
+ version=(props.get("build_info") or "").strip() or None,
+ n_ctx=self.max_context,
+ model=os.path.basename(props.get("model_path") or self.model_path or ""),
+ modalities=",".join(k for k, v in (props.get("modalities") or {}).items() if v) or "text",
+ )
+ if not self.model_path:
+ self.model_path = props.get("model_path") or ""
+ return
+ except (urllib.error.URLError, http.client.HTTPException, OSError, ValueError):
+ pass
+ if time.time() - t0 > READY_TIMEOUT_S:
+ raise EngineDied(f"llama-server at {self.url} did not become ready in {READY_TIMEOUT_S // 60} min")
+ elapsed = int(time.time() - t0)
+ if elapsed and elapsed % 30 == 0:
+ print(f"[strata] still loading ({elapsed} s) - a 30 GB model takes a few minutes ...", flush=True)
+ time.sleep(1)
+
+ # ------------------------------------------------------------------------------------------ http
+ def _get(self, path: str, timeout: float = 30):
+ with urllib.request.urlopen(self.url + path, timeout=timeout) as r:
+ return json.loads(r.read().decode("utf-8"))
+
+ def _post_stream(self, path: str, body: dict, timeout: float):
+ req = urllib.request.Request(self.url + path, data=json.dumps(body).encode("utf-8"),
+ headers={"content-type": "application/json"})
+ return urllib.request.urlopen(req, timeout=timeout)
+
+ # -------------------------------------------------------------------------------------- generate
+ @staticmethod
+ def sampling_body(sampling: dict) -> dict:
+ """Strata's request sampling -> llama-server's field names. Absent temperature means greedy, as the
+ CUDA engine does; a client's own values always win over the server's defaults."""
+ out: dict = {}
+ t = sampling.get("temperature")
+ out["temperature"] = float(t) if isinstance(t, (int, float)) and not isinstance(t, bool) and t > 0 else 0.0
+ for src, dst, lo in (("top_p", "top_p", 0.0), ("min_p", "min_p", 0.0),
+ ("presence_penalty", "presence_penalty", None),
+ ("frequency_penalty", "frequency_penalty", None),
+ ("repetition_penalty", "repeat_penalty", None)):
+ v = sampling.get(src)
+ if isinstance(v, (int, float)) and not isinstance(v, bool):
+ out[dst] = float(v)
+ tk = sampling.get("top_k")
+ if isinstance(tk, int) and not isinstance(tk, bool) and tk >= 0:
+ out["top_k"] = tk
+ n = sampling.get("penalty_last_n")
+ if isinstance(n, int) and not isinstance(n, bool):
+ out["repeat_last_n"] = n
+ seed = sampling.get("seed")
+ if isinstance(seed, int) and not isinstance(seed, bool):
+ out["seed"] = seed
+ return out
+
+ def generate(self, ids: list[int], max_new: int, sampling: dict, cancel: threading.Event,
+ embeddings=None) -> Iterator[int | None]:
+ """Yields token ids as llama-server streams them, and None every HEARTBEAT_S while it is quiet.
+ A consumer that stops early, or `cancel`, closes the response, which stops the generation server-side."""
+ if embeddings:
+ raise ValueError("images are not supported on the llama.cpp engine yet (start without vision)")
+ body = {"prompt": [int(t) for t in ids], "n_predict": int(max_new), "stream": True,
+ "return_tokens": True, "cache_prompt": True, **self.sampling_body(sampling or {})}
+ self.last = {}
+ self.progress = None
+ try:
+ resp = self._post_stream("/completion", body, timeout=HEARTBEAT_S)
+ except urllib.error.HTTPError as e:
+ detail = e.read().decode("utf-8", "replace")[:300]
+ raise ValueError(f"llama-server refused the request ({e.code}): {detail}") from None
+ except (urllib.error.URLError, OSError) as e:
+ raise EngineDied(f"llama-server is not answering: {e}") from None
+ finished = False
+ try:
+ buf = b""
+ while True:
+ if cancel.is_set():
+ return
+ try:
+ chunk = resp.read1(65536) if hasattr(resp, "read1") else resp.read(65536)
+ except TimeoutError:
+ yield None # quiet: reading a long prompt
+ continue
+ except (http.client.IncompleteRead, OSError) as e:
+ raise EngineDied(f"the connection to llama-server broke mid-answer: {e}") from None
+ if not chunk:
+ raise EngineDied("llama-server closed the stream without finishing")
+ buf += chunk
+ while b"\n" in buf:
+ line, buf = buf.split(b"\n", 1)
+ line = line.strip()
+ if not line.startswith(b"data: "):
+ continue
+ ev = json.loads(line[6:])
+ if "error" in ev:
+ raise ValueError(str(ev["error"]))
+ for t in ev.get("tokens") or []:
+ if cancel.is_set():
+ return
+ yield int(t)
+ if ev.get("stop"):
+ tm = ev.get("timings") or {}
+ self.last = {"generated": int(tm.get("predicted_n") or ev.get("tokens_predicted") or 0),
+ "prompt_tokens": int(tm.get("prompt_n") or ev.get("tokens_evaluated") or 0),
+ "prompt_ms": float(tm.get("prompt_ms") or 0.0),
+ "decode_ms": float(tm.get("predicted_ms") or 0.0),
+ "finish": STOP_TYPES.get(ev.get("stop_type"), "stop"),
+ "reused": int(tm.get("cache_n") or 0)}
+ finished = True
+ return
+ finally:
+ try:
+ resp.close() # early stop: dropping the connection stops the server
+ except Exception:
+ pass
+ if not finished and not self.last:
+ self.last = {"finish": "cancelled"}
+
+ # ------------------------------------------------------------------------------------- lifecycle
+ def alive(self) -> bool:
+ if self.ended:
+ return False
+ if self.proc is not None:
+ return self.proc.poll() is None
+ try:
+ self._get("/health", timeout=3)
+ return True
+ except Exception:
+ return False
+
+ def exit_code(self):
+ if self.proc is None:
+ return None
+ try:
+ return self.proc.wait(timeout=5)
+ except subprocess.TimeoutExpired:
+ return None
+
+ def death_note(self) -> str:
+ tail = ""
+ try:
+ with open(self.log_path, "rb") as f:
+ f.seek(0, 2)
+ f.seek(max(0, f.tell() - 4096))
+ tail = f.read().decode("utf-8", "replace")
+ except (OSError, TypeError):
+ pass
+ low = tail.lower()
+ if "out of memory" in low or "sycl_out_of_memory" in low or "failed to allocate" in low:
+ return "the GPU ran out of memory (a smaller context, or fewer layers on the card)"
+ if "sigsegv" in low or "segmentation fault" in low:
+ return "llama-server crashed (if SYCL_CACHE_PERSISTENT is set, unset it: the JIT cache segfaults on Xe2)"
+ if self.proc is None:
+ return f"the llama-server at {self.url} stopped answering (it is not managed by this server)"
+ return "llama-server stopped; the end of its log:\n" + tail[-1200:]
+
+ def restart(self):
+ if self.proc is None:
+ self.ended = False
+ self._wait_ready() # attached: wait for whoever runs it to bring it back
+ return
+ try:
+ self.proc.kill()
+ except OSError:
+ pass
+ info = dict(self.info)
+ self._start()
+ self._wait_ready()
+ self.info = {**info, **self.info}
+
+ def close(self):
+ self.ended = True
+ if self.proc is None:
+ return
+ try:
+ self.proc.terminate()
+ self.proc.wait(timeout=15)
+ except Exception:
+ try:
+ self.proc.kill()
+ except OSError:
+ pass
+
+
+def engine_from_config(cfg: dict, env: dict | None = None) -> LlamaEngine:
+ """Build the engine the run config describes.
+
+ "llama": {"url": "http://127.0.0.1:8081"} attach to a running llama-server
+ "llama": {"exe": ".../llama-server", "args": [...], "port": 8081} spawn it (args after -m/--port are ours)
+
+ `model` (the GGUF) is optional and only labels the Monitor tab when attached."""
+ L = cfg.get("llama") or {}
+ model = L.get("model") or cfg.get("model") or ""
+ if L.get("exe"):
+ port = int(L.get("port") or 8081)
+ host = L.get("host") or "127.0.0.1"
+ cmd = [L["exe"], *[str(a) for a in (L.get("args") or [])], "--host", host, "--port", str(port)]
+ if model and "-m" not in cmd and "--model" not in cmd:
+ cmd += ["-m", model]
+ return LlamaEngine.spawn(cmd, f"http://{host}:{port}", cwd=cfg.get("cwd"), log=cfg.get("log"), env=env,
+ model_path=model)
+ url = L.get("url")
+ if not url:
+ raise ValueError('the config\'s "llama" block needs either "url" (attach) or "exe" (spawn)')
+ return LlamaEngine.attach(url, model_path=model)
diff --git a/serve/server.py b/serve/server.py
index 7054d799922..b79047b0c84 100644
--- a/serve/server.py
+++ b/serve/server.py
@@ -1636,7 +1636,8 @@ def sampling_defaults_from_config(cfg: dict) -> dict:
def main() -> int:
ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
- ap.add_argument("--engine", choices=["mock", "strata"], default="mock")
+ ap.add_argument("--engine", choices=["mock", "strata", "llama"], default="mock",
+ help="strata = the CUDA engine; llama = llama.cpp's llama-server (any GPU llama.cpp runs on, e.g. Intel Arc via SYCL); see serve/engine_llama.py")
ap.add_argument("--config", help="strata engine config (JSON: exe, args, cwd, tokenizer, model_name), "
"written by setup.py")
ap.add_argument("--host", default=None,
@@ -1673,7 +1674,7 @@ def main() -> int:
a.tokenizer = cfg["tokenizer"]
tok = ByteTokenizer()
tpath = Path(a.tokenizer)
- if a.engine == "strata" and not (tpath / "vocab.json").exists():
+ if a.engine in ("strata", "llama") and not (tpath / "vocab.json").exists():
ap.error(f"the model's tokenizer is missing ({tpath / 'vocab.json'}); run setup again")
if (tpath / "vocab.json").exists():
import strata_tokenizer as ST
@@ -1703,6 +1704,22 @@ def main() -> int:
print(f"[strata] layer split across GPUs {gpu_list(cfg)} ({cfg.get('layer_split') or 'auto'})", flush=True)
engine = StrataEngine(cfg["exe"], engine_args(cfg), cwd=cfg.get("cwd"), log=cfg.get("log"), env=env)
warn_tight_ram(engine.info.get("arena_mib"))
+ elif a.engine == "llama":
+ if not cfg:
+ ap.error("--engine llama needs --config (with a \"llama\" block: url to attach, or exe to spawn)")
+ from serve.engine_llama import engine_from_config as llama_engine
+ vision = None # images: not on this engine yet (engine_llama.py)
+ sampling_defaults = sampling_defaults_from_config(cfg)
+ if sampling_defaults:
+ pretty = ", ".join(f"{k}={v}" for k, v in sampling_defaults.items())
+ print(f"[strata] sampling defaults from the config: {pretty}", flush=True)
+ if cfg.get("vision"):
+ print("[strata] note: \"vision\" in the config is ignored on the llama engine (images not wired yet)",
+ flush=True)
+ print("connecting to llama-server (or starting it: a 30 GB model takes a few minutes) ...", flush=True)
+ engine = llama_engine(cfg, env=child_env(cfg))
+ print(f"[strata] llama.cpp engine up: {engine.info.get('model') or engine.model_path}, "
+ f"context {engine.max_context}, build {engine.info.get('version') or '?'}", flush=True)
else:
engine, vision, sampling_defaults = MockEngine(tok, a.script or [
"Thinking about it.\n\nHello from the mock engine."]), None, {}
diff --git a/serve/test_engine_llama.py b/serve/test_engine_llama.py
new file mode 100644
index 00000000000..77369ff32e1
--- /dev/null
+++ b/serve/test_engine_llama.py
@@ -0,0 +1,197 @@
+"""serve/test_engine_llama.py - the llama.cpp engine boundary, against a scripted llama-server (no GPU).
+
+ python -m unittest serve.test_engine_llama -v
+ STRATA_LLAMA_URL=http://:8081 python -m unittest serve.test_engine_llama -v (+ the live test)
+
+The stub speaks exactly what the real llama-server was observed to send on 2026-09-29 (llama.cpp b29c606):
+`/props` with default_generation_settings.n_ctx, and `/completion` streaming `data: {...}` lines with
+`tokens`, `stop`, and on the last line `stop_type` and `timings`.
+"""
+from __future__ import annotations
+
+import json
+import os
+import sys
+import threading
+import time
+import unittest
+import urllib.request
+from http.server import BaseHTTPRequestHandler, ThreadingHTTPServer
+from pathlib import Path
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1]))
+from serve.engine_llama import EngineDied, LlamaEngine, engine_from_config # noqa: E402
+
+N_CTX = 4096
+
+
+class Stub(BaseHTTPRequestHandler):
+ """A llama-server that streams a scripted list of ids, one per chunk, slowly enough to test cancel."""
+ script: list[int] = []
+ delay_s = 0.0
+ seen: list[dict] = []
+ error: str | None = None
+ stop_type = "eos"
+
+ def log_message(self, *a):
+ pass
+
+ def do_GET(self):
+ if self.path == "/props":
+ body = {"default_generation_settings": {"n_ctx": N_CTX}, "build_info": "stub-1",
+ "model_path": "/models/stub.gguf", "modalities": {"vision": False}}
+ elif self.path == "/health":
+ body = {"status": "ok"}
+ else:
+ self.send_response(404); self.end_headers(); return
+ b = json.dumps(body).encode()
+ self.send_response(200); self.send_header("content-type", "application/json")
+ self.send_header("content-length", str(len(b))); self.end_headers(); self.wfile.write(b)
+
+ def do_POST(self):
+ n = int(self.headers.get("content-length") or 0)
+ req = json.loads(self.rfile.read(n))
+ Stub.seen.append(req)
+ self.send_response(200); self.send_header("content-type", "text/event-stream"); self.end_headers()
+ if Stub.error:
+ self.wfile.write(f"data: {json.dumps({'error': Stub.error})}\n\n".encode()); self.wfile.flush(); return
+ ids = Stub.script[: int(req["n_predict"])]
+ for i, t in enumerate(ids):
+ last = i == len(ids) - 1
+ ev = {"tokens": [t], "content": "x", "stop": last, "tokens_predicted": i + 1,
+ "tokens_evaluated": len(req["prompt"])}
+ if last:
+ ev["stop_type"] = Stub.stop_type if len(ids) == len(Stub.script) else "limit"
+ ev["timings"] = {"prompt_n": len(req["prompt"]), "prompt_ms": 12.5, "predicted_n": len(ids),
+ "predicted_ms": 40.0, "cache_n": 3}
+ try:
+ self.wfile.write(f"data: {json.dumps(ev)}\n\n".encode()); self.wfile.flush()
+ except (BrokenPipeError, ConnectionResetError):
+ Stub.seen[-1]["broken_at"] = i
+ return
+ if Stub.delay_s and not last:
+ time.sleep(Stub.delay_s)
+
+
+class Offline(unittest.TestCase):
+ @classmethod
+ def setUpClass(cls):
+ cls.httpd = ThreadingHTTPServer(("127.0.0.1", 0), Stub)
+ threading.Thread(target=cls.httpd.serve_forever, daemon=True).start()
+ cls.url = f"http://127.0.0.1:{cls.httpd.server_address[1]}"
+
+ @classmethod
+ def tearDownClass(cls):
+ cls.httpd.shutdown(); cls.httpd.server_close()
+
+ def setUp(self):
+ Stub.script, Stub.delay_s, Stub.seen, Stub.error, Stub.stop_type = [11, 22, 33, 44, 55], 0.0, [], None, "eos"
+
+ def test_attach_reads_context_and_info(self):
+ e = LlamaEngine.attach(self.url)
+ self.assertEqual(e.max_context, N_CTX)
+ self.assertEqual(e.info["version"], "stub-1")
+ self.assertEqual(e.info["model"], "stub.gguf")
+ self.assertTrue(e.alive())
+
+ def test_generate_streams_ids_and_reports_timings(self):
+ e = LlamaEngine.attach(self.url)
+ out = [t for t in e.generate([1, 2, 3], 10, {"temperature": 0.7, "top_k": 20}, threading.Event())
+ if t is not None]
+ self.assertEqual(out, [11, 22, 33, 44, 55])
+ self.assertEqual(e.last["generated"], 5)
+ self.assertEqual(e.last["prompt_tokens"], 3)
+ self.assertEqual(e.last["finish"], "stop")
+ self.assertEqual(e.last["reused"], 3)
+ self.assertAlmostEqual(e.last["prompt_ms"], 12.5)
+ req = Stub.seen[-1]
+ self.assertEqual(req["prompt"], [1, 2, 3]) # ids, not text
+ self.assertTrue(req["return_tokens"] and req["stream"] and req["cache_prompt"])
+ self.assertAlmostEqual(req["temperature"], 0.7)
+ self.assertEqual(req["top_k"], 20)
+
+ def test_max_new_caps_and_maps_to_length(self):
+ e = LlamaEngine.attach(self.url)
+ out = [t for t in e.generate([1], 2, {}, threading.Event()) if t is not None]
+ self.assertEqual(out, [11, 22])
+ self.assertEqual(e.last["finish"], "length")
+
+ def test_greedy_when_no_temperature(self):
+ e = LlamaEngine.attach(self.url)
+ list(e.generate([1], 1, {}, threading.Event()))
+ self.assertEqual(Stub.seen[-1]["temperature"], 0.0)
+
+ def test_cancel_stops_early_and_closes_the_stream(self):
+ Stub.delay_s = 0.05
+ e = LlamaEngine.attach(self.url)
+ cancel = threading.Event()
+ got = []
+ for t in e.generate([1], 5, {}, cancel):
+ if t is not None:
+ got.append(t)
+ if len(got) == 2:
+ cancel.set()
+ self.assertEqual(got, [11, 22])
+ self.assertEqual(e.last, {"finish": "cancelled"})
+ time.sleep(0.3) # the stub notices the closed socket on its next write
+ self.assertIn("broken_at", Stub.seen[-1])
+
+ def test_consumer_stopping_early_closes_the_stream(self):
+ Stub.delay_s = 0.05
+ e = LlamaEngine.attach(self.url)
+ gen = e.generate([1], 5, {}, threading.Event())
+ self.assertEqual(next(t for t in gen if t is not None), 11)
+ gen.close()
+ time.sleep(0.3)
+ self.assertIn("broken_at", Stub.seen[-1])
+
+ def test_server_error_event_is_a_value_error(self):
+ Stub.error = "context too long"
+ e = LlamaEngine.attach(self.url)
+ with self.assertRaises(ValueError):
+ list(e.generate([1], 1, {}, threading.Event()))
+
+ def test_images_refused_clearly(self):
+ e = LlamaEngine.attach(self.url)
+ with self.assertRaises(ValueError):
+ list(e.generate([1], 1, {}, threading.Event(), embeddings="/tmp/x.sve"))
+
+ def test_sampling_body_mapping(self):
+ b = LlamaEngine.sampling_body({"temperature": 0, "top_p": 0.9, "min_p": 0.05, "repetition_penalty": 1.1,
+ "presence_penalty": 1.5, "penalty_last_n": 64, "seed": 7, "top_k": 0})
+ self.assertEqual(b, {"temperature": 0.0, "top_p": 0.9, "min_p": 0.05, "repeat_penalty": 1.1,
+ "presence_penalty": 1.5, "repeat_last_n": 64, "seed": 7, "top_k": 0})
+
+ def test_engine_from_config_attach(self):
+ e = engine_from_config({"llama": {"url": self.url, "model": "/m/x.gguf"}})
+ self.assertEqual(e.model_path, "/m/x.gguf")
+ with self.assertRaises(ValueError):
+ engine_from_config({"llama": {}})
+
+ def test_attach_to_nothing_dies_quickly(self):
+ import serve.engine_llama as M
+ old = M.READY_TIMEOUT_S
+ M.READY_TIMEOUT_S = 2
+ try:
+ with self.assertRaises(EngineDied):
+ LlamaEngine.attach("http://127.0.0.1:9")
+ finally:
+ M.READY_TIMEOUT_S = old
+
+
+@unittest.skipUnless(os.environ.get("STRATA_LLAMA_URL"), "set STRATA_LLAMA_URL to test against a live llama-server")
+class Live(unittest.TestCase):
+ def test_round_trip(self):
+ e = LlamaEngine.attach(os.environ["STRATA_LLAMA_URL"])
+ self.assertGreater(e.max_context, 0)
+ ids = json.loads(urllib.request.urlopen(urllib.request.Request(
+ e.url + "/tokenize", data=json.dumps({"content": "def fib(n):", "add_special": True}).encode(),
+ headers={"content-type": "application/json"}), timeout=30).read())["tokens"]
+ out = [t for t in e.generate(ids, 8, {"temperature": 0}, threading.Event()) if t is not None]
+ self.assertEqual(len(out), 8)
+ self.assertEqual(e.last["generated"], 8)
+ self.assertGreater(e.last["decode_ms"], 0)
+
+
+if __name__ == "__main__":
+ unittest.main()
From 43cd6bdb5833b98ce8df0ff3ebf2ae51be06b47a Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 10:12:32 -0500
Subject: [PATCH 002/201] setup: an Intel Arc path (no CUDA), and the run
script starts llama-server itself
Without nvidia-smi, setup now looks for an Intel discrete GPU in sysfs (vendor 0x8086 under xe or
i915, named by PCI id - lspci's database does not know an Arc Pro B70 yet) and takes a path with no
compiler and no CUDA toolkit: download the model, export the tokenizer, find a llama-server
(--llama-server PATH, one on PATH, or the B70 container image via docker), write a config with a
"llama" block and a start-llama.sh, and a run script that uses --engine llama.
The CUDA engine's rules that do not apply are stepped around rather than deleted: the compute
capability and driver-version checks, the experts-in-RAM gates (on llama.cpp the fit is VRAM,
shard 1 = the download minus the 28.8 GB lookup table), the MTP draft download (5 GB the llama
engine cannot use), and the engine's own argument list.
engine_from_config runs the config's start script when nothing answers at the url, so the run
script is one click and closing the server stops the container.
Verified on an Arc Pro B70: setup runs steps 1-7 clean against the downloaded Coder IQ1_M.
docs/INTEL.md has the whole picture and the measurements.
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
README.md | 2 +-
docs/INTEL.md | 93 +++++++++++++
serve/engine_llama.py | 18 ++-
setup.py | 307 +++++++++++++++++++++++++++++++-----------
4 files changed, 340 insertions(+), 80 deletions(-)
create mode 100644 docs/INTEL.md
diff --git a/README.md b/README.md
index b72f6470c6e..1218a71bed6 100644
--- a/README.md
+++ b/README.md
@@ -1,7 +1,7 @@
Strata
Run a 125-billion-parameter AI model on a normal gaming PC
-one NVIDIA card (12-24 GB) + 64 GB of RAM · Windows or Linux · one click to install
+one NVIDIA card (12-24 GB) + 64 GB of RAM · Windows or Linux · one click to install
or an Intel Arc (32 GB) with llama.cpp underneath: docs/INTEL.md

A voxel pagoda garden, 1 shot prompt running on an RTX 5070 with Strata (IQ3_S, 128K context) ·
diff --git a/docs/INTEL.md b/docs/INTEL.md
new file mode 100644
index 00000000000..5bb8ea6e949
--- /dev/null
+++ b/docs/INTEL.md
@@ -0,0 +1,93 @@
+# Strata on an Intel Arc
+
+Strata's engine is CUDA. On an Intel Arc the same model runs behind the same Strata server, with
+llama.cpp's SYCL backend doing the arithmetic. Everything above the engine - the OpenAI and Anthropic
+APIs, streaming, tool calls, MCP, the web app - is unchanged.
+
+Written for and measured on an **Arc Pro B70 (32 GB)** running the Coder (IQ1_M) on Ubuntu 24.04.
+
+## What you get
+
+| | NVIDIA (Strata engine) | Intel Arc (llama.cpp engine) |
+|---|---|---|
+| model files | the same GGUFs | the same GGUFs |
+| where the model lives | experts in RAM, hot ones on the card | **all of shard 1 on the card** (it fits a 32 GB card); shard 2's lookup table paged from disk |
+| RAM needed | 32-64 GB | little: the model is on the card (23 GB was fine) |
+| decode, Coder IQ1_M | 44-51 tok/s on an RTX 5070 | 23-25 tok/s on the B70 |
+| prompt reading | ~1,870 tok/s | ~150 tok/s |
+| speculative decoding (MTP) | yes | no (the GGUF carries no draft layer llama.cpp can use) |
+| images | yes | not yet |
+| context | up to 262K | 32K by default; more as VRAM allows |
+
+The speed gap is the SYCL backend's known state on Battlemage: general matrix multiplies do not use
+the XMX units yet (only the oneDNN flash-attention path does). It is upstream llama.cpp work, not
+something in Strata. Decode is GPU-bound at 90%+ busy, so nothing is falling back to the CPU.
+
+## Setup
+
+ ./setup.sh
+
+Setup notices the Arc (no `nvidia-smi`; an Intel GPU under the `xe` or `i915` driver in sysfs), says
+so, and takes the Intel path: no compiler, no CUDA toolkit. It downloads the model, exports the
+tokenizer, finds a `llama-server`, and writes a config with a `"llama"` block and a start script.
+
+Where the `llama-server` comes from, in order:
+
+1. `--llama-server PATH` - a llama.cpp you built yourself with `-DGGML_SYCL=ON` (needs oneAPI).
+2. `llama-server` on PATH.
+3. A container: `ghcr.io/snailium/llama.cpp-sycl-intel-b70/llama-sycl-b70:stable`, a community SYCL
+ build for Battlemage (llama.cpp 0.4.1, IntelLLVM 2026.1, AOT for bmg-g31). Needs docker and the
+ `render` and `video` groups. Setup writes `start-llama.sh`, which runs it with the GPU passed
+ through and the right environment. The run script starts it when nothing is answering yet.
+
+Then `run-.sh` (or `./setup.sh` again) starts the model. First start of a 30 GB model: about
+two minutes.
+
+## The config
+
+```json
+{
+ "engine": "llama",
+ "tokenizer": ".../packs/coder-iq1_m/tokenizer",
+ "llama": {
+ "url": "http://127.0.0.1:8081",
+ "model": ".../Qwen3.8-Flash-Next-GSQ-RCO-IQ1_M-00001-of-00002.gguf",
+ "start": ".../start-llama.sh",
+ "image": "ghcr.io/snailium/llama.cpp-sycl-intel-b70/llama-sycl-b70:stable"
+ }
+}
+```
+
+`url` attaches to a running llama-server; `start` is run first if nothing answers there. A config
+with `"exe"` and `"args"` instead spawns a native binary directly. See `serve/engine_llama.py`.
+
+## Things that matter on this GPU
+
+- **`SYCL_CACHE_PERSISTENT` must be 0.** The persistent JIT cache segfaults on Xe2 during the first
+ compile. The start script sets it; if you run llama-server by hand, do too.
+- **The whole model goes on the card** (`--n-gpu-layers 999`), except `per_layer_token_embd.weight`,
+ the single 28.8 GB tensor of shard 2. `--override-tensor per_layer_token_embd=CPU` keeps it in host
+ memory, mmapped and paged from the SSD by row, which is how the model's authors serve it. Host RSS
+ stays around 2 GB.
+- **Thinking.** The model reasons before it answers. Strata's web app has the setting; for an API
+ client, put `{"reasoning_effort": "none"}` in `strata-.shared-settings.json` next to the
+ config, or send `chat_template_kwargs: {"enable_thinking": false}` per request. Without one of
+ those a short `max_tokens` is spent entirely inside the think block and the answer looks empty.
+- **`/health` says 503 while loading**; the server polls `/props` instead.
+
+## Not done
+
+- Images. Strata's vision path encodes with `strata-vision` into embeddings the CUDA engine reads;
+ llama-server takes the image itself (`--mmproj`). Wiring that is the next piece.
+- Speculative decoding. Strata's MTP draft is its own packed format; the GGUF has none.
+- The expert-residency profile and n-gram lookup: moot with the model on the card.
+
+## Measured, 2026-09-29, Arc Pro B70, Coder IQ1_M, 32K context, q8_0 KV
+
+| | |
+|---|---|
+| load | ~100 s |
+| VRAM | 28.4 of 32 GB |
+| decode | 23-25 tok/s; GPU 92% busy at 165 W, CPU idle |
+| prefill | 149 tok/s on a 2,701-token code prompt |
+| quality | correct code on every test; matches the NVIDIA path token for token in spirit, not measured |
diff --git a/serve/engine_llama.py b/serve/engine_llama.py
index 948c8086349..efd99bd2910 100644
--- a/serve/engine_llama.py
+++ b/serve/engine_llama.py
@@ -269,7 +269,7 @@ def close(self):
return
try:
self.proc.terminate()
- self.proc.wait(timeout=15)
+ self.proc.wait(timeout=60) # a container unloading 30 GB takes a moment
except Exception:
try:
self.proc.kill()
@@ -297,4 +297,20 @@ def engine_from_config(cfg: dict, env: dict | None = None) -> LlamaEngine:
url = L.get("url")
if not url:
raise ValueError('the config\'s "llama" block needs either "url" (attach) or "exe" (spawn)')
+ start = L.get("start")
+ if start and not _answering(url):
+ # setup wrote a start script (a container that owns the GPU): run it as our child, so the run script is
+ # one click and closing the server stops the container. If something already answers at the url, use it.
+ print(f"[strata] starting llama-server: {start}", flush=True)
+ return LlamaEngine.spawn([start], url, cwd=cfg.get("cwd"), log=cfg.get("log"), env=env, model_path=model)
return LlamaEngine.attach(url, model_path=model)
+
+
+def _answering(url: str) -> bool:
+ try:
+ with urllib.request.urlopen(url.rstrip("/") + "/health", timeout=3) as r:
+ return r.status < 500
+ except urllib.error.HTTPError as e:
+ return e.code < 500 # 503 = loading: someone is already bringing it up
+ except Exception:
+ return False
diff --git a/setup.py b/setup.py
index e9b07f50b73..477e550ad1b 100644
--- a/setup.py
+++ b/setup.py
@@ -274,6 +274,93 @@ def gpus():
return found
+# ---------------------------------------------------------------- Intel Arc: the llama.cpp backend
+# Strata's engine is CUDA. On an Intel Arc the same GGUF runs through llama.cpp's SYCL backend behind the same
+# server (serve/engine_llama.py). Setup then needs no compiler and no CUDA: it downloads the model, exports the
+# tokenizer, finds a llama-server and writes a config that attaches to (or spawns) it. Measured on an Arc Pro
+# B70 with the Coder IQ1_M: the whole 29.6 GB shard on the card, 23-25 tok/s decode. docs/INTEL.md.
+LLAMA_IMAGE = "ghcr.io/snailium/llama.cpp-sycl-intel-b70/llama-sycl-b70:stable" # a community SYCL build
+
+
+# Battlemage / Alchemist PCI device ids -> (name, VRAM GB). lspci's database lags new cards (an Arc Pro B70
+# reads "Intel Corporation Device [8086:e223]"), so sysfs ids are the reliable signal, and the name table is ours.
+INTEL_ARC = {"e223": ("Arc Pro B70", 32.0), "e221": ("Arc Pro B60", 24.0), "e20b": ("Arc B580", 12.0),
+ "e20c": ("Arc B570", 10.0), "e212": ("Arc B50", 16.0),
+ "56a0": ("Arc A770", 16.0), "56a1": ("Arc A750", 8.0), "56a2": ("Arc A580", 8.0),
+ "56a5": ("Arc A380", 6.0), "56a6": ("Arc A310", 4.0), "5690": ("Arc A770M", 16.0)}
+
+
+def intel_gpus():
+ """Intel discrete GPUs from sysfs: vendor 0x8086 under the xe or i915 driver, named by PCI device id. VRAM
+ comes from the id table, else from clinfo's global memory size when it is installed."""
+ found = []
+ for card in sorted(Path("/sys/class/drm").glob("card[0-9]*")):
+ dev = card / "device"
+ try:
+ vendor = (dev / "vendor").read_text().strip().lower()
+ devid = (dev / "device").read_text().strip().lower().replace("0x", "")
+ driver = os.path.basename(os.path.realpath(dev / "driver")) if (dev / "driver").exists() else ""
+ except OSError:
+ continue
+ if vendor != "0x8086" or driver not in ("xe", "i915"):
+ continue
+ name, vram = INTEL_ARC.get(devid, (None, 0.0))
+ if name is None:
+ if driver != "xe": # i915 without a known Arc id is an integrated GPU: not for this
+ continue
+ name = f"Intel GPU {devid} (xe)"
+ if not vram:
+ m = re.search(r"Global memory size\s+(\d+)", out(["clinfo"]))
+ vram = int(m.group(1)) / 2**30 if m else 0.0
+ found.append({"index": len(found), "name": name, "vram_gb": vram, "arch": "sycl", "driver": driver,
+ "pci_id": devid})
+ return found
+
+
+def llama_block(kind, server, gguf, context, port):
+ """The config's "llama" block: how serve/engine_llama.py gets its llama-server.
+
+ A binary is spawned by the server itself. A container is attached to: docker owns the GPU passthrough
+ and the SYCL environment (the persistent JIT cache must be OFF on Xe2, it segfaults), and this writes a
+ start-llama script beside the config that runs it. Either way the whole model goes on the card
+ (--n-gpu-layers 999) except the per-layer embedding table, the single tensor of shard 2, which stays
+ mmapped in host memory and is paged from disk by row, as the model's authors serve it."""
+ common = ["--n-gpu-layers", "999", "--override-tensor", "per_layer_token_embd=CPU",
+ "--ctx-size", str(int(context)), "--cache-type-k", "q8_0", "--cache-type-v", "q8_0",
+ "--flash-attn", "on", "--parallel", "1"]
+ if kind == "exe":
+ return {"exe": server, "args": common, "port": port, "host": "127.0.0.1", "model": str(gguf)}
+ gguf = Path(gguf).resolve()
+ script = ROOT / "start-llama.sh"
+ script.write_text(
+ "#!/bin/sh\n# llama-server for Strata on an Intel Arc, in a container that owns the GPU.\n"
+ f"exec docker run --rm --name strata-llama --device /dev/dri \\\n"
+ f" --group-add $(getent group render | cut -d: -f3) --group-add $(getent group video | cut -d: -f3) \\\n"
+ f" -v \"{gguf.parent}\":/models:ro -p 127.0.0.1:{port}:8080 \\\n"
+ " -e ONEAPI_DEVICE_SELECTOR=level_zero:0 -e SYCL_CACHE_PERSISTENT=0 -e ZES_ENABLE_SYSMAN=1 \\\n"
+ f" {server} -m /models/{gguf.name} " + " ".join(common) + " --port 8080 --host 0.0.0.0\n",
+ encoding="utf-8")
+ script.chmod(0o755)
+ return {"url": f"http://127.0.0.1:{port}", "model": str(gguf), "start": str(script), "image": server}
+
+
+def find_llama_server(explicit=None):
+ """A llama-server for the Intel path: --llama-server PATH, else one on PATH, else the container image if
+ docker is here. Returns (kind, value): ("exe", path) to spawn, ("image", name) to run in docker."""
+ if explicit:
+ if explicit.startswith(("ghcr.io/", "docker.io/")) or "/" in explicit and not Path(explicit).exists():
+ return "image", explicit
+ return "exe", str(Path(explicit).resolve())
+ exe = shutil.which("llama-server")
+ if exe:
+ return "exe", exe
+ if shutil.which("docker") and out(["docker", "image", "inspect", LLAMA_IMAGE, "--format", "{{.Id}}"]).strip():
+ return "image", LLAMA_IMAGE
+ if shutil.which("docker"):
+ return "image", LLAMA_IMAGE # not pulled yet: docker run pulls it (9 GB) on first start
+ return None, None
+
+
GPU_PICK = None # --gpu N (issue #51); None: the card with the most VRAM
SPLIT_MIN_VRAM_GB = 8 # a card sharing a model holds the dense weights and its own
# prompt buffers too (docs/MULTI_GPU.md)
@@ -1262,8 +1349,8 @@ def start(cfg_path: Path, port: int | None, gpu: int | list | None = None, open_
return subprocess.call(cmd)
-def write_run_script(model, cfg_path, port):
- serve = [sys.executable, str(ROOT / "serve" / "server.py"), "--engine", "strata", "--config", str(cfg_path),
+def write_run_script(model, cfg_path, port, engine="strata"):
+ serve = [sys.executable, str(ROOT / "serve" / "server.py"), "--engine", engine, "--config", str(cfg_path),
"--port", str(port), "--open"]
if WIN:
script = ROOT / f"run-{model.lower()}.bat"
@@ -1286,6 +1373,9 @@ def main() -> int:
ap.add_argument("--kv", choices=["int8", "q4_0"],
help="KV cache precision above 8K context: int8 (default) or q4_0 (half the memory, a little less "
"precise)")
+ ap.add_argument("--llama-server", metavar="PATH|IMAGE",
+ help="Intel Arc only: the llama-server to run the model with (a binary built with -DGGML_SYCL=ON, "
+ "or a container image); default: one on PATH, else the B70 image via docker")
ap.add_argument("--vision", choices=["yes", "no", "none", "gpu", "cpu"],
help="let the model read images (yes = the encoder on the GPU)")
ap.add_argument("--experimental-speed-projection", metavar="on|off|GGUF",
@@ -1385,29 +1475,63 @@ def main() -> int:
# ---- 1. the PC
step(1, "checking your PC")
found = gpus()
+ backend = "cuda"
if not found:
- fail("no NVIDIA GPU found (nvidia-smi did not answer)",
- "install the NVIDIA driver from https://www.nvidia.com/drivers and restart the PC")
- if len(found) > 1 or gpu_problem(found[0]) is not None:
- gpu_table(found)
- sel = choose_gpus(a, found) # asked when two or more cards can share the model
- multi = sel if len(sel) > 1 else []
- a.gpu = sel[0] # the main GPU: the checks and the sizing below are its
- GPU_PICK = a.gpu
- gpu = gpu_info(a.gpu)
- chosen = [gpu_info(i) for i in sel]
- gpu["archs"] = sorted({x["arch"] for x in chosen}) # the engine needs code for every one of them
+ intel = intel_gpus()
+ if intel:
+ backend = "llama"
+ g = max(intel, key=lambda x: x["vram_gb"])
+ ok(f"Intel {g['name']} ({g['vram_gb']:.0f} GB): using llama.cpp's SYCL backend instead of the CUDA engine")
+ if a.vision not in (None, "no", "none"):
+ warn("images are not wired on the llama.cpp engine yet: continuing without vision")
+ a.vision = "none"
+ found = intel
+ else:
+ fail("no NVIDIA or Intel Arc GPU found (nvidia-smi did not answer, sycl-ls/lspci show no Arc)",
+ "install the NVIDIA driver from https://www.nvidia.com/drivers and restart the PC")
+ if backend != "llama" and (len(found) > 1 or gpu_problem(found[0]) is not None):
+ gpu_table(found) # CUDA compute-capability checks; not for an Intel card
+ if backend == "llama":
+ sel, multi = [found[0]["index"]], []
+ a.gpu = sel[0]
+ GPU_PICK = a.gpu
+ gpu = {**found[0], "count": len(found), "archs": ["sycl"]}
+ chosen = [gpu] # the one card the model runs on (no layer split here)
+ else:
+ sel = choose_gpus(a, found) # asked when two or more cards can share the model
+ multi = sel if len(sel) > 1 else []
+ a.gpu = sel[0] # the main GPU: the checks and the sizing below are its
+ GPU_PICK = a.gpu
+ gpu = gpu_info(a.gpu)
+ chosen = [gpu_info(i) for i in sel]
+ gpu["archs"] = sorted({x["arch"] for x in chosen}) # the engine needs code for every one of them
if multi:
ok("GPUs: " + " + ".join(gpu_name(x) for x in chosen) + " together (the model's layers are split across them)")
- ok(f"GPU: {gpu['name']}, {gpu['vram_gb']:.1f} GB VRAM, compute capability {cc(gpu)}, driver {gpu['driver']}")
- if driver_major(gpu) < MIN_DRIVER:
- fail(f"the NVIDIA driver is too old ({gpu['driver']}; {MIN_DRIVER} or newer is needed)",
- "update it with the NVIDIA App or from https://www.nvidia.com/drivers, restart, and run this again")
+ if backend == "llama":
+ ok(f"GPU: {gpu['name']}, {gpu['vram_gb']:.1f} GB VRAM, {gpu['driver']} driver, SYCL via llama.cpp")
+ else:
+ ok(f"GPU: {gpu['name']}, {gpu['vram_gb']:.1f} GB VRAM, compute capability {cc(gpu)}, driver {gpu['driver']}")
+ if driver_major(gpu) < MIN_DRIVER:
+ fail(f"the NVIDIA driver is too old ({gpu['driver']}; {MIN_DRIVER} or newer is needed)",
+ "update it with the NVIDIA App or from https://www.nvidia.com/drivers, restart, and run this again")
if gpu["vram_gb"] < 11:
warn("less than 12 GB of VRAM: Strata will run, but most experts stay on the CPU and it will be slow")
ram = ram_gb()
cpu, avx2, avx512 = cpu_info()
need = min(d["ram_gb"] for d in MODELS.values())
+ if backend == "llama":
+ # The CUDA engine keeps every expert in RAM. llama.cpp puts the whole first shard on the card and only
+ # the second shard's lookup table (28.8 GB, read by row) in host memory as page cache, so the fit is a
+ # VRAM question: shard 1 = the download minus that table. Below that, experts spill to RAM (slower).
+ need = 8
+ fit = {m: round(d["download_gb"] - 28.8, 1) for m, d in MODELS.items()}
+ best = [m for m, g in fit.items() if g + 1.5 <= gpu["vram_gb"]]
+ if best:
+ ok(f"VRAM: {gpu['vram_gb']:.0f} GB fits the whole model for " + ", ".join(best) +
+ f" (shard 1: " + ", ".join(f"{m} {g} GB" for m, g in fit.items() if m in best) + ")")
+ else:
+ warn(f"VRAM: {gpu['vram_gb']:.0f} GB is less than any shard 1 (" +
+ ", ".join(f"{m} {g} GB" for m, g in fit.items()) + "): experts will spill to RAM and it will be slow")
if ram < need - 4 and not a.check:
# every model keeps ALL its experts in RAM (23+ GB); VRAM only holds a copy of the most-used ones, so a
# bigger GPU does not lower this
@@ -1451,21 +1575,22 @@ def main() -> int:
fail(f"{fam['title']} has no {a.model} model file", "choose one of: " + ", ".join(names))
for i, m in enumerate(names, 1):
d = MODELS[m]
- fit = "" if ram >= d["ram_gb"] else f" <- needs {d['ram_gb']} GB RAM, you have {ram:.0f}"
+ fit = "" if ram >= d["ram_gb"] or backend == "llama" else f" <- needs {d['ram_gb']} GB RAM, you have {ram:.0f}"
say(f" {i}) {m:8s} {d['about']}; download {d['download_gb']:.0f} GB, uses ~{d['arena_gb']:.0f} GB of RAM{fit}")
rec = str(names.index("IQ3_XXS") + 1) if ram >= 60 and "IQ3_XXS" in names else "1"
model = a.model or names[int(ask("Which size?", [str(i) for i in range(1, len(names) + 1)], rec, a.yes)) - 1]
- if ram < MODELS[model]["ram_gb"] - 4:
- # #125: a warning and a question, not a stop: the user may accept paging (asked, "no" by default, so an
- # unattended --yes install still stops here)
- need_gb, arena = MODELS[model]["ram_gb"], MODELS[model]["arena_gb"]
- warn(f"{model} needs about {need_gb} GB of RAM and this PC has {ram:.0f} GB: its experts alone are "
- f"{arena:.0f} GB and must stay in RAM, so Windows/Linux will page part of them from disk. Expect it "
- "to be much slower, and it may not start at all.")
- say(" A smaller size (Q2_0 or IQ2_XS) fits; more RAM fixes it.")
- if ask(" Install it anyway?", ["y", "n"], "n", a.yes) != "y":
- fail(f"{model} needs about {need_gb} GB of RAM; this PC has {ram:.0f} GB",
- "choose Q2_0 or IQ2_XS, or add RAM")
+ if backend != "llama": # the CUDA engine's experts-in-RAM rule; on llama.cpp the fit is VRAM (step 1)
+ if ram < MODELS[model]["ram_gb"] - 4:
+ # #125: a warning and a question, not a stop: the user may accept paging (asked, "no" by default, so an
+ # unattended --yes install still stops here)
+ need_gb, arena = MODELS[model]["ram_gb"], MODELS[model]["arena_gb"]
+ warn(f"{model} needs about {need_gb} GB of RAM and this PC has {ram:.0f} GB: its experts alone are "
+ f"{arena:.0f} GB and must stay in RAM, so Windows/Linux will page part of them from disk. Expect it "
+ "to be much slower, and it may not start at all.")
+ say(" A smaller size (Q2_0 or IQ2_XS) fits; more RAM fixes it.")
+ if ask(" Install it anyway?", ["y", "n"], "n", a.yes) != "y":
+ fail(f"{model} needs about {need_gb} GB of RAM; this PC has {ram:.0f} GB",
+ "choose Q2_0 or IQ2_XS, or add RAM")
warn(f"installing {model} with {ram:.0f} GB of RAM, as you chose")
ok(f"size: {model}")
tag = fam["tag"] + model # names of the pack, config and start script
@@ -1546,17 +1671,28 @@ def main() -> int:
step(4, "the Strata engine")
llama = get_llama_cpp()
ok(f"llama.cpp {LLAMA_CPP_COMMIT[:7]} (gguf-py, ggml, mtmd)")
- eng = None if a.build else get_prebuilt(a.prebuilt, gpu, vision)
+ if backend == "llama":
+ kind, server = find_llama_server(a.llama_server)
+ if kind is None:
+ fail("no llama-server found for the Intel card",
+ "build llama.cpp with -DGGML_SYCL=ON and put llama-server on PATH, pass --llama-server PATH, "
+ f"or install docker so the image {LLAMA_IMAGE} can run it")
+ ok(f"engine: llama-server ({'container ' if kind == 'image' else ''}{server})")
+ eng, lib_dirs = None, []
+ else:
+ kind = server = None
+ eng = None if a.build else get_prebuilt(a.prebuilt, gpu, vision)
if eng is not None and json.loads((eng / "BUILD.json").read_text()).get("source") != "local":
pip_install(CUDA_WHEELS, "NVIDIA CUDA libraries (cuBLAS, CUDA runtime; ~0.4 GB)")
if vision != "none" and not (eng / VEXE).exists():
warn("the ready-made engine has no image encoder: compiling it")
eng = None
- if eng is None:
- eng = build_engine(gpu, vision, a.yes, llama)
- meta = json.loads((eng / "BUILD.json").read_text())
- lib_dirs = meta.get("cuda_dirs") or cuda_lib_dirs()
- ok(f"engine: {eng / EXE}")
+ if backend != "llama":
+ if eng is None:
+ eng = build_engine(gpu, vision, a.yes, llama)
+ meta = json.loads((eng / "BUILD.json").read_text())
+ lib_dirs = meta.get("cuda_dirs") or cuda_lib_dirs()
+ ok(f"engine: {eng / EXE}")
# ---- 5. the model files
step(5, f"downloading {fam['title']} {model}")
@@ -1602,24 +1738,31 @@ def main() -> int:
if not (pack / "tokenizer" / "vocab.json").exists():
run([sys.executable, str(ROOT / "tools" / "strata_tokenizer.py"), "--gguf", str(shards[0]),
"--out", str(pack)], env=env) # writes /tokenizer/
+ elif backend == "llama":
+ # llama-server reads the GGUF as it is; Strata's server only needs the tokenizer beside it
+ if not (pack / "tokenizer" / "vocab.json").exists():
+ run([sys.executable, str(ROOT / "tools" / "strata_tokenizer.py"), "--gguf", str(shards[0]),
+ "--out", str(pack)], env=env)
elif not (pack / "native_experts.txt").exists() or not (pack / "tokenizer" / "vocab.json").exists():
# every tensor as the GGUF stores it; the experts are read from the GGUF at start (seconds to build)
run([sys.executable, str(ROOT / "tools" / "iq_pack.py"), "--gguf", str(shards[0]), "--out", str(pack)], env=env)
ok(f"model prepared: {pack}")
- mtp = (find_in(roots, "mtp/rt/experts.bin") or data / "mtp/rt/experts.bin").parent.parent
- rt = mtp / "rt"
- if not (rt / "experts.bin").exists():
- say(" The MTP draft layer (speculative decoding, ~2x faster output) comes from the original Qwen checkpoint:")
- say(" only its ~5 GB of MTP tensors are downloaded.")
- run([sys.executable, str(ROOT / "tools" / "mtp_fetch.py"), "fetch", "--out", str(mtp)], env=env)
- run([sys.executable, str(ROOT / "tools" / "mtp_pack.py"), "--src", str(mtp), "--experts", "q2_0",
- "--out", str(mtp / "mtp-q2_0.gguf")], env=env)
- run([sys.executable, str(ROOT / "tools" / "mtp_rt.py"), "--gguf", str(mtp / "mtp-q2_0.gguf"), "--out", str(rt)],
- env=env)
- if not (rt / "draft_vocab.bin").exists():
- shutil.copyfile(ROOT / "data" / "draft_vocab.bin", rt / "draft_vocab.bin")
- ok(f"MTP draft layer: {rt}")
-
+ if backend == "llama":
+ ok("MTP draft layer: not used by the llama.cpp engine (no speculative decoding for this model there)")
+ else:
+ mtp = (find_in(roots, "mtp/rt/experts.bin") or data / "mtp/rt/experts.bin").parent.parent
+ rt = mtp / "rt"
+ if not (rt / "experts.bin").exists():
+ say(" The MTP draft layer (speculative decoding, ~2x faster output) comes from the original Qwen checkpoint:")
+ say(" only its ~5 GB of MTP tensors are downloaded.")
+ run([sys.executable, str(ROOT / "tools" / "mtp_fetch.py"), "fetch", "--out", str(mtp)], env=env)
+ run([sys.executable, str(ROOT / "tools" / "mtp_pack.py"), "--src", str(mtp), "--experts", "q2_0",
+ "--out", str(mtp / "mtp-q2_0.gguf")], env=env)
+ run([sys.executable, str(ROOT / "tools" / "mtp_rt.py"), "--gguf", str(mtp / "mtp-q2_0.gguf"), "--out", str(rt)],
+ env=env)
+ if not (rt / "draft_vocab.bin").exists():
+ shutil.copyfile(ROOT / "data" / "draft_vocab.bin", rt / "draft_vocab.bin")
+ ok(f"MTP draft layer: {rt}")
# ---- 7. the start script
step(7, "writing the start script")
sys.path.insert(0, str(ROOT / "tools"))
@@ -1627,31 +1770,39 @@ def main() -> int:
ple = next((s for s in shards if any(t.name == "per_layer_token_embd.weight" for t in GGUFFile(s).tensors)), None)
if ple is None:
fail("the model has no per_layer_token_embd tensor (is this a Qwen3.8-Flash-Next GGUF?)")
- args = ["--pack", str(pack), "--native", str(shards[0]), "--ple-gguf", str(ple),
- "--expert-profile", str(ROOT / "data" / fam.get("profile", "expert-profile.bin")), "--expert-cache", "auto",
- "--prefill", "auto", "--spec", "4", "--spec-min-p", "0.5", "--mtp", str(rt),
- "--max-context", str(ctx)]
- if ctx > 8192:
- args += ["--kv", kv]
- # KV streaming: from 64K up the whole KV cache lives in RAM and only the part the attention reads (32K positions
- # per layer) stays in VRAM; the VRAM it frees holds more experts (+6% at 128K, +23% at 262K with Q2_0). It
- # costs ~13.7 KB of RAM per context token with 8-bit KV (1.7 GB at 128K), 7.5 KB with 4-bit, so only when it fits.
- kv_ram_gb = ctx * (13 * (576 if kv == "q4_0" else 1056)) / 1e9 # 12 QSA layers + the draft layer
- # the RAM copy must be pinned, and under WSL the NVIDIA driver pins only about 1 GB in all
- if is_wsl() and ctx >= 65536:
- ok("WSL: KV streaming off (the driver pins only about 1 GB of RAM); the KV cache stays in VRAM")
- elif ctx >= 65536 and ram >= MODELS[model]["ram_gb"] + kv_ram_gb + 1:
- args += ["--kv-resident", "32768"]
- ok(f"KV streaming on: the context's KV cache lives in RAM ({kv_ram_gb:.1f} GB), more experts fit in VRAM")
- if vision != "none":
- args += ["--vision", "--vram-reserve-mib", str(VISION[vision]["reserve_mib"])]
- if esp is not None:
- # the package's profile, with llama.cpp's flags (the engine takes the same ones)
- args += ["--control-vector-scaled", f"{esp}:1.0", "--control-vector-layer-range", "4", "44",
- "--cvec-mode", "project", "--cvec-dir", "per-layer"]
- cfg = {"exe": str(eng / EXE), "args": args, "cwd": str(ROOT), "tokenizer": str(pack / "tokenizer"),
- "model_name": f"{fam['name']}-{model.lower()}", "log": str(ROOT / f"strata-{tag.lower()}.log"),
- "lib_dirs": lib_dirs, "port": port}
+ args = []
+ if backend != "llama": # the CUDA engine's command line; nothing here applies to llama-server
+ args = ["--pack", str(pack), "--native", str(shards[0]), "--ple-gguf", str(ple),
+ "--expert-profile", str(ROOT / "data" / fam.get("profile", "expert-profile.bin")), "--expert-cache", "auto",
+ "--prefill", "auto", "--spec", "4", "--spec-min-p", "0.5", "--mtp", str(rt),
+ "--max-context", str(ctx)]
+ if ctx > 8192:
+ args += ["--kv", kv]
+ # KV streaming: from 64K up the whole KV cache lives in RAM and only the part the attention reads (32K positions
+ # per layer) stays in VRAM; the VRAM it frees holds more experts (+6% at 128K, +23% at 262K with Q2_0). It
+ # costs ~13.7 KB of RAM per context token with 8-bit KV (1.7 GB at 128K), 7.5 KB with 4-bit, so only when it fits.
+ kv_ram_gb = ctx * (13 * (576 if kv == "q4_0" else 1056)) / 1e9 # 12 QSA layers + the draft layer
+ # the RAM copy must be pinned, and under WSL the NVIDIA driver pins only about 1 GB in all
+ if is_wsl() and ctx >= 65536:
+ ok("WSL: KV streaming off (the driver pins only about 1 GB of RAM); the KV cache stays in VRAM")
+ elif ctx >= 65536 and ram >= MODELS[model]["ram_gb"] + kv_ram_gb + 1:
+ args += ["--kv-resident", "32768"]
+ ok(f"KV streaming on: the context's KV cache lives in RAM ({kv_ram_gb:.1f} GB), more experts fit in VRAM")
+ if vision != "none":
+ args += ["--vision", "--vram-reserve-mib", str(VISION[vision]["reserve_mib"])]
+ if esp is not None:
+ # the package's profile, with llama.cpp's flags (the engine takes the same ones)
+ args += ["--control-vector-scaled", f"{esp}:1.0", "--control-vector-layer-range", "4", "44",
+ "--cvec-mode", "project", "--cvec-dir", "per-layer"]
+
+ if backend == "llama":
+ cfg = {"engine": "llama", "cwd": str(ROOT), "tokenizer": str(pack / "tokenizer"),
+ "model_name": f"{fam['name']}-{model.lower()}", "log": str(ROOT / f"strata-{tag.lower()}.log"),
+ "port": port, "llama": llama_block(kind, server, shards[0], a.context or 32768, port + 1)}
+ else:
+ cfg = {"exe": str(eng / EXE), "args": args, "cwd": str(ROOT), "tokenizer": str(pack / "tokenizer"),
+ "model_name": f"{fam['name']}-{model.lower()}", "log": str(ROOT / f"strata-{tag.lower()}.log"),
+ "lib_dirs": lib_dirs, "port": port}
if gpu["count"] > 1 or a.gpu is not None:
cfg["gpu"] = gpu["index"] # the engine is told this card (issue #51)
cfg["gpus_asked"] = True # chosen at setup: not asked again at start
@@ -1663,22 +1814,22 @@ def main() -> int:
cfg["host"] = a.host
if a.api_key:
cfg["api_key"] = a.api_key
- if vision != "none":
+ if vision != "none" and backend != "llama":
cfg["vision"] = {"exe": str(eng / VEXE), "mmproj": str(mmproj), "model": str(shards[0]),
"gpu": vision == "gpu", "max_tokens": VISION[vision]["max_tokens"]}
if vision == "cpu":
cfg["vision"]["threads"] = max(1, (os.cpu_count() or 8) // 2)
cfg_path = ROOT / f"strata-{tag.lower()}.json"
- cal = saved_calibration(cfg)
+ cal = None if backend == "llama" else saved_calibration(cfg)
if cal is not None:
sys.path.insert(0, str(ROOT / "tools"))
import calibrate as CAL
cfg["args"] = CAL.apply(cfg["args"], cal.get("settings") or {})
ok("the settings tuned for this PC earlier are used" + (f" ({cal['date']})" if cal.get("date") else ""))
cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8")
- script = write_run_script(tag, cfg_path, port)
+ script = write_run_script(tag, cfg_path, port, engine="llama" if backend == "llama" else "strata")
# offered only when someone answers: --yes installs and adopted earlier installs are not held up by it
- if cal is None and not a.no_start and not a.yes and ask(
+ if cal is None and backend != "llama" and not a.no_start and not a.yes and ask(
"Tune Strata for this PC now? It measures a few engine settings (about 5-10 minutes; the PC is busy "
"meanwhile; later: START-HERE --calibrate)", ["y", "n"], "y", a.yes) == "y":
calibrate_config(cfg_path)
From 8e1354bb08294ad9cc5b68b0a3c2de03424faa26 Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 10:19:23 -0500
Subject: [PATCH 003/201] docs/INTEL.md: --port and --host, learned the hard
way (8080 is open-webui here)
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
docs/INTEL.md | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/docs/INTEL.md b/docs/INTEL.md
index 5bb8ea6e949..72fbdf41718 100644
--- a/docs/INTEL.md
+++ b/docs/INTEL.md
@@ -43,6 +43,10 @@ Where the `llama-server` comes from, in order:
Then `run-.sh` (or `./setup.sh` again) starts the model. First start of a 30 GB model: about
two minutes.
+Two flags worth knowing on a box that already runs things: `--port N` (the default 8080 is also what
+open-webui takes; the llama-server container gets N+1) and `--host 0.0.0.0` to reach it from other
+machines (the default binds localhost only, as upstream does).
+
## The config
```json
From 56653152b31f1d65ab5eb29716d49d0b5d0c4944 Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 11:10:24 -0500
Subject: [PATCH 004/201] README: the Intel Arc port and the rig it was tested
on, up top
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
README.md | 18 ++++++++++++++++++
1 file changed, 18 insertions(+)
diff --git a/README.md b/README.md
index 1218a71bed6..265b6cee072 100644
--- a/README.md
+++ b/README.md
@@ -19,6 +19,24 @@ of a word): faster than you can read.
---
+
+> **Intel Arc port (this fork).** Strata's engine is CUDA. This fork adds a path for Intel Arc cards
+> that runs the same GGUF weights through llama.cpp's SYCL backend behind Strata's own server, web app
+> and APIs. It was developed and tested on this rig, so those are the numbers you will see quoted:
+>
+> | | |
+> |---|---|
+> | GPU | Intel Arc Pro B70, 32 GB (Battlemage G31, PCI `8086:e223`), `xe` kernel driver |
+> | CPU / RAM | AMD Ryzen 7 1700X (8 cores, AVX2, no AVX-512) / 23 GiB |
+> | OS | Ubuntu 24.04.5 LTS, kernel 7.0, Docker 29, Python 3.12 |
+> | Storage | 2 TB PCIe NVMe (the model's 28.8 GB lookup table is paged from it) |
+> | Engine | llama.cpp 0.4.1-dev (b29c606), community SYCL build for the B70, oneAPI 2026.1, `ghcr.io/snailium/llama.cpp-sycl-intel-b70` |
+> | Model | Qwen3.8-Flash-Next GSQ-RCO **Coder**, IQ1_M: all 29.6 GB of shard 1 on the card |
+> | Measured | 23-25 tok/s decode (GPU-bound), ~150 tok/s prompt reading, correct code; see [docs/INTEL.md](docs/INTEL.md) |
+>
+> No CUDA, no compiler: `./setup.sh` detects the Arc and does the rest. Not yet on this path: images
+> and MTP speculative decoding. Everything below this box is the original README for NVIDIA cards.
+
## How fast is it?
Measured on an RTX 5070 (12 GB), a Ryzen 5 7600 and 64 GB of RAM:
From bcfcf40cc8b0cc951da2ee1f4a752562fc2be369 Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 11:14:27 -0500
Subject: [PATCH 005/201] gitignore: start-llama.sh, generated by setup on the
Intel path
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
.gitignore | 3 +++
1 file changed, 3 insertions(+)
diff --git a/.gitignore b/.gitignore
index 16f92c6b471..e714ec2c3f7 100644
--- a/.gitignore
+++ b/.gitignore
@@ -47,3 +47,6 @@ build-vision.bat
*.done
/third_party/*.zip
+
+# Intel Arc path (setup.py): the llama-server container start script
+/start-llama.sh
From 34cae9ec3f70d9d4404bd2020b885b893a0a1b6a Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 11:29:19 -0500
Subject: [PATCH 006/201] docs/INTEL.md: how much context fits on a 32 GB card,
measured
About 20 KiB of KV per token (only 12 of 48 layers are full attention). 32k, 64k, 96k and 128k all
load with the whole model on the card, 128k with 1.6 GB to spare; 160k is under the safety margin
and 256k does not fit - asking for it spilled VRAM into host RAM and took the machine down until its
watchdog reset it, so the doc says so in bold.
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
docs/INTEL.md | 25 ++++++++++++++++++++++++-
1 file changed, 24 insertions(+), 1 deletion(-)
diff --git a/docs/INTEL.md b/docs/INTEL.md
index 72fbdf41718..4e8c42946f1 100644
--- a/docs/INTEL.md
+++ b/docs/INTEL.md
@@ -17,7 +17,7 @@ Written for and measured on an **Arc Pro B70 (32 GB)** running the Coder (IQ1_M)
| prompt reading | ~1,870 tok/s | ~150 tok/s |
| speculative decoding (MTP) | yes | no (the GGUF carries no draft layer llama.cpp can use) |
| images | yes | not yet |
-| context | up to 262K | 32K by default; more as VRAM allows |
+| context | up to 262K | 32K by default; 131K is the measured ceiling on a 32 GB card (see below); 262K does not fit |
The speed gap is the SYCL backend's known state on Battlemage: general matrix multiplies do not use
the XMX units yet (only the oneDNN flash-attention path does). It is upstream llama.cpp work, not
@@ -79,6 +79,29 @@ with `"exe"` and `"args"` instead spawns a native binary directly. See `serve/en
those a short `max_tokens` is spent entirely inside the think block and the answer looks empty.
- **`/health` says 503 while loading**; the server polls `/props` instead.
+## How much context fits
+
+The architecture keeps this small: only every fourth layer is full attention (12 of 48); the other
+36 are gated-delta-net layers with a fixed-size recurrent state that does not grow with context.
+Per token, the 12 attention layers keep 2 KV heads x 256 x (K + V) at q8_0, plus the
+sparse-attention indexer's keys. **Measured: about 20 KiB per token** (VRAM grows 0.6 GB per 32k
+of context with everything else unchanged).
+
+| context | VRAM in use, model loaded | headroom on 32 GB | status |
+|---|---|---|---|
+| 32,768 | 28.4 GB | 3.5 GB | measured, the default |
+| 65,536 | 29.0 GB | 2.9 GB | measured, loads |
+| 98,304 | 29.6 GB | 2.3 GB | measured, loads |
+| 131,072 | 30.3 GB | 1.6 GB | measured, loads: the practical ceiling |
+| 163,840 | ~31.0 GB | ~0.9 GB | not attempted: under the 1.2 GB safety margin |
+| 262,144 | ~32.6 GB | none | **does not fit - asking for it took the host down** |
+
+**Do not ask for more than fits.** On this driver a GPU allocation past VRAM does not fail: the xe
+driver evicts buffers into host RAM, the kernel runs out of memory, and the machine livelocks until
+its hardware watchdog resets it. A 262K request did exactly that on 2026-09-29, and llama.cpp's own
+"failed to fit params" check fired too late to prevent it. Compute the KV size first and leave
+1.5 GB free. Measured ceilings are in the table at the end.
+
## Not done
- Images. Strata's vision path encodes with `strata-vision` into embeddings the CUDA engine reads;
From 0741381c8ec5d927ac48a9aab1d4319c10d1d40f Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 11:36:51 -0500
Subject: [PATCH 007/201] docs: 128k context proven on the B70 - a 105k-token
prompt read at 424 tok/s
Prompt reading scales up with prompt size (150 tok/s at 2.7k, 424 tok/s at 105k), so the full
window costs about five minutes, not fifteen. 30.3 GB in use with 1.5 GB free at 131,072.
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
README.md | 2 +-
docs/INTEL.md | 10 +++++++---
2 files changed, 8 insertions(+), 4 deletions(-)
diff --git a/README.md b/README.md
index 265b6cee072..a2d19ac0960 100644
--- a/README.md
+++ b/README.md
@@ -32,7 +32,7 @@ of a word): faster than you can read.
> | Storage | 2 TB PCIe NVMe (the model's 28.8 GB lookup table is paged from it) |
> | Engine | llama.cpp 0.4.1-dev (b29c606), community SYCL build for the B70, oneAPI 2026.1, `ghcr.io/snailium/llama.cpp-sycl-intel-b70` |
> | Model | Qwen3.8-Flash-Next GSQ-RCO **Coder**, IQ1_M: all 29.6 GB of shard 1 on the card |
-> | Measured | 23-25 tok/s decode (GPU-bound), ~150 tok/s prompt reading, correct code; see [docs/INTEL.md](docs/INTEL.md) |
+> | Measured | 23-25 tok/s decode (GPU-bound); prompt reading 150 tok/s on short prompts, 424 tok/s at 105k tokens; 128k context fits on the card; correct code; see [docs/INTEL.md](docs/INTEL.md) |
>
> No CUDA, no compiler: `./setup.sh` detects the Arc and does the rest. Not yet on this path: images
> and MTP speculative decoding. Everything below this box is the original README for NVIDIA cards.
diff --git a/docs/INTEL.md b/docs/INTEL.md
index 4e8c42946f1..17442d663e2 100644
--- a/docs/INTEL.md
+++ b/docs/INTEL.md
@@ -92,7 +92,7 @@ of context with everything else unchanged).
| 32,768 | 28.4 GB | 3.5 GB | measured, the default |
| 65,536 | 29.0 GB | 2.9 GB | measured, loads |
| 98,304 | 29.6 GB | 2.3 GB | measured, loads |
-| 131,072 | 30.3 GB | 1.6 GB | measured, loads: the practical ceiling |
+| 131,072 | 30.3 GB | 1.5 GB | **measured, the practical ceiling**: a 104,798-token prompt (79% of the window) read in 247 s at 424 tok/s, then answered |
| 163,840 | ~31.0 GB | ~0.9 GB | not attempted: under the 1.2 GB safety margin |
| 262,144 | ~32.6 GB | none | **does not fit - asking for it took the host down** |
@@ -100,7 +100,11 @@ of context with everything else unchanged).
driver evicts buffers into host RAM, the kernel runs out of memory, and the machine livelocks until
its hardware watchdog resets it. A 262K request did exactly that on 2026-09-29, and llama.cpp's own
"failed to fit params" check fired too late to prevent it. Compute the KV size first and leave
-1.5 GB free. Measured ceilings are in the table at the end.
+1.5 GB free.
+
+Prompt reading gets faster with size, because the work batches better: ~150 tok/s on a 2.7k-token
+prompt, 424 tok/s on a 105k-token one. So a full 128k window costs about five minutes to read,
+not fifteen.
## Not done
@@ -116,5 +120,5 @@ its hardware watchdog resets it. A 262K request did exactly that on 2026-09-29,
| load | ~100 s |
| VRAM | 28.4 of 32 GB |
| decode | 23-25 tok/s; GPU 92% busy at 165 W, CPU idle |
-| prefill | 149 tok/s on a 2,701-token code prompt |
+| prefill | 149 tok/s on a 2,701-token prompt; 424 tok/s on a 104,798-token prompt at 131k context |
| quality | correct code on every test; matches the NVIDIA path token for token in spirit, not measured |
From 7021c80ff819ace16d55e2b7246a051eb5d3e4bd Mon Sep 17 00:00:00 2001
From: maxfridbe
Date: Tue, 29 Sep 2026 14:05:57 -0500
Subject: [PATCH 008/201] sycl: the CUDA engine migrated to SYCL - builds with
oneAPI, 19 of 22 kernel parity tests pass on the Arc Pro B70
The port lives in sycl/ and changes no upstream file: sycl/src and sycl/include hold only the files the
migration touched, sycl/CMakeLists.txt takes everything else from the original tree.
How it was made, reproducibly (sycl/tools/): a dev image (the llama.cpp SYCL image + SYCLomatic 2025.3 +
CUDA 12.8 headers from the pip wheels), migrate.sh (compilation database + dpct over 86 translation
units), fixups.py (every hand fix with its reason, idempotent), build.sh (icpx, ninja).
What dpct got wrong and the script repairs: CUDA's null stream became a null sycl::queue* (strata::q_of
maps it to the default queue); __ldg((const float*) p) lost its cast and read one byte; __fadd_rn lost
parentheses; ggml's lookup tables were threaded through kernel parameters with the wrong table per
template; helper headers renamed since 2025.3; cudaGraphUpload, graph introspection and %globaltimer
have no equivalents. Two compiler flags are load-bearing: -fp-model=precise, and the device compiler's
-cl-fp32-correctly-rounded-divide-sqrt - the Arc's fp32 divide is not correctly rounded by default, and
quantize_act_parity went from 303k mismatches to byte-exact with it.
Not ported: the three kernels with inline PTX (mma.sync, ldmatrix, cp.async) take their existing
pre-sm_80 fallback; the launchers refuse the device. The XMX joint_matrix versions are next.
The engine itself starts on the card, loads the pack, the projections, the PLE table and the 23.4 GiB
expert arena, and is then stopped by the test rig's memory guard: Strata keeps the experts as a resident
host copy, and the rig has 23 GiB of RAM (upstream asks for 32). A streaming fill of the VRAM cache
from the GGUF is the next piece; the whole model fits in the 32 GB card.
Co-Authored-By: Claude Fable 5.1
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
---
.gitignore | 4 +
README.md | 1 +
docs/INTEL.md | 68 +
sycl/CMakeLists.txt | 175 +
sycl/dpct_report.apis.csv | 167 +
sycl/dpct_report.stats.csv | 101 +
sycl/include/dpct/atomic.hpp | 868 +++
sycl/include/dpct/bindless_images.hpp | 2061 ++++++
sycl/include/dpct/blas_gemm_utils.hpp | 896 +++
sycl/include/dpct/blas_utils.hpp | 2599 ++++++++
sycl/include/dpct/ccl_utils.hpp | 266 +
sycl/include/dpct/codepin/codepin.hpp | 383 ++
.../dpct/codepin/serialization/basic.hpp | 501 ++
sycl/include/dpct/compat_service.hpp | 122 +
.../dpct/detail/blas_gemm_utils_detail.hpp | 426 ++
.../include/dpct/detail/blas_utils_detail.hpp | 649 ++
sycl/include/dpct/detail/ccl_utils_detail.hpp | 37 +
.../include/dpct/detail/dnnl_utils_detail.hpp | 98 +
.../dpct/detail/group_utils_detail.hpp | 236 +
.../dpct/detail/lapack_utils_detail.hpp | 625 ++
.../dpct/detail/lib_common_utils_detail.hpp | 149 +
sycl/include/dpct/detail/math_detail.hpp | 262 +
sycl/include/dpct/detail/memory_detail.hpp | 984 +++
sycl/include/dpct/detail/rng_utils_detail.hpp | 382 ++
.../dpct/detail/sparse_utils_detail.hpp | 585 ++
sycl/include/dpct/device.hpp | 1010 +++
sycl/include/dpct/dnnl_utils.hpp | 4854 ++++++++++++++
sycl/include/dpct/dpct.hpp | 106 +
sycl/include/dpct/dpl_extras/algorithm.h | 2562 ++++++++
.../dpct/dpl_extras/dpcpp_extensions.h | 422 ++
sycl/include/dpct/dpl_extras/functional.h | 486 ++
.../dpct/dpl_extras/iterator_adaptor.h | 54 +
sycl/include/dpct/dpl_extras/iterators.h | 347 +
sycl/include/dpct/dpl_extras/memory.h | 976 +++
sycl/include/dpct/dpl_extras/numeric.h | 32 +
sycl/include/dpct/dpl_extras/vector.h | 990 +++
sycl/include/dpct/dpl_utils.hpp | 33 +
sycl/include/dpct/fft_utils.hpp | 1249 ++++
sycl/include/dpct/graph.hpp | 195 +
sycl/include/dpct/group_utils.hpp | 1566 +++++
sycl/include/dpct/image.hpp | 1075 ++++
sycl/include/dpct/kernel.hpp | 632 ++
sycl/include/dpct/lapack_utils.hpp | 1286 ++++
sycl/include/dpct/lib_common_utils.hpp | 121 +
sycl/include/dpct/math.hpp | 2807 +++++++++
sycl/include/dpct/memory.hpp | 872 +++
sycl/include/dpct/rng_utils.hpp | 276 +
sycl/include/dpct/shmem_utils.hpp | 56 +
sycl/include/dpct/sparse_utils.hpp | 1838 ++++++
sycl/include/dpct/util.hpp | 1358 ++++
sycl/include/dpct/ze_utils.hpp | 85 +
sycl/include/strata/core/graph.hpp | 118 +
sycl/include/strata/core/mtp.hpp | 149 +
sycl/include/strata/core/on_device.hpp | 38 +
sycl/include/strata/core/remote_experts.hpp | 85 +
sycl/include/strata/core/session.hpp | 417 ++
sycl/include/strata/core/verify.hpp | 239 +
sycl/include/strata/kernels/bf16_bits.hpp | 49 +
sycl/include/strata/kernels/f16_bits.hpp | 99 +
sycl/include/strata/kernels/mrope.hpp | 41 +
sycl/include/strata/kernels/rope.hpp | 53 +
sycl/include/strata/kernels/sampler.hpp | 64 +
sycl/include/strata/sycl_queue.hpp | 14 +
sycl/src/core/concurrent_main.cpp | 272 +
sycl/src/core/device.dp.cpp | 201 +
sycl/src/core/expert_cache.cpp | 465 ++
sycl/src/core/expert_source.cpp | 830 +++
sycl/src/core/graph.cpp | 249 +
sycl/src/core/layer.cpp | 1482 +++++
sycl/src/core/mtp.cpp | 997 +++
sycl/src/core/native_dense.cpp | 235 +
sycl/src/core/native_head.cpp | 221 +
sycl/src/core/overlap_main.cpp | 239 +
sycl/src/core/pinned.dp.cpp | 448 ++
sycl/src/core/remote_experts.cpp | 477 ++
sycl/src/core/session.cpp | 1404 +++++
sycl/src/core/verify.cpp | 1413 +++++
sycl/src/core/weights.cpp | 504 ++
sycl/src/kernels/bf16_gemv_parity.cpp | 191 +
sycl/src/kernels/cuda/bf16_gemv.dp.cpp | 251 +
sycl/src/kernels/cuda/cvec.dp.cpp | 302 +
sycl/src/kernels/cuda/dequant_bf16.dp.cpp | 302 +
sycl/src/kernels/cuda/dequant_s2.dp.cpp | 72 +
sycl/src/kernels/cuda/elementwise.dp.cpp | 658 ++
sycl/src/kernels/cuda/fused_gdn.dp.cpp | 341 +
sycl/src/kernels/cuda/fused_gr.dp.cpp | 633 ++
sycl/src/kernels/cuda/gdn.dp.cpp | 401 ++
sycl/src/kernels/cuda/gr.dp.cpp | 735 +++
sycl/src/kernels/cuda/iq_kernels.dp.cpp | 1401 +++++
sycl/src/kernels/cuda/kv_q4.dp.cpp | 353 ++
sycl/src/kernels/cuda/kv_q8.dp.cpp | 204 +
sycl/src/kernels/cuda/kv_stream.dp.cpp | 459 ++
sycl/src/kernels/cuda/native_bf16.dp.cpp | 345 +
.../src/kernels/cuda/native_flash_attn.dp.cpp | 382 ++
sycl/src/kernels/cuda/native_gdn.dp.cpp | 167 +
.../kernels/cuda/native_gdn_preprocess.dp.cpp | 356 ++
sycl/src/kernels/cuda/native_gr_norm.dp.cpp | 165 +
.../src/kernels/cuda/native_gr_postops.dp.cpp | 247 +
sycl/src/kernels/cuda/native_mmvq.dp.cpp | 2048 ++++++
sycl/src/kernels/cuda/native_moe.dp.cpp | 156 +
.../kernels/cuda/native_ple_postops.dp.cpp | 516 ++
sycl/src/kernels/cuda/native_qsa.dp.cpp | 209 +
.../kernels/cuda/native_qsa_indexer.dp.cpp | 442 ++
sycl/src/kernels/cuda/native_qsa_score.dp.cpp | 287 +
sycl/src/kernels/cuda/native_rope.dp.cpp | 145 +
sycl/src/kernels/cuda/native_router.dp.cpp | 232 +
sycl/src/kernels/cuda/ple.dp.cpp | 674 ++
sycl/src/kernels/cuda/qsa.dp.cpp | 1267 ++++
sycl/src/kernels/cuda/qsa_decode_attn.dp.cpp | 466 ++
sycl/src/kernels/cuda/qsa_prompt_attn.dp.cpp | 1169 ++++
sycl/src/kernels/cuda/qsa_select.dp.cpp | 917 +++
sycl/src/kernels/cuda/quantize_act.dp.cpp | 424 ++
sycl/src/kernels/cuda/rope.dp.cpp | 124 +
sycl/src/kernels/cuda/router_top10.dp.cpp | 394 ++
.../src/kernels/cuda/s2_expert_grouped.dp.cpp | 1352 ++++
sycl/src/kernels/cuda/s2_gemv.dp.cpp | 85 +
sycl/src/kernels/cuda/s2_gemv_fast.dp.cpp | 250 +
sycl/src/kernels/cuda/s2_gemv_q8.dp.cpp | 141 +
sycl/src/kernels/cuda/s2_gemv_quads.dp.cpp | 138 +
sycl/src/kernels/cuda/s_gemv.dp.cpp | 1007 +++
sycl/src/kernels/cuda/sampler.dp.cpp | 547 ++
sycl/src/kernels/cuda/shared_expert.dp.cpp | 596 ++
sycl/src/kernels/cuda/verify_kernels.dp.cpp | 1197 ++++
sycl/src/kernels/cvec_parity.cpp | 242 +
sycl/src/kernels/dequant_bf16_test.cpp | 116 +
sycl/src/kernels/dequant_s2_parity.cpp | 165 +
sycl/src/kernels/elementwise_parity.cpp | 623 ++
sycl/src/kernels/gdn_parity.cpp | 594 ++
sycl/src/kernels/gr_parity.cpp | 930 +++
sycl/src/kernels/iq_parity.cpp | 111 +
sycl/src/kernels/kv_q4_parity.cpp | 429 ++
sycl/src/kernels/kv_q8_parity.cpp | 230 +
sycl/src/kernels/kv_stream_parity.cpp | 344 +
sycl/src/kernels/native_expert_parity.cpp | 262 +
sycl/src/kernels/ple_parity.cpp | 1569 +++++
sycl/src/kernels/qsa_parity.cpp | 1459 +++++
sycl/src/kernels/qsa_prompt_attn_parity.cpp | 263 +
sycl/src/kernels/qsa_select_bench.cpp | 197 +
sycl/src/kernels/quantize_act_parity.cpp | 450 ++
sycl/src/kernels/rope_parity.cpp | 229 +
sycl/src/kernels/router_top10_parity.cpp | 186 +
sycl/src/kernels/s2_gemv_parity.cpp | 205 +
sycl/src/kernels/s2_gemv_q8_parity.cpp | 236 +
sycl/src/kernels/s_gemv_parity.cpp | 754 +++
sycl/src/kernels/s_gemv_q8k_parity.cpp | 321 +
sycl/src/kernels/sampler_parity.cpp | 929 +++
sycl/src/kernels/shared_expert_parity.cpp | 740 +++
sycl/src/prefill/gemm.dp.cpp | 155 +
sycl/src/prefill/ggml_cuda_host.dp.cpp | 104 +
sycl/src/prefill/kernels.dp.cpp | 1428 +++++
sycl/src/prefill/moe_mmq.dp.cpp | 311 +
sycl/src/prefill/prefill.cpp | 2004 ++++++
sycl/src/program/generate.cpp | 5561 +++++++++++++++++
sycl/third_party/ggml/ggml-common.h | 2193 +++++++
sycl/third_party/ggml/ggml-common.h.yaml | 468 ++
sycl/tools/Dockerfile | 9 +
sycl/tools/build.sh | 12 +
sycl/tools/fixups.py | 172 +
sycl/tools/get-cuda-headers.sh | 28 +
sycl/tools/migrate.sh | 45 +
160 files changed, 95558 insertions(+)
create mode 100644 sycl/CMakeLists.txt
create mode 100644 sycl/dpct_report.apis.csv
create mode 100644 sycl/dpct_report.stats.csv
create mode 100644 sycl/include/dpct/atomic.hpp
create mode 100644 sycl/include/dpct/bindless_images.hpp
create mode 100644 sycl/include/dpct/blas_gemm_utils.hpp
create mode 100644 sycl/include/dpct/blas_utils.hpp
create mode 100644 sycl/include/dpct/ccl_utils.hpp
create mode 100644 sycl/include/dpct/codepin/codepin.hpp
create mode 100644 sycl/include/dpct/codepin/serialization/basic.hpp
create mode 100644 sycl/include/dpct/compat_service.hpp
create mode 100644 sycl/include/dpct/detail/blas_gemm_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/blas_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/ccl_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/dnnl_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/group_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/lapack_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/lib_common_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/math_detail.hpp
create mode 100644 sycl/include/dpct/detail/memory_detail.hpp
create mode 100644 sycl/include/dpct/detail/rng_utils_detail.hpp
create mode 100644 sycl/include/dpct/detail/sparse_utils_detail.hpp
create mode 100644 sycl/include/dpct/device.hpp
create mode 100644 sycl/include/dpct/dnnl_utils.hpp
create mode 100644 sycl/include/dpct/dpct.hpp
create mode 100644 sycl/include/dpct/dpl_extras/algorithm.h
create mode 100644 sycl/include/dpct/dpl_extras/dpcpp_extensions.h
create mode 100644 sycl/include/dpct/dpl_extras/functional.h
create mode 100644 sycl/include/dpct/dpl_extras/iterator_adaptor.h
create mode 100644 sycl/include/dpct/dpl_extras/iterators.h
create mode 100644 sycl/include/dpct/dpl_extras/memory.h
create mode 100644 sycl/include/dpct/dpl_extras/numeric.h
create mode 100644 sycl/include/dpct/dpl_extras/vector.h
create mode 100644 sycl/include/dpct/dpl_utils.hpp
create mode 100644 sycl/include/dpct/fft_utils.hpp
create mode 100644 sycl/include/dpct/graph.hpp
create mode 100644 sycl/include/dpct/group_utils.hpp
create mode 100644 sycl/include/dpct/image.hpp
create mode 100644 sycl/include/dpct/kernel.hpp
create mode 100644 sycl/include/dpct/lapack_utils.hpp
create mode 100644 sycl/include/dpct/lib_common_utils.hpp
create mode 100644 sycl/include/dpct/math.hpp
create mode 100644 sycl/include/dpct/memory.hpp
create mode 100644 sycl/include/dpct/rng_utils.hpp
create mode 100644 sycl/include/dpct/shmem_utils.hpp
create mode 100644 sycl/include/dpct/sparse_utils.hpp
create mode 100644 sycl/include/dpct/util.hpp
create mode 100644 sycl/include/dpct/ze_utils.hpp
create mode 100644 sycl/include/strata/core/graph.hpp
create mode 100644 sycl/include/strata/core/mtp.hpp
create mode 100644 sycl/include/strata/core/on_device.hpp
create mode 100644 sycl/include/strata/core/remote_experts.hpp
create mode 100644 sycl/include/strata/core/session.hpp
create mode 100644 sycl/include/strata/core/verify.hpp
create mode 100644 sycl/include/strata/kernels/bf16_bits.hpp
create mode 100644 sycl/include/strata/kernels/f16_bits.hpp
create mode 100644 sycl/include/strata/kernels/mrope.hpp
create mode 100644 sycl/include/strata/kernels/rope.hpp
create mode 100644 sycl/include/strata/kernels/sampler.hpp
create mode 100644 sycl/include/strata/sycl_queue.hpp
create mode 100644 sycl/src/core/concurrent_main.cpp
create mode 100644 sycl/src/core/device.dp.cpp
create mode 100644 sycl/src/core/expert_cache.cpp
create mode 100644 sycl/src/core/expert_source.cpp
create mode 100644 sycl/src/core/graph.cpp
create mode 100644 sycl/src/core/layer.cpp
create mode 100644 sycl/src/core/mtp.cpp
create mode 100644 sycl/src/core/native_dense.cpp
create mode 100644 sycl/src/core/native_head.cpp
create mode 100644 sycl/src/core/overlap_main.cpp
create mode 100644 sycl/src/core/pinned.dp.cpp
create mode 100644 sycl/src/core/remote_experts.cpp
create mode 100644 sycl/src/core/session.cpp
create mode 100644 sycl/src/core/verify.cpp
create mode 100644 sycl/src/core/weights.cpp
create mode 100644 sycl/src/kernels/bf16_gemv_parity.cpp
create mode 100644 sycl/src/kernels/cuda/bf16_gemv.dp.cpp
create mode 100644 sycl/src/kernels/cuda/cvec.dp.cpp
create mode 100644 sycl/src/kernels/cuda/dequant_bf16.dp.cpp
create mode 100644 sycl/src/kernels/cuda/dequant_s2.dp.cpp
create mode 100644 sycl/src/kernels/cuda/elementwise.dp.cpp
create mode 100644 sycl/src/kernels/cuda/fused_gdn.dp.cpp
create mode 100644 sycl/src/kernels/cuda/fused_gr.dp.cpp
create mode 100644 sycl/src/kernels/cuda/gdn.dp.cpp
create mode 100644 sycl/src/kernels/cuda/gr.dp.cpp
create mode 100644 sycl/src/kernels/cuda/iq_kernels.dp.cpp
create mode 100644 sycl/src/kernels/cuda/kv_q4.dp.cpp
create mode 100644 sycl/src/kernels/cuda/kv_q8.dp.cpp
create mode 100644 sycl/src/kernels/cuda/kv_stream.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_bf16.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_flash_attn.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_gdn.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_gdn_preprocess.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_gr_norm.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_gr_postops.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_mmvq.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_moe.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_ple_postops.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_qsa.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_qsa_indexer.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_qsa_score.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_rope.dp.cpp
create mode 100644 sycl/src/kernels/cuda/native_router.dp.cpp
create mode 100644 sycl/src/kernels/cuda/ple.dp.cpp
create mode 100644 sycl/src/kernels/cuda/qsa.dp.cpp
create mode 100644 sycl/src/kernels/cuda/qsa_decode_attn.dp.cpp
create mode 100644 sycl/src/kernels/cuda/qsa_prompt_attn.dp.cpp
create mode 100644 sycl/src/kernels/cuda/qsa_select.dp.cpp
create mode 100644 sycl/src/kernels/cuda/quantize_act.dp.cpp
create mode 100644 sycl/src/kernels/cuda/rope.dp.cpp
create mode 100644 sycl/src/kernels/cuda/router_top10.dp.cpp
create mode 100644 sycl/src/kernels/cuda/s2_expert_grouped.dp.cpp
create mode 100644 sycl/src/kernels/cuda/s2_gemv.dp.cpp
create mode 100644 sycl/src/kernels/cuda/s2_gemv_fast.dp.cpp
create mode 100644 sycl/src/kernels/cuda/s2_gemv_q8.dp.cpp
create mode 100644 sycl/src/kernels/cuda/s2_gemv_quads.dp.cpp
create mode 100644 sycl/src/kernels/cuda/s_gemv.dp.cpp
create mode 100644 sycl/src/kernels/cuda/sampler.dp.cpp
create mode 100644 sycl/src/kernels/cuda/shared_expert.dp.cpp
create mode 100644 sycl/src/kernels/cuda/verify_kernels.dp.cpp
create mode 100644 sycl/src/kernels/cvec_parity.cpp
create mode 100644 sycl/src/kernels/dequant_bf16_test.cpp
create mode 100644 sycl/src/kernels/dequant_s2_parity.cpp
create mode 100644 sycl/src/kernels/elementwise_parity.cpp
create mode 100644 sycl/src/kernels/gdn_parity.cpp
create mode 100644 sycl/src/kernels/gr_parity.cpp
create mode 100644 sycl/src/kernels/iq_parity.cpp
create mode 100644 sycl/src/kernels/kv_q4_parity.cpp
create mode 100644 sycl/src/kernels/kv_q8_parity.cpp
create mode 100644 sycl/src/kernels/kv_stream_parity.cpp
create mode 100644 sycl/src/kernels/native_expert_parity.cpp
create mode 100644 sycl/src/kernels/ple_parity.cpp
create mode 100644 sycl/src/kernels/qsa_parity.cpp
create mode 100644 sycl/src/kernels/qsa_prompt_attn_parity.cpp
create mode 100644 sycl/src/kernels/qsa_select_bench.cpp
create mode 100644 sycl/src/kernels/quantize_act_parity.cpp
create mode 100644 sycl/src/kernels/rope_parity.cpp
create mode 100644 sycl/src/kernels/router_top10_parity.cpp
create mode 100644 sycl/src/kernels/s2_gemv_parity.cpp
create mode 100644 sycl/src/kernels/s2_gemv_q8_parity.cpp
create mode 100644 sycl/src/kernels/s_gemv_parity.cpp
create mode 100644 sycl/src/kernels/s_gemv_q8k_parity.cpp
create mode 100644 sycl/src/kernels/sampler_parity.cpp
create mode 100644 sycl/src/kernels/shared_expert_parity.cpp
create mode 100644 sycl/src/prefill/gemm.dp.cpp
create mode 100644 sycl/src/prefill/ggml_cuda_host.dp.cpp
create mode 100644 sycl/src/prefill/kernels.dp.cpp
create mode 100644 sycl/src/prefill/moe_mmq.dp.cpp
create mode 100644 sycl/src/prefill/prefill.cpp
create mode 100644 sycl/src/program/generate.cpp
create mode 100644 sycl/third_party/ggml/ggml-common.h
create mode 100644 sycl/third_party/ggml/ggml-common.h.yaml
create mode 100644 sycl/tools/Dockerfile
create mode 100644 sycl/tools/build.sh
create mode 100644 sycl/tools/fixups.py
create mode 100755 sycl/tools/get-cuda-headers.sh
create mode 100755 sycl/tools/migrate.sh
diff --git a/.gitignore b/.gitignore
index e714ec2c3f7..7b0f1918882 100644
--- a/.gitignore
+++ b/.gitignore
@@ -50,3 +50,7 @@ build-vision.bat
# Intel Arc path (setup.py): the llama-server container start script
/start-llama.sh
+/build-sycl/
+/compile_commands.json
+/sycl-migrate.log
+/sycl/migrate.log
diff --git a/README.md b/README.md
index a2d19ac0960..d706bea01ec 100644
--- a/README.md
+++ b/README.md
@@ -33,6 +33,7 @@ of a word): faster than you can read.
> | Engine | llama.cpp 0.4.1-dev (b29c606), community SYCL build for the B70, oneAPI 2026.1, `ghcr.io/snailium/llama.cpp-sycl-intel-b70` |
> | Model | Qwen3.8-Flash-Next GSQ-RCO **Coder**, IQ1_M: all 29.6 GB of shard 1 on the card |
> | Measured | 23-25 tok/s decode (GPU-bound); prompt reading 150 tok/s on short prompts, 424 tok/s at 105k tokens; 128k context fits on the card; correct code; see [docs/INTEL.md](docs/INTEL.md) |
+> | SYCL port of the engine | `sycl/`: the CUDA engine migrated to SYCL, builds and links with oneAPI; 19 of 22 kernel parity tests pass on the B70. Not yet generating: the engine's resident host copy of the experts (23.4 GiB) needs more RAM than this rig has - [details](docs/INTEL.md#the-engine-itself-on-intel-the-sycl-port-sycl) |
>
> No CUDA, no compiler: `./setup.sh` detects the Arc and does the rest. Not yet on this path: images
> and MTP speculative decoding. Everything below this box is the original README for NVIDIA cards.
diff --git a/docs/INTEL.md b/docs/INTEL.md
index 17442d663e2..681432bce48 100644
--- a/docs/INTEL.md
+++ b/docs/INTEL.md
@@ -106,6 +106,74 @@ Prompt reading gets faster with size, because the work batches better: ~150 tok/
prompt, 424 tok/s on a 105k-token one. So a full 128k window costs about five minutes to read,
not fifteen.
+## The engine itself on Intel: the SYCL port (`sycl/`)
+
+Everything above runs the model through llama.cpp. This is the other half: Strata's own engine - the 50 CUDA
+kernels and the host code that drives them (streams, events, graph capture, pinned memory) - built for the
+Arc with oneAPI. It is a migration of the tree, not a new backend: the engine has no backend seam to slot
+into.
+
+**Layout.** `sycl/src` and `sycl/include` mirror the tree and hold only the files the port changes; the
+SYCL build (`sycl/CMakeLists.txt`) takes every other source from the original location. No upstream file
+is edited. `sycl/include/dpct/` is the vendored SYCLomatic helper library, so the port builds without the
+migration tool.
+
+**How it was made, so it can be redone.**
+
+1. `sycl/tools/Dockerfile` - the dev image: the llama.cpp SYCL image plus SYCLomatic (`dpct` 2025.3),
+ ninja, and the CUDA 12.8 headers that `sycl/tools/get-cuda-headers.sh` pulls out of NVIDIA's pip wheels
+ (dpct parses CUDA; it needs the headers, not the toolkit).
+2. `sycl/tools/migrate.sh` - writes a compilation database for the 86 CUDA-touching translation units and
+ runs dpct over them. 85 migrate; dpct reports no line it could not migrate, and about 1,400 advisory
+ notes.
+3. `sycl/tools/fixups.py` - what dpct got wrong or could not do, as an idempotent script with a reason per
+ item. The ones that mattered:
+ - CUDA's null stream means the default stream; dpct turned it into a null `sycl::queue*`. Every
+ stream cast now goes through `strata::q_of()` (`sycl/include/strata/sycl_queue.hpp`).
+ - `__ldg((const float*) p)` came out as `*p`, reading one byte of a float scale (two sites, s_gemv).
+ - `__fadd_rn(a, b ? c : d)` lost its parentheses (two sites).
+ - the ggml lookup tables were threaded through kernel parameters with the wrong table per template;
+ they are plain `static const` arrays read from device code, as ggml-sycl does.
+ - helper headers renamed since dpct 2025.3 (`entangle`, `chunked_partition`), graph introspection and
+ `cudaGraphUpload` (no SYCL equivalents), `%globaltimer` (the stage profiler reads zeros).
+4. `sycl/tools/build.sh` - configure + build with icpx inside the image. Two compiler flags are load-bearing:
+ `-fp-model=precise` (icpx defaults to a fast FP model) and `-cl-fp32-correctly-rounded-divide-sqrt`
+ for the device compiler - the Arc's fp32 divide is not correctly rounded by default (OpenCL allows
+ 2.5 ulp), and Strata's quantizers are byte-exact against ggml through `amax / 127`. Measured: without
+ it `quantize_act_parity` has 303k mismatches, with it none.
+
+**Where it stands (2026-09-29).** The whole tree builds and links: the `strata` binary and 22 kernel
+parity tests. On the B70, 19 of the 22 pass byte-for-byte or within their tolerances; the other three
+need model fixtures (`iq_parity`, `ple_parity`) or the tensor-core kernel below (`qsa_prompt_attn`).
+
+| parity test | result |
+|---|---|
+| bf16_gemv, cvec, dequant_s2, elementwise, gdn, gr, kv_q4, kv_q8, kv_stream, qsa, quantize_act, rope, router_top10, s2_gemv, s2_gemv_q8, s_gemv, s_gemv_q8k, sampler, shared_expert | pass |
+| iq_parity, ple_parity | need fixtures (a `logs/iq_fixture` directory, a Q2_0 shard) |
+| qsa_prompt_attn_parity | needs the tensor-core kernel (below) |
+
+**The engine end to end.** `build-sycl/strata` starts on the card, measures the PCIe link (5.3 GB/s
+host-to-device through SYCL), loads the dense weights (1.4 GiB), the 300 native projection matrices
+(2.0 GiB), maps the 28.8 GB PLE table and reads the 23.4 GiB expert arena at 0.89 GiB/s - and is then
+stopped by the test rig's memory guard, because Strata keeps that arena as a resident host copy (its
+"R2.1" design: the CPU computes cache misses from it) and the rig has 23 GiB of RAM. That is upstream's
+own requirement (`ram_gb: 32` for IQ1_M in `setup.py`), not a port defect, and on a 32 GB card it is
+pure redundancy: every one of the 12,288 experts fits in the VRAM cache (23.4 GiB of 32), so nothing
+would ever be computed from the host copy. The missing piece is a streaming fill - load each layer's
+experts from the GGUF into the VRAM slots through a small staging buffer and keep no host arena - which
+is the next piece of engine work for machines with less RAM than experts. No token has been generated by
+the SYCL engine yet; the kernels have been validated one by one.
+
+**Not ported yet.**
+
+- Three kernels carry inline PTX (`mma.sync` tensor-core matrix ops, `ldmatrix`, `cp.async`):
+ `qsa_prompt_attn`, `qsa_select`'s block scores, `native_qsa_score`. Each already has the "older card"
+ fallback the CUDA build uses below sm_80, and the SYCL build takes that path: the launchers refuse the
+ device and their callers use the older kernels. The XMX versions (`joint_matrix`, bf16/f16 on Xe2; no
+ tf32) are the next piece of work and the one that decides prompt speed.
+- The ggml MMQ prefill path (`moe_mmq.cu`) needs llama.cpp's ggml-cuda sources; not built.
+- AOT device code (`-DSTRATA_SYCL_AOT=bmg-g31`) is wired but untested; the port JITs from SPIR-V.
+
## Not done
- Images. Strata's vision path encodes with `strata-vision` into embeddings the CUDA engine reads;
diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt
new file mode 100644
index 00000000000..01698a51d43
--- /dev/null
+++ b/sycl/CMakeLists.txt
@@ -0,0 +1,175 @@
+# Strata on Intel GPUs - the SYCL build of the engine (Phase 4 of the Intel port, docs/INTEL.md).
+#
+# This is a separate project on purpose. The top-level build describes the CUDA engine; every source it lists is
+# taken from ../src unless this directory carries a migrated copy: `src/X.cu` -> `sycl/src/X.dp.cpp`,
+# `src/X.cpp` -> `sycl/src/X.cpp` (only host files that called the CUDA runtime were migrated). Headers resolve
+# the same way through the include order. Upstream files are never edited, so the port is the content of this
+# directory plus the tool that produced it (../../sycl/migrate.sh).
+#
+# cmake -S sycl -B build-sycl -G Ninja -DCMAKE_CXX_COMPILER=icpx && cmake --build build-sycl
+#
+# Device code is JIT-compiled from SPIR-V unless STRATA_SYCL_AOT names a device (bmg-g31 for the Arc Pro B70).
+cmake_minimum_required(VERSION 3.24)
+project(strata_sycl VERSION 0.1.24 LANGUAGES C CXX)
+set(STRATA_ROOT ${CMAKE_CURRENT_SOURCE_DIR}/..)
+set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_EXTENSIONS OFF)
+if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
+ set(CMAKE_BUILD_TYPE Release CACHE STRING "" FORCE)
+endif()
+add_compile_definitions(STRATA_VERSION="${PROJECT_VERSION}-sycl")
+
+set(STRATA_SYCL_AOT "" CACHE STRING "AOT device for the GPU code, e.g. bmg-g31 (empty: SPIR-V, JIT at run time)")
+set(_sycl_flags -fsycl)
+if(STRATA_SYCL_AOT)
+ list(APPEND _sycl_flags -fsycl-targets=spir64_gen)
+ set(_sycl_link -fsycl -fsycl-targets=spir64_gen "-Xs=-device ${STRATA_SYCL_AOT}")
+else()
+ set(_sycl_link -fsycl)
+endif()
+# The Intel GPU's fp32 divide is not correctly rounded by default (OpenCL allows 2.5 ulp). Strata's quantizers
+# are byte-exact against ggml through `amax / 127`, so the device compiler is told to round divides and sqrt
+# properly; measured: quantize_act_parity goes from 303k mismatches to byte-exact. (Same as
+# SYCL_PROGRAM_COMPILE_OPTIONS=-cl-fp32-correctly-rounded-divide-sqrt at run time.)
+list(APPEND _sycl_link "-Xsycl-target-backend=spir64" "-cl-fp32-correctly-rounded-divide-sqrt")
+# Every translation unit gets -fsycl: the migrated host files construct queues and USM, not only the kernels.
+# icpx defaults to a fast floating-point model (reciprocal approximations, reassociation); the kernels are
+# compared byte-for-byte against ggml's, which nvcc built value-safe, so ask for the same here.
+add_compile_options("$<$:${_sycl_flags}>" -fp-model=precise -Wno-unused-parameter
+ -Wno-unused-variable -Wno-deprecated-declarations)
+add_link_options(${_sycl_link})
+
+# --- source resolution: the migrated copy when there is one, the original otherwise -------------------------------
+function(strata_resolve out)
+ set(r "")
+ foreach(f IN LISTS ARGN)
+ if(f MATCHES "\\.cu$")
+ string(REGEX REPLACE "\\.cu$" ".dp.cpp" m "${f}")
+ if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${m})
+ list(APPEND r ${CMAKE_CURRENT_SOURCE_DIR}/${m})
+ else()
+ message(FATAL_ERROR "no migrated copy of ${f} (expected ${m})")
+ endif()
+ elseif(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${f})
+ list(APPEND r ${CMAKE_CURRENT_SOURCE_DIR}/${f})
+ elseif(EXISTS ${STRATA_ROOT}/${f})
+ list(APPEND r ${STRATA_ROOT}/${f})
+ else()
+ message(FATAL_ERROR "no such source: ${f}")
+ endif()
+ endforeach()
+ set(${out} ${r} PARENT_SCOPE)
+endfunction()
+
+add_library(strata_includes INTERFACE)
+target_include_directories(strata_includes INTERFACE
+ ${CMAKE_CURRENT_SOURCE_DIR}/include # migrated headers + dpct/ helpers, first
+ ${STRATA_ROOT}/include)
+target_compile_options(strata_includes INTERFACE -Wall -Wextra)
+
+# --- CPU-side pieces, identical to the top-level build ------------------------------------------------------------
+strata_resolve(_spec src/spec/suffix_drafter.cpp src/spec/controller.cpp src/spec/draft_policy.cpp)
+add_library(strata_spec STATIC ${_spec})
+target_link_libraries(strata_spec PUBLIC strata_includes)
+
+strata_resolve(_cpu src/kernels/cpu/expert.cpp src/kernels/cpu/pool.cpp src/kernels/cpu/expert_layout.cpp
+ src/kernels/cpu/q2_avx2.cpp)
+add_library(strata_kernels_cpu STATIC ${_cpu})
+set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/expert.cpp PROPERTIES COMPILE_OPTIONS
+ "-mavx512f;-mavx512bw;-mavx512vl;-mavx512dq;-mavx512vnni;-mavx512vbmi;-mf16c;-mfma")
+set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/q2_avx2.cpp PROPERTIES COMPILE_OPTIONS "-mavx2;-mfma;-mf16c")
+find_package(Threads REQUIRED)
+target_link_libraries(strata_kernels_cpu PUBLIC strata_includes Threads::Threads)
+
+option(STRATA_NATIVE_EXPERTS "Native (IQ) expert support through ggml-cpu" ON)
+set(STRATA_GGML_DIR "" CACHE PATH "llama.cpp checkout to take ggml from (empty: fetch the pinned commit)")
+if(STRATA_NATIVE_EXPERTS)
+ set(BUILD_SHARED_LIBS OFF CACHE BOOL "" FORCE)
+ set(GGML_STATIC ON CACHE BOOL "" FORCE)
+ set(GGML_CUDA OFF CACHE BOOL "" FORCE)
+ set(GGML_SYCL OFF CACHE BOOL "" FORCE)
+ set(GGML_OPENMP OFF CACHE BOOL "" FORCE)
+ set(GGML_BACKEND_DL OFF CACHE BOOL "" FORCE)
+ set(GGML_BUILD_TESTS OFF CACHE BOOL "" FORCE)
+ set(GGML_BUILD_EXAMPLES OFF CACHE BOOL "" FORCE)
+ set(GGML_NATIVE ON CACHE BOOL "" FORCE)
+ if(STRATA_GGML_DIR)
+ add_subdirectory(${STRATA_GGML_DIR}/ggml ${CMAKE_BINARY_DIR}/ggml EXCLUDE_FROM_ALL)
+ else()
+ include(FetchContent)
+ FetchContent_Declare(strata_llamacpp
+ GIT_REPOSITORY https://github.com/ggml-org/llama.cpp.git
+ GIT_TAG 3cf03257f219afbe7334045ff7c6a06ac68c627d
+ GIT_SHALLOW FALSE)
+ FetchContent_GetProperties(strata_llamacpp)
+ if(NOT strata_llamacpp_POPULATED)
+ FetchContent_Populate(strata_llamacpp)
+ endif()
+ add_subdirectory(${strata_llamacpp_SOURCE_DIR}/ggml ${CMAKE_BINARY_DIR}/ggml EXCLUDE_FROM_ALL)
+ endif()
+ strata_resolve(_cpu_native src/kernels/cpu/native_expert.cpp src/kernels/cpu/iq_avx512.cpp src/kernels/cpu/iq_avx2.cpp)
+ target_sources(strata_kernels_cpu PRIVATE ${_cpu_native})
+ target_include_directories(strata_kernels_cpu PRIVATE ${STRATA_ROOT}/third_party/ggml)
+ set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/iq_avx512.cpp PROPERTIES COMPILE_OPTIONS
+ "-mavx512f;-mavx512bw;-mavx512vl;-mavx512dq;-mavx512vnni;-mavx512vbmi;-mf16c;-mfma")
+ set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/iq_avx2.cpp PROPERTIES COMPILE_OPTIONS "-mavx2;-mfma;-mf16c")
+ target_link_libraries(strata_kernels_cpu PUBLIC ggml-cpu ggml-base)
+ target_compile_definitions(strata_kernels_cpu PUBLIC STRATA_NATIVE_EXPERTS=1)
+endif()
+
+# --- the GPU engine ------------------------------------------------------------------------------------------------
+strata_resolve(_core src/core/device.cu src/core/pinned.cu src/platform/memory.cpp src/core/graph.cpp
+ src/core/weights.cpp src/core/layout.cpp)
+add_library(strata_core STATIC ${_core})
+target_link_libraries(strata_core PUBLIC strata_includes)
+
+set(_kernel_cu dequant_s2 s2_gemv s_gemv s2_gemv_quads s2_gemv_fast router_top10 quantize_act rope sampler s2_gemv_q8
+ s2_expert_grouped shared_expert gr gdn qsa elementwise bf16_gemv native_mmvq native_bf16 native_gr_norm native_moe
+ native_gdn native_gr_postops native_router native_gdn_preprocess native_qsa native_rope native_ple_postops
+ native_flash_attn native_qsa_indexer native_qsa_score ple kv_q8 kv_stream kv_q4 fused_gr cvec qsa_decode_attn
+ qsa_prompt_attn fused_gdn dequant_bf16 qsa_select verify_kernels iq_kernels)
+set(_kernel_files "")
+foreach(k IN LISTS _kernel_cu)
+ list(APPEND _kernel_files src/kernels/cuda/${k}.cu)
+endforeach()
+strata_resolve(_kernels ${_kernel_files} src/kernels/ngram.cpp src/ngram/ple_reader.cpp src/platform/direct_file.cpp)
+add_library(strata_kernels STATIC ${_kernels})
+target_include_directories(strata_kernels PRIVATE ${STRATA_ROOT}/third_party/ggml)
+target_link_libraries(strata_kernels PUBLIC strata_core)
+
+strata_resolve(_engine src/core/layer.cpp src/core/session.cpp src/core/expert_source.cpp src/core/remote_experts.cpp
+ src/core/expert_cache.cpp src/core/native_head.cpp src/core/native_dense.cpp src/core/verify.cpp
+ src/core/mtp.cpp)
+add_library(strata_engine STATIC ${_engine})
+target_link_libraries(strata_engine PUBLIC strata_core strata_kernels strata_kernels_cpu)
+
+# cuBLAS -> oneMKL through dpct::blas (needs the MKL SYCL BLAS domain)
+strata_resolve(_prefill src/prefill/gemm.cu src/prefill/kernels.cu src/prefill/prefill.cpp)
+add_library(strata_prefill STATIC ${_prefill})
+target_compile_options(strata_prefill PUBLIC -qmkl=sequential)
+target_link_options(strata_prefill PUBLIC -qmkl=sequential)
+target_link_libraries(strata_prefill PUBLIC strata_engine)
+
+strata_resolve(_main src/program/generate.cpp)
+add_executable(strata ${_main})
+target_link_libraries(strata PRIVATE strata_engine strata_prefill strata_spec)
+
+# --- the kernel parity self-tests (each runs on the card and compares against a CPU reference) ---------------------
+option(STRATA_SYCL_PARITY "Build the kernel parity tests" ON)
+if(STRATA_SYCL_PARITY)
+ enable_testing()
+ foreach(p IN ITEMS dequant_s2 s2_gemv iq shared_expert gr gdn s2_gemv_q8 sampler rope quantize_act router_top10
+ s_gemv elementwise bf16_gemv s_gemv_q8k qsa ple kv_q8 kv_stream kv_q4 cvec)
+ strata_resolve(_src src/kernels/${p}_parity.cpp)
+ add_executable(${p}_parity ${_src})
+ target_include_directories(${p}_parity PRIVATE ${STRATA_ROOT}/src/kernels) # the oracle .inc fixtures
+ target_link_libraries(${p}_parity PRIVATE strata_kernels)
+ add_test(NAME ${p}_parity COMMAND ${p}_parity --selftest)
+ endforeach()
+ # the prompt-attention parity takes a context length, not --selftest, and needs the tensor-core kernel (not
+ # ported yet: it reports "refused the pools" here). Built, not registered.
+ strata_resolve(_src src/kernels/qsa_prompt_attn_parity.cpp)
+ add_executable(qsa_prompt_attn_parity ${_src})
+ target_link_libraries(qsa_prompt_attn_parity PRIVATE strata_kernels)
+endif()
diff --git a/sycl/dpct_report.apis.csv b/sycl/dpct_report.apis.csv
new file mode 100644
index 00000000000..2f07be623bc
--- /dev/null
+++ b/sycl/dpct_report.apis.csv
@@ -0,0 +1,167 @@
+ API name, Frequency
+"_Bool isnan(float __x)",1
+"__half",18
+"__half __float2half(const float a)",4
+"__half __float2half_rn(const float a)",10
+"__half __ushort_as_half(const unsigned short i)",22
+"__half2",27
+"__half2 __floats2half2_rn(const float a,const float b)",3
+"__half2 __halves2half2(const __half a,const __half b)",3
+"__half2 __hsub2(const __half2 a,const __half2 b)",1
+"__half2 make_half2(const __half x,const __half y)",2
+"char4",4
+"const char * cudaGetErrorString(cudaError_t error)",142
+"cublasHandle_t",6
+"cublasStatus_t",1
+"cublasStatus_t cublasCreate_v2(cublasHandle_t * handle)",2
+"cublasStatus_t cublasDestroy_v2(cublasHandle_t handle)",1
+"cublasStatus_t cublasGemmEx(cublasHandle_t handle,cublasOperation_t transa,cublasOperation_t transb,int m,int n,int k,const void * alpha,const void * A,cudaDataType Atype,int lda,const void * B,cudaDataType Btype,int ldb,const void * beta,void * C,cudaDataType Ctype,int ldc,cublasComputeType_t computeType,cublasGemmAlgo_t algo)",2
+"cublasStatus_t cublasSetMathMode(cublasHandle_t handle,cublasMath_t mode)",2
+"cublasStatus_t cublasSetStream_v2(cublasHandle_t handle,cudaStream_t streamId)",2
+"cublasStatus_t cublasSetWorkspace_v2(cublasHandle_t handle,void * workspace,size_t workspaceSizeInBytes)",3
+"cudaDeviceProp",2
+"cudaError_t",113
+"cudaError_t cudaDeviceGetAttribute(int * value,enum cudaDeviceAttr attr,int device)",6
+"cudaError_t cudaDeviceSynchronize()",70
+"cudaError_t cudaDriverGetVersion(int * driverVersion)",1
+"cudaError_t cudaEventCreate(cudaEvent_t * event)",17
+"cudaError_t cudaEventCreateWithFlags(cudaEvent_t * event,unsigned int flags)",9
+"cudaError_t cudaEventDestroy(cudaEvent_t event)",21
+"cudaError_t cudaEventElapsedTime(float * ms,cudaEvent_t start,cudaEvent_t end)",17
+"cudaError_t cudaEventQuery(cudaEvent_t event)",6
+"cudaError_t cudaEventRecord(cudaEvent_t event,cudaStream_t stream)",39
+"cudaError_t cudaEventSynchronize(cudaEvent_t event)",10
+"cudaError_t cudaFree(void * devPtr)",176
+"cudaError_t cudaFreeHost(void * ptr)",32
+"cudaError_t cudaFuncSetAttribute(void (*)(const float *, const float *, const int *, long, long, long, float *) func,enum cudaFuncAttribute attr,int value)",1
+"cudaError_t cudaFuncSetAttribute(void (*)(const float *, const unsigned short *, const unsigned short *, const int *, int, int, int, float *, float *) func,enum cudaFuncAttribute attr,int value)",1
+"cudaError_t cudaFuncSetAttribute(void (*)(const float *, struct strata::kernels::QsaAttnPools, const int *, const int *, int, int, float, float *, int) func,enum cudaFuncAttribute attr,int value)",2
+"cudaError_t cudaFuncSetAttribute(void (*)(struct strata::kernels::(anonymous namespace)::GrMulti) func,enum cudaFuncAttribute attr,int value)",1
+"cudaError_t cudaGetDevice(int * device)",19
+"cudaError_t cudaGetDeviceCount(int * count)",5
+"cudaError_t cudaGetDeviceProperties_v2(struct cudaDeviceProp * prop,int device)",2
+"cudaError_t cudaGetLastError()",89
+"cudaError_t cudaGraphDestroy(cudaGraph_t graph)",11
+"cudaError_t cudaGraphExecDestroy(cudaGraphExec_t graphExec)",14
+"cudaError_t cudaGraphGetNodes(cudaGraph_t graph,cudaGraphNode_t * nodes,size_t * numNodes)",1
+"cudaError_t cudaGraphInstantiate(cudaGraphExec_t * pGraphExec,cudaGraph_t graph,cudaGraphNode_t * pErrorNode,char * pLogBuffer,size_t bufferSize)",5
+"cudaError_t cudaGraphInstantiate(cudaGraphExec_t * pGraphExec,cudaGraph_t graph,unsigned long long flags)",5
+"cudaError_t cudaGraphLaunch(cudaGraphExec_t graphExec,cudaStream_t stream)",21
+"cudaError_t cudaHostAlloc(void ** pHost,size_t size,unsigned int flags)",15
+"cudaError_t cudaHostGetDevicePointer(void ** pDevice,void * pHost,unsigned int flags)",9
+"cudaError_t cudaHostRegister(void * ptr,size_t size,unsigned int flags)",2
+"cudaError_t cudaHostUnregister(void * ptr)",2
+"cudaError_t cudaMalloc(float ** devPtr,size_t size)",130
+"cudaError_t cudaMalloc(int ** devPtr,size_t size)",17
+"cudaError_t cudaMalloc(signed char ** devPtr,size_t size)",1
+"cudaError_t cudaMalloc(unsigned char ** devPtr,size_t size)",22
+"cudaError_t cudaMalloc(unsigned short ** devPtr,size_t size)",26
+"cudaError_t cudaMalloc(void ** devPtr,size_t size)",42
+"cudaError_t cudaMallocHost(void ** ptr,size_t size)",1
+"cudaError_t cudaMemGetInfo(size_t * free,size_t * total)",11
+"cudaError_t cudaMemcpy(void * dst,const void * src,size_t count,enum cudaMemcpyKind kind)",339
+"cudaError_t cudaMemcpy2DAsync(void * dst,size_t dpitch,const void * src,size_t spitch,size_t width,size_t height,enum cudaMemcpyKind kind,cudaStream_t stream)",1
+"cudaError_t cudaMemcpyAsync(void * dst,const void * src,size_t count,enum cudaMemcpyKind kind,cudaStream_t stream)",70
+"cudaError_t cudaMemcpyToSymbol(const float &[256][4] symbol,const void * src,size_t count,size_t offset,enum cudaMemcpyKind kind)",1
+"cudaError_t cudaMemset(void * devPtr,int value,size_t count)",13
+"cudaError_t cudaMemsetAsync(void * devPtr,int value,size_t count,cudaStream_t stream)",22
+"cudaError_t cudaPeekAtLastError()",3
+"cudaError_t cudaRuntimeGetVersion(int * runtimeVersion)",1
+"cudaError_t cudaSetDevice(int device)",11
+"cudaError_t cudaStreamBeginCapture(cudaStream_t stream,enum cudaStreamCaptureMode mode)",10
+"cudaError_t cudaStreamCreate(cudaStream_t * pStream)",11
+"cudaError_t cudaStreamCreateWithFlags(cudaStream_t * pStream,unsigned int flags)",12
+"cudaError_t cudaStreamDestroy(cudaStream_t stream)",17
+"cudaError_t cudaStreamEndCapture(cudaStream_t stream,cudaGraph_t * pGraph)",10
+"cudaError_t cudaStreamQuery(cudaStream_t stream)",2
+"cudaError_t cudaStreamSynchronize(cudaStream_t stream)",57
+"cudaError_t cudaStreamWaitEvent(cudaStream_t stream,cudaEvent_t event,unsigned int flags)",8
+"cudaEvent_t",30
+"cudaStream_t",306
+"dim3",55
+"double __dadd_rn(double __a,double __b)",6
+"double __ddiv_rn(double __a,double __b)",2
+"double __dmul_rn(double __a,double __b)",3
+"double __dsqrt_rn(double __a)",2
+"double __shfl_down_sync(unsigned int __mask,double __val,unsigned int __offset,int __width)",6
+"double __shfl_sync(unsigned int __mask,double __val,int __offset,int __width)",2
+"double __shfl_xor_sync(unsigned int __mask,double __val,int __offset,int __width)",1
+"double exp(double __a)",10
+"double fmax(double __a,double __b)",1
+"double rint(double __a)",1
+"double sqrt(double __a)",2
+"float __expf(float __a)",33
+"float __fadd_rn(float __a,float __b)",29
+"float __fdiv_rn(float __a,float __b)",3
+"float __fdividef(float __a,float __b)",3
+"float __fmaf_rn(float __a,float __b,float __c)",18
+"float __fmul_rn(float __a,float __b)",24
+"float __fsub_rn(float __a,float __b)",1
+"float __half2float(const __half a)",45
+"float __high2float(const __half2 a)",6
+"float __int_as_float(int __a)",21
+"float __ldg(const float * ptr)",2
+"float __low2float(const __half2 a)",29
+"float __shfl_down_sync(unsigned int __mask,float __val,unsigned int __offset,int __width)",16
+"float __shfl_sync(unsigned int __mask,float __val,int __offset,int __width)",2
+"float __shfl_xor_sync(unsigned int __mask,float __val,int __offset,int __width)",51
+"float __uint_as_float(unsigned int __a)",24
+"float cosf(float __a)",4
+"float exp2f(float __a)",4
+"float expf(float __a)",26
+"float fabsf(float __a)",13
+"float fmaf(float __a,float __b,float __c)",63
+"float fmaxf(float __a,float __b)",48
+"float frexpf(float __a,int * __b)",2
+"float ldexpf(float __a,int __b)",4
+"float log1pf(float __a)",5
+"float logf(float __a)",2
+"float max(const float a,const float b)",2
+"float powf(float __a,float __b)",4
+"float roundf(float __a)",2
+"float rsqrtf(float __a)",25
+"float sinf(float __a)",4
+"float sqrtf(float __a)",4
+"float2",23
+"float2 __half22float2(const __half2 a)",16
+"float2 make_float2(float x,float y)",6
+"float4",64
+"float4 make_float4(float x,float y,float z,float w)",6
+"half",15
+"half2",8
+"int __dp4a(int srcA,int srcB,int c)",54
+"int __float2int_rn(float __a)",2
+"int __float2int_rz(float __a)",1
+"int __ldg(const int * ptr)",22
+"int __popc(unsigned int __a)",5
+"int __shfl_down_sync(unsigned int __mask,int __val,unsigned int __offset,int __width)",7
+"int __shfl_up_sync(unsigned int __mask,int __val,unsigned int __offset,int __width)",4
+"int __shfl_xor_sync(unsigned int __mask,int __val,int __offset,int __width)",2
+"int atomicAdd(int * address,int val)",4
+"int atomicCAS(int * address,int compare,int val)",1
+"int max(int __a,int __b)",1
+"int min(int __a,int __b)",7
+"int2",12
+"int2 make_int2(int x,int y)",7
+"long long __shfl_sync(unsigned int __mask,long long __val,int __offset,int __width)",1
+"size_t __cvta_generic_to_shared(const void * p)",1
+"uint2",19
+"uint2 make_uint2(unsigned int x,unsigned int y)",1
+"uint4",50
+"uint4 __ldg(const uint4 * ptr)",12
+"uint4 make_uint4(unsigned int x,unsigned int y,unsigned int z,unsigned int w)",2
+"unsigned int __ballot_sync(unsigned int mask,int pred)",2
+"unsigned int __byte_perm(unsigned int __a,unsigned int __b,unsigned int __c)",22
+"unsigned int __float_as_uint(float __a)",3
+"unsigned int __umulhi(unsigned int __a,unsigned int __b)",2
+"unsigned int __vcmpne4(unsigned int __a,unsigned int __b)",10
+"unsigned int __vsub4(unsigned int __a,unsigned int __b)",10
+"unsigned int __vsubss4(unsigned int __a,unsigned int __b)",2
+"unsigned int atomicOr(unsigned int * address,unsigned int val)",2
+"unsigned short __half_as_ushort(const __half h)",3
+"unsigned short __ldg(const unsigned short * ptr)",2
+"ushort4",3
+"void __nanosleep(unsigned int ns)",3
+"void __syncwarp(unsigned int mask)",10
+"void __threadfence()",1
+"void __threadfence_system()",6
diff --git a/sycl/dpct_report.stats.csv b/sycl/dpct_report.stats.csv
new file mode 100644
index 00000000000..8bee6c52d6e
--- /dev/null
+++ b/sycl/dpct_report.stats.csv
@@ -0,0 +1,101 @@
+
+File name, LOC migrated to SYCL, LOC migrated to helper functions, LOC not needed to migrate, LOC not able to migrate
+/work/Strata_B70/src/core/session.cpp, 68, 91, 781, 0
+/work/Strata_B70/src/kernels/cuda/native_gdn_preprocess.cu, 43, 2, 155, 0
+/work/Strata_B70/src/kernels/cuda/rope.cu, 14, 4, 78, 0
+/work/Strata_B70/src/kernels/cuda/kv_q4.cu, 49, 12, 171, 0
+/work/Strata_B70/src/kernels/kv_q8_parity.cpp, 2, 18, 122, 0
+/work/Strata_B70/src/kernels/s_gemv_parity.cpp, 6, 44, 585, 0
+/work/Strata_B70/src/kernels/cuda/native_qsa_score.cu, 33, 5, 134, 0
+/work/Strata_B70/src/kernels/cuda/gdn.cu, 75, 15, 159, 0
+/work/Strata_B70/src/kernels/shared_expert_parity.cpp, 12, 60, 411, 0
+/work/Strata_B70/src/kernels/cuda/sampler.cu, 64, 15, 259, 0
+/work/Strata_B70/src/kernels/cuda/qsa_select.cu, 105, 17, 399, 0
+/work/Strata_B70/src/kernels/cuda/native_flash_attn.cu, 47, 5, 161, 0
+/work/Strata_B70/src/kernels/router_top10_parity.cpp, 6, 11, 151, 0
+/work/Strata_B70/src/kernels/bf16_gemv_parity.cpp, 6, 13, 130, 0
+/work/Strata_B70/src/kernels/qsa_prompt_attn_parity.cpp, 11, 18, 181, 0
+/work/Strata_B70/src/kernels/sampler_parity.cpp, 6, 26, 822, 0
+/work/Strata_B70/src/kernels/cuda/router_top10.cu, 35, 11, 181, 0
+/work/Strata_B70/src/kernels/cuda/native_qsa_indexer.cu, 63, 5, 176, 0
+/work/Strata_B70/src/kernels/cuda/qsa_decode_attn.cu, 94, 11, 170, 0
+/work/Strata_B70/src/kernels/cuda/fused_gr.cu, 97, 9, 262, 0
+/work/Strata_B70/src/kernels/cuda/quantize_act.cu, 53, 15, 246, 0
+/work/Strata_B70/src/kernels/cuda/native_gr_postops.cu, 27, 2, 87, 0
+/work/Strata_B70/src/core/expert_source.cpp, 6, 4, 783, 0
+/work/Strata_B70/src/kernels/kv_stream_parity.cpp, 3, 21, 211, 0
+/work/Strata_B70/src/kernels/cuda/gr.cu, 94, 12, 327, 0
+/work/Strata_B70/src/kernels/cuda/native_ple_postops.cu, 60, 4, 158, 0
+/work/Strata_B70/src/kernels/native_expert_parity.cpp, 1, 2, 259, 0
+/work/Strata_B70/src/prefill/ggml_cuda_host.cu, 7, 5, 71, 0
+/work/Strata_B70/src/kernels/cuda/s2_gemv_fast.cu, 45, 6, 103, 0
+/work/Strata_B70/src/kernels/cuda/native_qsa.cu, 19, 2, 107, 0
+/work/Strata_B70/src/kernels/qsa_parity.cpp, 20, 37, 1327, 0
+/work/Strata_B70/src/kernels/cuda/dequant_bf16.cu, 19, 2, 225, 0
+/work/Strata_B70/src/kernels/cuda/kv_stream.cu, 46, 12, 201, 0
+/work/Strata_B70/src/prefill/moe_mmq.cu, 38, 6, 143, 0
+/work/Strata_B70/src/kernels/cvec_parity.cpp, 6, 12, 206, 0
+/work/Strata_B70/src/core/expert_cache.cpp, 8, 16, 300, 0
+/work/Strata_B70/src/kernels/cuda/verify_kernels.cu, 177, 19, 359, 0
+/work/Strata_B70/include/strata/kernels/mrope.hpp, 3, 0, 31, 0
+/work/Strata_B70/src/kernels/cuda/iq_kernels.cu, 167, 40, 541, 0
+/work/Strata_B70/src/kernels/cuda/fused_gdn.cu, 65, 9, 93, 0
+/work/Strata_B70/src/kernels/cuda/native_mmvq.cu, 242, 59, 1192, 0
+/work/Strata_B70/include/strata/core/verify.hpp, 1, 5, 232, 0
+/work/Strata_B70/src/kernels/elementwise_parity.cpp, 19, 42, 356, 0
+/work/Strata_B70/src/kernels/cuda/qsa.cu, 138, 21, 653, 0
+/work/Strata_B70/src/kernels/dequant_bf16_test.cpp, 1, 9, 99, 0
+/work/Strata_B70/src/kernels/cuda/native_rope.cu, 10, 4, 84, 0
+/work/Strata_B70/src/core/device.cu, 11, 13, 95, 0
+/work/Strata_B70/src/kernels/quantize_act_parity.cpp, 6, 21, 381, 0
+/work/Strata_B70/src/kernels/cuda/s_gemv.cu, 127, 18, 542, 0
+/work/Strata_B70/src/kernels/cuda/bf16_gemv.cu, 31, 6, 105, 0
+/work/Strata_B70/include/strata/core/session.hpp, 1, 7, 404, 0
+/work/Strata_B70/src/kernels/cuda/shared_expert.cu, 60, 10, 277, 0
+/work/Strata_B70/src/core/layer.cpp, 29, 27, 1215, 0
+/work/Strata_B70/src/kernels/dequant_s2_parity.cpp, 6, 11, 127, 0
+/work/Strata_B70/src/core/pinned.cu, 13, 8, 346, 0
+/work/Strata_B70/src/core/concurrent_main.cpp, 13, 22, 163, 0
+/work/Strata_B70/src/kernels/s_gemv_q8k_parity.cpp, 6, 24, 207, 0
+/work/Strata_B70/src/core/weights.cpp, 2, 15, 440, 0
+/work/Strata_B70/src/kernels/cuda/elementwise.cu, 87, 9, 225, 0
+/work/Strata_B70/src/kernels/kv_q4_parity.cpp, 6, 32, 309, 0
+/work/Strata_B70/src/kernels/ple_parity.cpp, 26, 108, 1051, 0
+/work/Strata_B70/src/kernels/qsa_select_bench.cpp, 5, 17, 118, 0
+/work/Strata_B70/src/core/native_dense.cpp, 4, 7, 172, 0
+/work/Strata_B70/src/kernels/cuda/s2_gemv_quads.cu, 25, 2, 67, 0
+/work/Strata_B70/include/strata/kernels/bf16_bits.hpp, 2, 0, 45, 0
+/work/Strata_B70/src/kernels/gdn_parity.cpp, 6, 45, 347, 0
+/work/Strata_B70/src/kernels/cuda/native_bf16.cu, 32, 7, 148, 0
+/work/Strata_B70/src/kernels/cuda/s2_gemv.cu, 13, 2, 49, 0
+/work/Strata_B70/src/kernels/cuda/kv_q8.cu, 43, 6, 79, 0
+/work/Strata_B70/include/strata/kernels/rope.hpp, 2, 0, 49, 0
+/work/Strata_B70/third_party/ggml/ggml-common.h, 4, 865, 1043, 0
+/work/Strata_B70/include/strata/core/on_device.hpp, 1, 2, 20, 0
+/work/Strata_B70/src/core/overlap_main.cpp, 16, 28, 126, 0
+/work/Strata_B70/include/strata/core/mtp.hpp, 1, 6, 142, 0
+/work/Strata_B70/src/kernels/cuda/native_gr_norm.cu, 15, 3, 85, 0
+/work/Strata_B70/src/kernels/gr_parity.cpp, 16, 80, 520, 0
+/work/Strata_B70/src/kernels/cuda/native_router.cu, 19, 4, 105, 0
+/work/Strata_B70/src/kernels/s2_gemv_parity.cpp, 6, 14, 154, 0
+/work/Strata_B70/src/kernels/rope_parity.cpp, 6, 15, 151, 0
+/work/Strata_B70/src/prefill/prefill.cpp, 60, 52, 1542, 0
+/work/Strata_B70/src/kernels/iq_parity.cpp, 2, 12, 74, 0
+/work/Strata_B70/src/kernels/cuda/qsa_prompt_attn.cu, 159, 26, 485, 0
+/work/Strata_B70/src/kernels/cuda/ple.cu, 73, 6, 288, 0
+/work/Strata_B70/src/kernels/s2_gemv_q8_parity.cpp, 6, 16, 160, 0
+/work/Strata_B70/src/kernels/cuda/cvec.cu, 25, 22, 120, 0
+/work/Strata_B70/src/kernels/cuda/dequant_s2.cu, 10, 2, 40, 0
+/work/Strata_B70/src/kernels/cuda/native_gdn.cu, 15, 2, 106, 0
+/work/Strata_B70/include/strata/kernels/f16_bits.hpp, 2, 0, 95, 0
+/work/Strata_B70/include/strata/core/remote_experts.hpp, 1, 1, 84, 0
+/work/Strata_B70/src/kernels/cuda/s2_expert_grouped.cu, 184, 21, 581, 0
+/work/Strata_B70/src/prefill/gemm.cu, 7, 25, 83, 0
+/work/Strata_B70/src/kernels/cuda/native_moe.cu, 16, 2, 69, 0
+/work/Strata_B70/src/program/generate.cpp, 48, 134, 4874, 0
+/work/Strata_B70/src/core/native_head.cpp, 6, 12, 134, 0
+/work/Strata_B70/src/prefill/kernels.cu, 211, 10, 401, 0
+/work/Strata_B70/src/kernels/cuda/s2_gemv_q8.cu, 20, 4, 73, 0
+
+Total migration time: 179566.969000 ms
+
diff --git a/sycl/include/dpct/atomic.hpp b/sycl/include/dpct/atomic.hpp
new file mode 100644
index 00000000000..2bbecc94850
--- /dev/null
+++ b/sycl/include/dpct/atomic.hpp
@@ -0,0 +1,868 @@
+//==---- atomic.hpp -------------------------------*- C++ -*----------------==//
+//
+// Copyright (C) Intel Corporation
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// See https://llvm.org/LICENSE.txt for license information.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __DPCT_ATOMIC_HPP__
+#define __DPCT_ATOMIC_HPP__
+
+#include
+
+namespace dpct {
+
+/// Atomically add the value operand to the value at the addr and assign the
+/// result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to add to the value at \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_add(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_add(operand);
+}
+
+template
+inline T1 atomic_fetch_add(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_add(operand);
+}
+
+template
+inline sycl::half2 atomic_fetch_add(sycl::half2 *addr, sycl::half2 operand) {
+ auto atm = sycl::atomic_ref(
+ *reinterpret_cast(addr));
+
+ union {
+ unsigned i;
+ sycl::half2 h;
+ } old{0}, output{0};
+
+ while (true) {
+ old.i = atm.load();
+ output.h = old.h + operand;
+ if (atm.compare_exchange_strong(old.i, output.i))
+ break;
+ }
+
+ return output.h;
+}
+
+/// Atomically add the value operand to the value at the addr and assign the
+/// result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to add to the value at \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_add(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_add(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_add(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_add(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_add(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_add(addr, operand, memoryOrder);
+}
+
+/// Atomically subtract the value operand from the value at the addr and assign
+/// the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to subtract from the value at \p addr
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_sub(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_sub(operand);
+}
+
+template
+inline T1 atomic_fetch_sub(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_sub(operand);
+}
+
+/// Atomically subtract the value operand from the value at the addr and assign
+/// the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to subtract from the value at \p addr
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_sub(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_sub(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_sub(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_sub(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_sub(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_sub(addr, operand, memoryOrder);
+}
+
+/// Atomically perform a bitwise AND between the value operand and the value at the addr
+/// and assign the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to use in bitwise AND operation with the value at the \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_and(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_and(operand);
+}
+
+template
+inline T1 atomic_fetch_and(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_and(operand);
+}
+
+/// Atomically perform a bitwise AND between the value operand and the value at the addr
+/// and assign the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to use in bitwise AND operation with the value at the \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_and(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_and(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_and(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_and(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_and(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_and(addr, operand, memoryOrder);
+}
+
+/// Atomically or the value at the addr with the value operand, and assign
+/// the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to use in bitwise OR operation with the value at the \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_or(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_or(operand);
+}
+
+template
+inline T1 atomic_fetch_or(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_or(operand);
+}
+
+/// Atomically or the value at the addr with the value operand, and assign
+/// the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to use in bitwise OR operation with the value at the \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_or(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_or(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_or(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_or(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_or(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_or(addr, operand, memoryOrder);
+}
+
+/// Atomically xor the value at the addr with the value operand, and assign
+/// the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to use in bitwise XOR operation with the value at the \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_xor(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_xor(operand);
+}
+
+template
+inline T1 atomic_fetch_xor(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_xor(operand);
+}
+
+/// Atomically xor the value at the addr with the value operand, and assign
+/// the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to use in bitwise XOR operation with the value at the \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_xor(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_xor(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_xor(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_xor(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_xor(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_xor(addr, operand, memoryOrder);
+}
+
+/// Atomically calculate the minimum of the value at addr and the value operand
+/// and assign the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_min(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_min(operand);
+}
+
+template
+inline T1 atomic_fetch_min(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_min(operand);
+}
+
+/// Atomically calculate the minimum of the value at addr and the value operand
+/// and assign the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_min(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_min(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_min(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_min(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_min(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_min(addr, operand, memoryOrder);
+}
+
+/// Atomically calculate the maximum of the value at addr and the value operand
+/// and assign the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_max(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_max(operand);
+}
+
+template
+inline T1 atomic_fetch_max(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.fetch_max(operand);
+}
+
+/// Atomically calculate the maximum of the value at addr and the value operand
+/// and assign the result to the value at addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_fetch_max(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_max(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_max(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_max(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_fetch_max(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_fetch_max(addr, operand, memoryOrder);
+}
+
+/// Atomically set \p operand to the value stored in \p addr, if old value stored in
+/// \p addr is equal to zero or greater than \p operand, else decrease the value stored
+/// in \p addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The threshold value.
+/// \param memoryOrder The memory ordering used.
+/// \returns The old value stored in \p addr.
+template
+inline unsigned int atomic_fetch_compare_dec(unsigned int *addr,
+ unsigned int operand) {
+ auto atm = sycl::atomic_ref(addr[0]);
+ unsigned int old;
+
+ while (true) {
+ old = atm.load();
+ if (old == 0 || old > operand) {
+ if (atm.compare_exchange_strong(old, operand))
+ break;
+ } else if (atm.compare_exchange_strong(old, old - 1))
+ break;
+ }
+
+ return old;
+}
+
+/// Atomically increment the value stored in \p addr if old value stored in \p
+/// addr is less than \p operand, else set 0 to the value stored in \p addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The threshold value.
+/// \param memoryOrder The memory ordering used.
+/// \returns The old value stored in \p addr.
+template
+inline unsigned int atomic_fetch_compare_inc(unsigned int *addr,
+ unsigned int operand) {
+ auto atm = sycl::atomic_ref(addr[0]);
+ unsigned int old;
+ while (true) {
+ old = atm.load();
+ if (old >= operand) {
+ if (atm.compare_exchange_strong(old, 0))
+ break;
+ } else if (atm.compare_exchange_strong(old, old + 1))
+ break;
+ }
+ return old;
+}
+
+/// Atomically increment the value stored in \p addr if old value stored in \p
+/// addr is less than \p operand, else set 0 to the value stored in \p addr.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The threshold value.
+/// \param memoryOrder The memory ordering used.
+/// \returns The old value stored in \p addr.
+template
+inline unsigned int
+atomic_fetch_compare_inc(unsigned int *addr, unsigned int operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_fetch_compare_inc(addr,
+ operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_fetch_compare_inc(addr,
+ operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_fetch_compare_inc(addr,
+ operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+/// Atomically exchange the value at the address addr with the value operand.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to be exchanged with the value pointed by \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_exchange(T *addr, T operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.exchange(operand);
+}
+
+template
+inline T1 atomic_exchange(T1 *addr, T2 operand) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ return atm.exchange(operand);
+}
+
+/// Atomically exchange the value at the address addr with the value operand.
+/// \param [in, out] addr The pointer to the data.
+/// \param operand The value to be exchanged with the value pointed by \p addr.
+/// \param memoryOrder The memory ordering used.
+/// \returns The value at the \p addr before the call.
+template
+inline T atomic_exchange(T *addr, T operand,
+ sycl::memory_order memoryOrder) {
+ switch (memoryOrder) {
+ case sycl::memory_order::relaxed:
+ return atomic_exchange(addr, operand);
+ case sycl::memory_order::acq_rel:
+ return atomic_exchange(addr, operand);
+ case sycl::memory_order::seq_cst:
+ return atomic_exchange(addr, operand);
+ default:
+ assert(false && "Invalid memory_order for atomics. Valid memory_order for "
+ "atomics are: sycl::memory_order::relaxed, "
+ "sycl::memory_order::acq_rel, sycl::memory_order::seq_cst!");
+ }
+}
+
+template
+inline T1 atomic_exchange(T1 *addr, T2 operand,
+ sycl::memory_order memoryOrder) {
+ atomic_exchange(addr, operand, memoryOrder);
+}
+
+/// Atomically compare the value at \p addr to the value expected and exchange
+/// with the value desired if the value at \p addr is equal to the value expected.
+/// Returns the value at the \p addr before the call.
+/// \param [in, out] addr Multi_ptr.
+/// \param expected The value to compare against the value at \p addr.
+/// \param desired The value to assign to \p addr if the value at \p addr is expected.
+/// \param success The memory ordering used when comparison succeeds.
+/// \param fail The memory ordering used when comparison fails.
+/// \returns The value at the \p addr before the call.
+template
+T atomic_compare_exchange_strong(
+ sycl::multi_ptr addr, T expected, T desired,
+ sycl::memory_order success = sycl::memory_order::relaxed,
+ sycl::memory_order fail = sycl::memory_order::relaxed) {
+ auto atm = sycl::atomic_ref(*addr);
+
+ atm.compare_exchange_strong(expected, desired, success, fail);
+ return expected;
+}
+
+template
+T1 atomic_compare_exchange_strong(
+ sycl::multi_ptr addr, T2 expected, T3 desired,
+ sycl::memory_order success = sycl::memory_order::relaxed,
+ sycl::memory_order fail = sycl::memory_order::relaxed) {
+ auto atm =
+ sycl::atomic_ref(*addr);
+ T1 expected_value = expected;
+ atm.compare_exchange_strong(expected_value, desired, success, fail);
+ return expected_value;
+}
+
+/// Atomically compare the value at \p addr to the value expected and exchange
+/// with the value desired if the value at \p addr is equal to the value expected.
+/// Returns the value at the \p addr before the call.
+/// \param [in] addr The pointer to the data.
+/// \param expected The value to compare against the value at \p addr.
+/// \param desired The value to assign to \p addr if the value at \p addr is expected.
+/// \param success The memory ordering used when comparison succeeds.
+/// \param fail The memory ordering used when comparison fails.
+/// \returns The value at the \p addr before the call.
+template
+T atomic_compare_exchange_strong(
+ T *addr, T expected, T desired,
+ sycl::memory_order success = sycl::memory_order::relaxed,
+ sycl::memory_order fail = sycl::memory_order::relaxed) {
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ atm.compare_exchange_strong(expected, desired, success, fail);
+ return expected;
+}
+
+template
+T1 atomic_compare_exchange_strong(
+ T1 *addr, T2 expected, T3 desired,
+ sycl::memory_order success = sycl::memory_order::relaxed,
+ sycl::memory_order fail = sycl::memory_order::relaxed) {
+ T1 expected_value = expected;
+ auto atm =
+ sycl::atomic_ref(addr[0]);
+ atm.compare_exchange_strong(expected_value, desired, success, fail);
+ return expected_value;
+}
+
+/// Atomic extension to implement standard APIs in std::atomic
+namespace detail{
+template struct IsValidAtomicType {
+ static constexpr bool value =
+ (std::is_same::value || std::is_same::value ||
+ std::is_same::value || std::is_same::value ||
+ std::is_same::value ||
+ std::is_same::value ||
+ std::is_same::value || std::is_same::value ||
+ std::is_pointer::value);
+};
+} // namespace detail
+
+template
+class atomic {
+ static_assert(
+ detail::IsValidAtomicType::value,
+ "Invalid atomic type. Valid types are int, unsigned int, long, "
+ "unsigned long, long long, unsigned long long, float, double "
+ "and pointer types");
+ T __d;
+
+public:
+ /// default memory synchronization order
+ static constexpr sycl::memory_order default_read_order =
+ sycl::atomic_ref::default_read_order;
+ static constexpr sycl::memory_order default_write_order =
+ sycl::atomic_ref::default_write_order;
+ static constexpr sycl::memory_scope default_scope = DefaultScope;
+ static constexpr sycl::memory_order default_read_modify_write_order =
+ DefaultOrder;
+
+ /// Default constructor.
+ constexpr atomic() noexcept = default;
+ /// Constructor with initialize value.
+ constexpr atomic(T d) noexcept : __d(d){};
+
+ /// atomically replaces the value of the referenced object with a non-atomic argument
+ /// \param operand The value to replace the pointed value.
+ /// \param memoryOrder The memory ordering used.
+ /// \param memoryScope The memory scope used.
+ void store(T operand, sycl::memory_order memoryOrder = default_write_order,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+ sycl::atomic_ref atm(__d);
+ atm.store(operand, memoryOrder, memoryScope);
+ }
+
+ /// atomically obtains the value of the referenced object
+ /// \param memoryOrder The memory ordering used.
+ /// \param memoryScope The memory scope used.
+ /// \returns The value of the referenced object
+ T load(sycl::memory_order memoryOrder = default_read_order,
+ sycl::memory_scope memoryScope = default_scope) const noexcept {
+ sycl::atomic_ref atm(
+ const_cast(__d));
+ return atm.load(memoryOrder, memoryScope);
+ }
+
+ /// atomically replaces the value of the referenced object and obtains the value held previously
+ /// \param operand The value to replace the pointed value.
+ /// \param memoryOrder The memory ordering used.
+ /// \param memoryScope The memory scope used.
+ /// \returns The value of the referenced object before the call.
+ T exchange(T operand,
+ sycl::memory_order memoryOrder = default_read_modify_write_order,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+
+ sycl::atomic_ref atm(__d);
+ return atm.exchange(operand, memoryOrder, memoryScope);
+ }
+
+ /// atomically compares the value of the referenced object with non-atomic argument
+ /// and performs atomic exchange if equal or atomic load if not
+ /// \param expected The value expected to be found in the object referenced by the atomic_ref object
+ /// \param desired The value to store in the referenced object if it is as expected
+ /// \param success The memory models for the read-modify-write
+ /// \param failure The memory models for load operations
+ /// \param memoryScope The memory scope used.
+ /// \returns true if the referenced object was successfully changed, false otherwise.
+ bool compare_exchange_weak(
+ T &expected, T desired,
+ sycl::memory_order success, sycl::memory_order failure,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+ sycl::atomic_ref atm(__d);
+ return atm.compare_exchange_weak(expected, desired, success, failure, memoryScope);
+ }
+ /// \param expected The value expected to be found in the object referenced by the atomic_ref object
+ /// \param desired The value to store in the referenced object if it is as expected
+ /// \param memoryOrder The memory synchronization ordering for operations
+ /// \param memoryScope The memory scope used.
+ /// \returns true if the referenced object was successfully changed, false otherwise.
+ bool compare_exchange_weak(T &expected, T desired,
+ sycl::memory_order memoryOrder = default_read_modify_write_order,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+ sycl::atomic_ref atm(__d);
+ return atm.compare_exchange_weak(expected, desired, memoryOrder, memoryScope);
+ }
+
+ /// atomically compares the value of the referenced object with non-atomic argument
+ /// and performs atomic exchange if equal or atomic load if not
+ /// \param expected The value expected to be found in the object referenced by the atomic_ref object
+ /// \param desired The value to store in the referenced object if it is as expected
+ /// \param success The memory models for the read-modify-write
+ /// \param failure The memory models for load operations
+ /// \param memoryScope The memory scope used.
+ /// \returns true if the referenced object was successfully changed, false otherwise.
+ bool compare_exchange_strong(
+ T &expected, T desired,
+ sycl::memory_order success, sycl::memory_order failure,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+
+ sycl::atomic_ref atm(__d);
+ return atm.compare_exchange_strong(expected, desired, success, failure, memoryScope);
+ }
+ /// \param expected The value expected to be found in the object referenced by the atomic_ref object
+ /// \param desired The value to store in the referenced object if it is as expected
+ /// \param memoryOrder The memory synchronization ordering for operations
+ /// \param memoryScope The memory scope used.
+ /// \returns true if the referenced object was successfully changed, false otherwise.
+ bool compare_exchange_strong(T &expected, T desired,
+ sycl::memory_order memoryOrder = default_read_modify_write_order,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+ sycl::atomic_ref atm(__d);
+ return atm.compare_exchange_strong(expected, desired, memoryOrder, memoryScope);
+ }
+
+ /// atomically adds the argument to the value stored in the atomic object and obtains the value held previously
+ /// \param operand The other argument of arithmetic addition
+ /// \param memoryOrder The memory ordering used.
+ /// \param memoryScope The memory scope used.
+ /// \returns The value of the referenced object before the call.
+ T fetch_add(T operand,
+ sycl::memory_order memoryOrder = default_read_modify_write_order,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+
+ sycl::atomic_ref atm(__d);
+ return atm.fetch_add(operand, memoryOrder, memoryScope);
+ }
+
+ /// atomically subtracts the argument from the value stored in the atomic object and obtains the value held previously
+ /// \param operand The other argument of arithmetic subtraction
+ /// \param memoryOrder The memory ordering used.
+ /// \param memoryScope The memory scope used.
+ /// \returns The value of the referenced object before the call.
+ T fetch_sub(T operand,
+ sycl::memory_order memoryOrder = default_read_modify_write_order,
+ sycl::memory_scope memoryScope = default_scope) noexcept {
+
+ sycl::atomic_ref atm(__d);
+ return atm.fetch_sub(operand, memoryOrder, memoryScope);
+ }
+};
+
+} // namespace dpct
+#endif // __DPCT_ATOMIC_HPP__
diff --git a/sycl/include/dpct/bindless_images.hpp b/sycl/include/dpct/bindless_images.hpp
new file mode 100644
index 00000000000..042321d328a
--- /dev/null
+++ b/sycl/include/dpct/bindless_images.hpp
@@ -0,0 +1,2061 @@
+//===--- bindless_images.hpp ----------------------*- C++ -*---------------===//
+//
+// Copyright (C) Intel Corporation
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// See https://llvm.org/LICENSE.txt for license information.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __DPCT_BINDLESS_IMAGE_HPP__
+#define __DPCT_BINDLESS_IMAGE_HPP__
+
+#ifdef _WIN32
+// DirectX headers
+#include
+#include
+#include
+#endif // _WIN32
+
+namespace dpct {
+namespace experimental {
+
+#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES
+
+/// The wrapper structure for external memory resource handle desc
+class external_mem_handle_desc {
+#ifdef _WIN32
+ void *handle;
+ const void *name;
+#else
+ int handle;
+#endif // _WIN32
+ sycl::ext::oneapi::experimental::external_mem_handle_type handle_type;
+ size_t size_in_bytes;
+ int flags;
+
+public:
+ external_mem_handle_desc() {
+#ifdef _WIN32
+ handle = nullptr;
+ name = nullptr;
+#else
+ handle = -1;
+#endif // _WIN32
+ handle_type =
+ sycl::ext::oneapi::experimental::external_mem_handle_type::opaque_fd;
+ size_in_bytes = 0;
+ flags = 0;
+ }
+
+#ifdef _WIN32
+ /// handle setter
+ /// \param [in] win32_handle The win32 NT handle of the external memory
+ /// resource
+ void set_win32_handle(void *win32_handle) { handle = win32_handle; }
+
+ /// name setter
+ /// \param [in] win32_obj_name The object of the external memory resource
+ void set_win32_obj_name(const void *win32_obj_name) { name = win32_obj_name; }
+#else
+ /// fd setter
+ /// \param [in] fd The file descriptor of the external memory resource
+ void set_fd_handle(int fd) { handle = fd; }
+#endif // _WIN32
+
+ /// handle_type setter
+ /// \param [in] type The type of the external memory resource handle
+ void set_handle_type(
+ sycl::ext::oneapi::experimental::external_mem_handle_type type) {
+ handle_type = type;
+ }
+
+ /// size setter
+ /// \param [in] size The size (in bytes) of the external memory resource
+ void set_res_size(unsigned long long size) { size_in_bytes = size; }
+
+ /// flags setter
+ /// \param [in] flags The flags used to handle importing of external memory
+ /// resource
+ void set_flags(int flags) { this->flags = flags; }
+
+#ifdef _WIN32
+ /// handle getter
+ /// \param [out] handle The win32 NT handle of the external memory resource
+ void *get_win32_handle() { return handle; }
+
+ /// name getter
+ /// \param [out] name The object of the external memory resource
+ const void *get_win32_obj_name() { return name; }
+#else
+ /// fd getter
+ /// \param [out] fd The file descriptor of the external memory resource
+ int get_fd_handle() { return handle; }
+#endif // _WIN32
+
+ /// handle_type getter
+ /// \param [out] handle_type The type of the external memory resource handle
+ sycl::ext::oneapi::experimental::external_mem_handle_type get_handle_type() {
+ return handle_type;
+ }
+
+ /// size getter
+ /// \param [out] size The size (in bytes) of the external memory resource
+ unsigned long long get_res_size() { return size_in_bytes; }
+
+ /// flags getter
+ /// \param [out] flags The flags used to handle importing of external memory
+ /// resource
+ int get_flags() { return flags; }
+};
+
+/// The wrapper structure for external image memory desc
+class external_mem_img_desc {
+ sycl::range<3> size;
+ image_channel channel;
+ unsigned int num_levels;
+ sycl::ext::oneapi::experimental::image_type type;
+
+public:
+ external_mem_img_desc() {
+ size = sycl::range<3>{0, 0, 0};
+ channel = image_channel::create();
+ num_levels = 1;
+ type = sycl::ext::oneapi::experimental::image_type::standard;
+ }
+
+ /// size setter
+ /// \param [in] img_dims The dimensions of the imported image resource
+ void set_size(sycl::range<3> img_dims) {
+ size = img_dims;
+
+ for (int i = 0; i < 3; i++) {
+ if (size[i] == 1) {
+ size[i] = 0;
+ }
+ }
+ }
+
+ /// image_channel setter
+ /// \param [in] img_ch The channel info of the imported image resource
+ void set_image_channel(image_channel img_ch) { channel = img_ch; }
+
+ /// num_levels setter
+ /// \param [in] numLevels The no. of levels in the imported image resource
+ void set_num_levels(unsigned int num_levels) {
+ this->num_levels = num_levels;
+
+ if (num_levels > 1) {
+ set_image_type(sycl::ext::oneapi::experimental::image_type::mipmap);
+ }
+ }
+
+ /// type setter
+ /// \param [in] img_type The type of the imported image resource
+ void set_image_type(sycl::ext::oneapi::experimental::image_type img_type) {
+ type = img_type;
+ }
+
+ /// size getter
+ /// \param [out] size The dimensions of the imported image resource
+ sycl::range<3> get_size() { return size; }
+
+ /// image_channel getter
+ /// \param [out] image_channel The channel info of the imported image resource
+ image_channel get_image_channel() { return channel; }
+
+ /// num_levels getter
+ /// \param [out] num_levels The no. of levels in the imported image resource
+ unsigned int get_num_levels() { return num_levels; }
+
+ /// type getter
+ /// \param [out] type The image type of the imported image resource
+ sycl::ext::oneapi::experimental::image_type get_image_type() { return type; }
+};
+
+/// The wrapper structure for external buffer memory desc
+class external_mem_buf_desc {
+ uint64_t size_in_bytes;
+ uint64_t offset;
+ unsigned int flags;
+
+public:
+ external_mem_buf_desc() {
+ size_in_bytes = 0;
+ offset = 0;
+ flags = 0;
+ }
+
+ /// size setter
+ /// \param [in] size The size (in bytes) of the external buffer resource
+ void set_res_size(uint64_t size) { size_in_bytes = size; }
+
+ /// offset setter
+ /// \param [in] offset The memory offset of the external buffer resource
+ void set_mem_offset(uint64_t offset) { this->offset = offset; }
+
+ /// flags setter
+ /// \param [in] flags The flags used to import external buffer resource
+ void set_flags(unsigned int flags) { this->flags = flags; }
+
+ /// size getter
+ /// \param [out] size_in_bytes The size (in bytes) of the external buffer res
+ uint64_t get_res_size() { return size_in_bytes; }
+
+ /// offset getter
+ /// \param [out] offset The memory offset of the external buffer resource
+ uint64_t get_mem_offset() { return offset; }
+
+ /// flags getter
+ /// \param [out] flags The flags used to import external buffer resource
+ unsigned int get_flags() { return flags; }
+};
+
+/// The wrapper class for external semaphore resource handle desc
+class external_sem_handle_desc {
+#ifdef _WIN32
+ void *handle;
+ const void *name;
+#else
+ int handle;
+#endif // _WIN32
+ sycl::ext::oneapi::experimental::external_semaphore_handle_type handle_type;
+ int flags;
+
+public:
+ external_sem_handle_desc() {
+#ifdef _WIN32
+ handle = nullptr;
+ name = nullptr;
+#else
+ handle = -1;
+#endif // _WIN32
+ handle_type = sycl::ext::oneapi::experimental::
+ external_semaphore_handle_type::opaque_fd;
+ flags = 0;
+ }
+
+#ifdef _WIN32
+ /// handle setter
+ /// \param [in] win32_handle The win32 NT handle of the external semaphore
+ /// resource
+ void set_win32_handle(void *win32_handle) { handle = win32_handle; }
+
+ /// name setter
+ /// \param [in] win32_obj_name The object of the external semaphore resource
+ void set_win32_obj_name(const void *win32_obj_name) { name = win32_obj_name; }
+#else
+ /// fd setter
+ /// \param [in] fd The file descriptor of the external semaphore resource
+ void set_fd_handle(int fd) { handle = fd; }
+#endif // _WIN32
+
+ /// handle_type setter
+ /// \param [in] type The type of the external semaphore resource handle
+ void set_handle_type(
+ sycl::ext::oneapi::experimental::external_semaphore_handle_type type) {
+ handle_type = type;
+ }
+
+ /// flags setter
+ /// \param [in] flags The flags used to handle importing of external semaphore
+ /// resource
+ void set_flags(int flags) { this->flags = flags; }
+
+#ifdef _WIN32
+ /// handle getter
+ /// \param [out] handle The win32 NT handle of the external semaphore resource
+ void *get_win32_handle() { return handle; }
+
+ /// name getter
+ /// \param [out] name The object of the external semaphore resource
+ const void *get_win32_obj_name() { return name; }
+#else
+ /// fd getter
+ /// \param [out] fd The file descriptor of the external semaphore resource
+ int get_fd_handle() { return handle; }
+#endif // _WIN32
+
+ /// handle_type getter
+ /// \param [out] handle_type The type of the external semaphore resource
+ /// handle
+ sycl::ext::oneapi::experimental::external_semaphore_handle_type
+ get_handle_type() {
+ return handle_type;
+ }
+
+ /// flags getter
+ /// \param [out] flags The flags used to handle importing of external
+ /// semaphore
+ int get_flags() { return flags; }
+};
+
+/// The wrapper structure for external semaphore signal/wait params
+class external_sem_params {
+ uint64_t value;
+
+public:
+ external_sem_params() { value = 0; }
+
+ /// value setter
+ /// \param [in] val The state value of the semaphore
+ void set_value(unsigned long long val) { value = val; }
+
+ /// value getter
+ /// \param [out] value The state value of the semaphore
+ unsigned long long get_value() { return value; }
+};
+
+/// The wrapper class of external semaphore
+class external_sem_wrapper {
+ sycl::ext::oneapi::experimental::external_semaphore ext_sem;
+ sycl::ext::oneapi::experimental::external_semaphore_handle_type
+ ext_sem_handle_type;
+
+public:
+ /// Create external semaphore wrapper.
+ /// \param [in] ext_sem The external semaphore resource object
+ /// \param [in] ext_sem_handle_type The external semaphore resource handle
+ /// type
+ external_sem_wrapper(
+ sycl::ext::oneapi::experimental::external_semaphore ext_sem,
+ sycl::ext::oneapi::experimental::external_semaphore_handle_type
+ ext_sem_handle_type) {
+ this->ext_sem = ext_sem;
+ this->ext_sem_handle_type = ext_sem_handle_type;
+ }
+
+ /// destory external semaphore resource
+ ~external_sem_wrapper() {
+ sycl::ext::oneapi::experimental::release_external_semaphore(
+ ext_sem, get_default_queue());
+ }
+
+ /// ext_sem getter
+ /// \param [out] ext_sem The external semaphore resource object
+ sycl::ext::oneapi::experimental::external_semaphore get() { return ext_sem; }
+
+ /// ext_sem_handle_type getter
+ /// \param [out] ext_sem_handle_type The external semaphore resource handle
+ /// type
+ sycl::ext::oneapi::experimental::external_semaphore_handle_type
+ get_handle_type() {
+ return ext_sem_handle_type;
+ }
+};
+
+/// The wrapper class of bindless image memory handle.
+class image_mem_wrapper {
+public:
+ /// Create bindless image memory wrapper.
+ /// \tparam dimensions The dimensions of image memory.
+ /// \param [in] channel The image channel used to create bindless image
+ /// \param [in] range The sizes of each dimension of bindless image memory.
+ /// memory.
+ template
+ image_mem_wrapper(image_channel channel, sycl::range range,
+ sycl::ext::oneapi::experimental::image_type type =
+ sycl::ext::oneapi::experimental::image_type::standard,
+ unsigned int num_levels = 1)
+ : _channel(channel) {
+ unsigned array_size = 1;
+ if (type == sycl::ext::oneapi::experimental::image_type::array) {
+ assert(dimensions == 3);
+ array_size = range[2];
+ range[2] = 0;
+ }
+ _desc = sycl::ext::oneapi::experimental::image_descriptor(
+ range, _channel.get_channel_num(), _channel.get_channel_type(), type,
+ num_levels, array_size);
+ auto q = get_default_queue();
+ _handle = alloc_image_mem(_desc, q);
+ init_mip_level_wrappers(q);
+ }
+ /// Create bindless image memory wrapper.
+ /// \param [in] channel The image channel used to create bindless image
+ /// \param [in] width The width of bindless image memory.
+ /// \param [in] height The height of bindless image memory.
+ /// \param [in] depth The depth of bindless image memory.
+ image_mem_wrapper(image_channel channel, size_t width, size_t height = 0,
+ size_t depth = 0)
+ : image_mem_wrapper(channel, {width, height, depth}) {}
+ /// Create bindless image memory wrapper.
+ /// \param [in] desc The image descriptor used to create bindless image
+ image_mem_wrapper(
+ const sycl::ext::oneapi::experimental::image_descriptor *desc)
+ : _desc(*desc) {
+ _channel.set_channel_type(desc->channel_type);
+ _channel.set_channel_num(desc->num_channels);
+ auto q = get_default_queue();
+ _handle = alloc_image_mem(_desc, q);
+ }
+ /// Create bindless image memory wrapper.
+ /// \param [in] desc The image descriptor of bindless image.
+ /// \param [in] img_mem_handle The memory handle of bindless image.
+ image_mem_wrapper(
+ sycl::ext::oneapi::experimental::image_descriptor desc,
+ sycl::ext::oneapi::experimental::image_mem_handle img_mem_handle)
+ : _desc(desc), _handle(img_mem_handle) {
+ _channel.set_channel_type(_desc.channel_type);
+#if (__SYCL_COMPILER_VERSION && __SYCL_COMPILER_VERSION >= 20240725)
+ _channel.set_channel_num(_desc.num_channels);
+#endif
+ auto q = get_default_queue();
+ init_mip_level_wrappers(q);
+ }
+ /// Create bindless image memory wrapper.
+ /// \param [in] extMem SYCL external memory resource object.
+ /// \param [in] ext_img_mem_desc The img desc of the imported image resource.
+ image_mem_wrapper(sycl::ext::oneapi::experimental::external_mem extMem,
+ external_mem_img_desc *ext_img_mem_desc) {
+ auto q = get_default_queue();
+ unsigned int array_size = 1;
+
+ _channel = ext_img_mem_desc->get_image_channel();
+ _desc = sycl::ext::oneapi::experimental::image_descriptor(
+ ext_img_mem_desc->get_size(), _channel.get_channel_num(),
+ _channel.get_channel_type(), ext_img_mem_desc->get_image_type(),
+ ext_img_mem_desc->get_num_levels(), array_size);
+ _handle = sycl::ext::oneapi::experimental::map_external_image_memory(
+ extMem, _desc, q);
+
+ init_mip_level_wrappers(q);
+ }
+ /// Create bindless image memory wrapper.
+ /// \param [in] desc The image descriptor of bindless image.
+ /// \param [in] num_levels Number of mipmap levels to allocate
+ image_mem_wrapper(sycl::ext::oneapi::experimental::image_descriptor desc,
+ unsigned int num_levels)
+ : _desc(desc) {
+ _desc.type = sycl::ext::oneapi::experimental::image_type::mipmap;
+ _desc.num_levels = num_levels;
+ auto q = get_default_queue();
+ _handle = alloc_image_mem(_desc, q);
+ init_mip_level_wrappers(q);
+ }
+
+ image_mem_wrapper(const image_mem_wrapper &) = delete;
+ image_mem_wrapper &operator=(const image_mem_wrapper &) = delete;
+ /// Destroy bindless image memory wrapper.
+ ~image_mem_wrapper() {
+ if (_sub_wrappers) {
+ std::destroy_n(_sub_wrappers, _desc.num_levels);
+ std::free(_sub_wrappers);
+ }
+ free_image_mem(_handle, _desc.type, get_default_queue());
+ }
+ /// Get the image channel of the bindless image memory.
+ /// \returns The image channel of bindless image memory.
+ image_channel get_channel() const noexcept { return _channel; }
+ /// Get the sizes of each dimension of the bindless image memory.
+ /// \returns The sizes of each dimension of bindless image memory.
+ sycl::range<3> get_range() const {
+ return {_desc.width, _desc.height, _desc.depth};
+ }
+ /// Get the image descriptor of the bindless image memory.
+ /// \returns The image descriptor of bindless image memory.
+ const sycl::ext::oneapi::experimental::image_descriptor &
+ get_desc() const noexcept {
+ return _desc;
+ }
+ /// Get the image handle of the bindless image memory.
+ /// \returns The image handle of bindless image memory.
+ sycl::ext::oneapi::experimental::image_mem_handle
+ get_handle() const noexcept {
+ return _handle;
+ }
+ /// Get the image mip level of the bindless image memory.
+ /// \returns The image mip level of the bindless image memory.
+ image_mem_wrapper *get_mip_level(unsigned int level) {
+ if (_desc.type == sycl::ext::oneapi::experimental::image_type::standard) {
+ assert(level == 0);
+ return this;
+ }
+
+ assert(_desc.type == sycl::ext::oneapi::experimental::image_type::mipmap);
+ return _sub_wrappers + level;
+ }
+
+ sycl::ext::oneapi::experimental::image_type get_image_type(void) {
+ return _desc.type;
+ }
+
+private:
+ image_mem_wrapper(
+ const image_channel &channel,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc,
+ const sycl::ext::oneapi::experimental::image_mem_handle &handle)
+ : _channel(channel), _desc(desc), _handle(handle) {}
+
+ /// Initialize image mem wrappers for all the levels of mipmap image
+ /// \param [in] q The Queue to be used to query the mip levels
+ void init_mip_level_wrappers(sycl::queue q) {
+ auto num_levels = _desc.num_levels;
+ if (_desc.type == sycl::ext::oneapi::experimental::image_type::mipmap) {
+ assert(num_levels > 1);
+ _sub_wrappers = (image_mem_wrapper *)std::malloc(
+ sizeof(image_mem_wrapper) * num_levels);
+ for (unsigned i = 0; i < num_levels; ++i)
+ new (_sub_wrappers + i) image_mem_wrapper(
+ _channel, _desc.get_mip_level_desc(i),
+ sycl::ext::oneapi::experimental::get_mip_level_mem_handle(
+ _handle, i, q.get_device(), q.get_context()));
+ }
+ }
+
+ image_channel _channel;
+ sycl::ext::oneapi::experimental::image_descriptor _desc;
+ sycl::ext::oneapi::experimental::image_mem_handle _handle;
+ image_mem_wrapper *_sub_wrappers{nullptr};
+};
+
+/// The base wrapper class of external memory resource handle.
+class external_mem_wrapper_base {
+public:
+ /// Cleans up the external mem wrapper by releasing the SYCL external
+ /// memory handle
+ virtual ~external_mem_wrapper_base() {}
+
+ external_mem_wrapper_base(const external_mem_wrapper_base &) = delete;
+
+ external_mem_wrapper_base &
+ operator=(const external_mem_wrapper_base &) = delete;
+
+ const sycl::ext::oneapi::experimental::external_mem
+ get_external_mem() const noexcept {
+ return _res_external_mem;
+ }
+
+ /// Map the buf/texture resource memory and get its pointer/mem handle.
+ /// \param [in] q The queue used to map the resource with.
+ void map_resource(sycl::queue q) {
+ if (_res_is_buffer) {
+ if (_res_buf_ptr) {
+ throw std::runtime_error("Resource is already mapped! "
+ "Resource cannot be mapped more than once.");
+ }
+
+#if (__SYCL_COMPILER_VERSION && __SYCL_COMPILER_VERSION >= 20240823)
+ _res_buf_ptr =
+ sycl::ext::oneapi::experimental::map_external_linear_memory(
+ get_external_mem(), 0, _res_size_bytes, q);
+#else
+ throw std::runtime_error(
+ "Mapping a buffer resoure is not yet supported!");
+#endif
+ } else {
+ if (_res_img_mem_wrapper_ptr) {
+ throw std::runtime_error("Resource is already mapped! "
+ "Resource cannot be mapped more than once.");
+ }
+
+ auto img_mem_handle =
+ sycl::ext::oneapi::experimental::map_external_image_memory(
+ get_external_mem(), _res_img_desc, q);
+
+ _res_img_mem_wrapper_ptr =
+ new image_mem_wrapper(_res_img_desc, img_mem_handle);
+ }
+ }
+
+ /// Gets the pointer and size of the mapped resource
+ /// \param [out] ptr The pointer to the resource memory.
+ /// \param [out] size The size of the resource memory.
+ /// and resource size.
+ void get_mapped_pointer(void **ptr, size_t *size) {
+ if (!_res_is_buffer) {
+ throw std::runtime_error(
+ "Non buffer resouce cannot be accessed as a pointer!");
+ }
+
+ *ptr = _res_buf_ptr;
+ if (!*ptr) {
+ throw std::runtime_error(
+ "Resource is not mapped! "
+ "Resource should be mapped before accessing its memory");
+ }
+
+ *size = _res_size_bytes;
+ }
+
+ /// Returns mapped memory as image_mem_wrapper*
+ image_mem_wrapper *get_mapped_mipmapped_array() {
+ if (_res_is_buffer) {
+ throw std::runtime_error(
+ "Buffer resource cannot be accessed as an array!");
+ }
+ if (!_res_img_mem_wrapper_ptr) {
+ throw std::runtime_error(
+ "Resource is not mapped! "
+ "Resource should be mapped before accessing its memory.");
+ }
+
+ return _res_img_mem_wrapper_ptr;
+ }
+
+ /// Returns mapped subresource memory as image_mem_wrapper*
+ /// \param [in] array_index The layer of the image array to be returned.
+ /// \param [in] mip_level The level of the mipmap image to be returned.
+ image_mem_wrapper *get_sub_resource_mapped_array(unsigned int array_index,
+ unsigned int mip_level) {
+ if (_res_is_buffer) {
+ throw std::runtime_error(
+ "Buffer resouce cannot be accessed as an array!");
+ }
+ if (array_index != 0) {
+ throw std::runtime_error("Retrieving a particular layer from image array "
+ "is not yet supported!");
+ }
+ if (array_index >= _res_img_desc.array_size) {
+ throw std::runtime_error(
+ "Array index requested exceeds total array size in resource!");
+ }
+ if (mip_level >= _res_img_desc.num_levels) {
+ throw std::runtime_error(
+ "Mip level requested exceeds total num of levels in resource!");
+ }
+ if (!_res_img_mem_wrapper_ptr) {
+ throw std::runtime_error(
+ "Resouce is not mapped! "
+ "Resource should be mapped before accessing its memory");
+ }
+
+ return (mip_level == 0)
+ ? _res_img_mem_wrapper_ptr
+ : _res_img_mem_wrapper_ptr->get_mip_level(mip_level);
+ }
+
+ /// Unmap the buf/texture resource memory and get its pointer/mem handle.
+ /// \param [in] q The queue used to unmap the resource with.
+ void unmap_resource(sycl::queue q) {
+ if (_res_is_buffer) {
+ if (!_res_buf_ptr) {
+ throw std::runtime_error("Resource is not mapped! "
+ "Resource should be mapped before unmapping.");
+ }
+
+ _res_buf_ptr = nullptr;
+
+ throw std::runtime_error(
+ "Unmapping a buffer resoure is not yet supported!");
+ } else {
+ if (!_res_img_mem_wrapper_ptr) {
+ throw std::runtime_error("Resource is not mapped! "
+ "Resource should be mapped before unmapping.");
+ }
+
+ delete _res_img_mem_wrapper_ptr;
+ _res_img_mem_wrapper_ptr = nullptr;
+ }
+ }
+
+protected:
+ /// Init external mem handle for a given resource.
+ external_mem_wrapper_base(
+ sycl::ext::oneapi::experimental::external_mem external_mem,
+ unsigned reg_flags)
+ : _res_external_mem(external_mem), _res_reg_flags(reg_flags) {}
+
+ unsigned _res_reg_flags = 0;
+ bool _res_is_buffer = false;
+ size_t _res_size_bytes = 0;
+
+ sycl::ext::oneapi::experimental::image_descriptor _res_img_desc;
+ sycl::ext::oneapi::experimental::external_mem _res_external_mem;
+
+private:
+ void *_res_buf_ptr = nullptr;
+ image_mem_wrapper *_res_img_mem_wrapper_ptr = nullptr;
+};
+
+#ifdef _WIN32
+class external_mem_wrapper : public external_mem_wrapper_base {
+public:
+ /// Creates external memory resource wrapper for a given DX11 resource.
+ /// \param [in] d3d11_res Pointer to the ID3D11Resource to be wrapped.
+ /// \param [in] reg_flags Registration flags used for resource registration.
+ external_mem_wrapper(ID3D11Resource *d3d11_res, unsigned reg_flags)
+ : _res_D3D11(d3d11_res),
+ external_mem_wrapper_base(
+ sycl::ext::oneapi::experimental::external_mem(), reg_flags) {
+ // Check if the resource is already registered to prevent double
+ // registration.
+ if (_res_D3D11_set.find(_res_D3D11) != _res_D3D11_set.end()) {
+ throw std::runtime_error("This D3D11 resource is already registered!");
+ }
+
+ // Query the information from D3D11 resource object
+ query_res_info(_res_D3D11);
+
+ // Create a shared handle for the DX11 resource to enable interoperability.
+ _res_win_nt_handle = create_shared_handle(_res_D3D11);
+
+ // Prepare the external memory resource descriptor for SYCL interop.
+ sycl::ext::oneapi::experimental::resource_win32_handle ext_mem_win_handle{
+ _res_win_nt_handle};
+
+ sycl::ext::oneapi::experimental::external_mem_descriptor<
+ sycl::ext::oneapi::experimental::resource_win32_handle>
+ ext_mem_desc{ext_mem_win_handle,
+ sycl::ext::oneapi::experimental::external_mem_handle_type::
+ win32_nt_dx12_resource,
+ _res_size_bytes};
+
+ // Import resource's external memory into SYCL for use with SYCL APIs.
+ _res_external_mem = sycl::ext::oneapi::experimental::external_mem(
+ sycl::ext::oneapi::experimental::import_external_memory(
+ ext_mem_desc, get_default_queue()));
+
+ // Register the resource to track its uniqueness.
+ _res_D3D11_set.insert(_res_D3D11);
+ }
+
+ /// Cleans up the external mem wrapper by releasing the SYCL external
+ /// memory handle and the shared handle created for the DX11 resource.
+ virtual ~external_mem_wrapper() {
+ // Release the mutex when done
+ _res_keyed_mutex->ReleaseSync(0);
+
+ // Release the keyed mutex interface
+ _res_keyed_mutex->Release();
+
+ // Close the shared handle to the DX11 resource.
+ CloseHandle(_res_win_nt_handle);
+
+ // Release the resource's external memory associated with the SYCL external
+ // mem handle.
+ sycl::ext::oneapi::experimental::release_external_memory(
+ get_external_mem(), get_default_queue());
+
+ // Remove the resource from the tracking set.
+ if (_res_D3D11_set.find(_res_D3D11) != _res_D3D11_set.end()) {
+ _res_D3D11_set.erase(_res_D3D11);
+ }
+ }
+
+private:
+ inline static std::set _res_D3D11_set;
+
+ ID3D11Resource *_res_D3D11 = nullptr;
+ IDXGIKeyedMutex *_res_keyed_mutex = nullptr;
+ HANDLE _res_win_nt_handle = nullptr;
+
+ /// Helper function to query the properties of DX11 resource
+ void query_res_info(ID3D11Resource *resource) {
+ unsigned int res_width = 0;
+ unsigned int res_height = 0;
+ unsigned int res_depth = 0;
+ unsigned int res_num_levels = 1;
+ unsigned int res_arr_size = 1;
+
+ image_channel channel;
+
+ // Get the dimension info of DX11 resource
+ D3D11_RESOURCE_DIMENSION dimension;
+ resource->GetType(&dimension);
+
+ switch (dimension) {
+ case D3D11_RESOURCE_DIMENSION_BUFFER: {
+ _res_is_buffer = true;
+
+ D3D11_BUFFER_DESC desc;
+ ((ID3D11Buffer *)resource)->GetDesc(&desc);
+
+ _res_size_bytes = desc.ByteWidth;
+ break;
+ }
+ case D3D11_RESOURCE_DIMENSION_TEXTURE1D: {
+ D3D11_TEXTURE1D_DESC desc;
+ ((ID3D11Texture1D *)resource)->GetDesc(&desc);
+
+ res_width = desc.Width;
+ res_arr_size = desc.ArraySize;
+ res_num_levels = (desc.MipLevels ? desc.MipLevels : 1);
+
+ channel = get_img_ch_info(desc.Format);
+ _res_size_bytes = res_width;
+ break;
+ }
+ case D3D11_RESOURCE_DIMENSION_TEXTURE2D: {
+ D3D11_TEXTURE2D_DESC desc;
+ ((ID3D11Texture2D *)resource)->GetDesc(&desc);
+
+ res_width = desc.Width;
+ res_height = desc.Height;
+ res_arr_size = desc.ArraySize;
+ res_num_levels = (desc.MipLevels ? desc.MipLevels : 1);
+
+ channel = get_img_ch_info(desc.Format);
+ _res_size_bytes = res_width * res_height;
+ break;
+ }
+ case D3D11_RESOURCE_DIMENSION_TEXTURE3D: {
+ D3D11_TEXTURE3D_DESC desc;
+ ((ID3D11Texture3D *)resource)->GetDesc(&desc);
+
+ res_width = desc.Width;
+ res_height = desc.Height;
+ res_depth = desc.Depth;
+ res_num_levels = (desc.MipLevels ? desc.MipLevels : 1);
+
+ channel = get_img_ch_info(desc.Format);
+ _res_size_bytes = res_width * res_height * res_depth;
+ break;
+ }
+ default:
+ throw std::runtime_error("Unsupported DX11 resource type!");
+ break;
+ }
+
+ if (!_res_is_buffer) {
+ if (res_arr_size > 1) {
+ throw std::runtime_error(
+ "Images with multiple array levels are not supported. "
+ "Only standard and mipmap images are supported!");
+ }
+
+ // Calculate the size of a resource in bytes by multiplyting its dims
+ // with arr size, num levels & channel size
+ _res_size_bytes *=
+ res_arr_size * res_num_levels * channel.get_total_size();
+
+ sycl::ext::oneapi::experimental::image_type img_type =
+ (res_num_levels > 1)
+ ? sycl::ext::oneapi::experimental::image_type::mipmap
+ : sycl::ext::oneapi::experimental::image_type::standard;
+
+ _res_img_desc = sycl::ext::oneapi::experimental::image_descriptor{
+ sycl::range<3>{res_width, res_height, res_depth},
+ channel.get_channel_num(),
+ channel.get_channel_type(),
+ img_type,
+ res_num_levels,
+ res_arr_size};
+ }
+ }
+
+ /// Helper function to create shared handle for DXD11 resource
+ HANDLE create_shared_handle(ID3D11Resource *resource) {
+ HANDLE win_nt_handle = nullptr;
+
+ IDXGIResource1 *pDXGIResource1;
+ resource->QueryInterface(__uuidof(IDXGIResource1),
+ reinterpret_cast(&pDXGIResource1));
+
+ if (_res_is_buffer) {
+ pDXGIResource1->GetSharedHandle(&win_nt_handle);
+ } else {
+ pDXGIResource1->CreateSharedHandle(
+ nullptr, DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE,
+ nullptr, &win_nt_handle);
+ }
+
+ pDXGIResource1->Release();
+
+ // Query for the keyed mutex interface
+ resource->QueryInterface(__uuidof(IDXGIKeyedMutex),
+ reinterpret_cast(&_res_keyed_mutex));
+
+ // Use the keyed mutex to synchronize access to the shared resource
+ _res_keyed_mutex->AcquireSync(0, INFINITE);
+
+ return win_nt_handle;
+ }
+
+ /// Helper function to get the texture channel info of DX11 resource using
+ /// DXGI format
+ image_channel get_img_ch_info(DXGI_FORMAT format) {
+ image_channel channel;
+
+ switch (format) {
+ case DXGI_FORMAT_R16G16B16A16_FLOAT:
+ channel = image_channel::create();
+ break;
+ case DXGI_FORMAT_R8G8B8A8_UNORM:
+ channel = image_channel::create();
+ break;
+ case DXGI_FORMAT_R16G16_FLOAT:
+ channel = image_channel::create();
+ break;
+ case DXGI_FORMAT_R32_FLOAT:
+ channel = image_channel::create();
+ break;
+ default:
+ throw std::runtime_error("Unsupported DX11 resource format!");
+ break;
+ }
+
+ return channel;
+ }
+};
+#endif // _WIN32
+
+namespace detail {
+struct sampled_image_handle_compare {
+ template bool operator()(T L, T R) const {
+ return L.raw_handle < R.raw_handle;
+ }
+};
+
+inline std::pair &get_img_info_map(
+ const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
+ static std::map,
+ sampled_image_handle_compare>
+ img_info_map;
+ return img_info_map[handle];
+}
+
+inline std::pair &get_img_info_map(
+ const sycl::ext::oneapi::experimental::unsampled_image_handle handle) {
+ static std::map,
+ sampled_image_handle_compare>
+ img_info_map;
+ return img_info_map[handle];
+}
+
+static inline size_t
+get_ele_size(const sycl::ext::oneapi::experimental::image_descriptor &decs) {
+ size_t channel_size;
+ switch (decs.channel_type) {
+ case sycl::image_channel_type::signed_int8:
+ case sycl::image_channel_type::unsigned_int8:
+ channel_size = 1;
+ break;
+ case sycl::image_channel_type::fp16:
+ case sycl::image_channel_type::signed_int16:
+ case sycl::image_channel_type::unsigned_int16:
+ channel_size = 2;
+ break;
+ case sycl::image_channel_type::fp32:
+ case sycl::image_channel_type::signed_int32:
+ case sycl::image_channel_type::unsigned_int32:
+ channel_size = 4;
+ break;
+ default:
+ throw std::runtime_error("Unsupported channel_type in get_ele_size!");
+ break;
+ }
+ return decs.num_channels * channel_size;
+}
+
+static inline sycl::event
+dpct_memcpy(sycl::ext::oneapi::experimental::image_mem_handle src,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc_src,
+ size_t w_offset_src, size_t h_offset_src, void *dest, size_t p,
+ size_t w, size_t h, sycl::queue q) {
+ const auto ele_size = get_ele_size(desc_src);
+ const auto src_offset =
+ sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
+ const auto dest_offset = sycl::range<3>(0, 0, 0);
+ const auto dest_extend = sycl::range<3>(p / ele_size, 0, 0);
+ const auto copy_extend = sycl::range<3>(w / ele_size, h, 0);
+ return q.ext_oneapi_copy(src, src_offset, desc_src, dest, dest_offset,
+ dest_extend, copy_extend);
+}
+
+static inline std::vector dpct_memcpy_to_host(
+ sycl::ext::oneapi::experimental::image_mem_handle src,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc_src,
+ size_t w_offset_src, size_t h_offset_src, void *dest_host_ptr, size_t s,
+ sycl::queue q) {
+ std::vector event_list;
+ const auto ele_size = get_ele_size(desc_src);
+ const auto w = desc_src.width * ele_size;
+ size_t offset_dest = 0;
+ while (s - offset_dest > w - w_offset_src) {
+ const auto src_offset =
+ sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
+ const auto dest_offset = sycl::range<3>(offset_dest / ele_size, 0, 0);
+ const auto dest_extend = sycl::range<3>(0, 0, 0);
+ const auto copy_extend =
+ sycl::range<3>((w - w_offset_src) / ele_size, 1, 0);
+ event_list.push_back(q.ext_oneapi_copy(src, src_offset, desc_src,
+ dest_host_ptr, dest_offset,
+ dest_extend, copy_extend));
+ offset_dest += w - w_offset_src;
+ w_offset_src = 0;
+ ++h_offset_src;
+ }
+ if (s - offset_dest > 0) {
+ const auto src_offset =
+ sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
+ const auto dest_offset = sycl::range<3>(offset_dest / ele_size, 0, 0);
+ const auto dest_extend = sycl::range<3>(0, 0, 0);
+ const auto copy_extend = sycl::range<3>((s - offset_dest) / ele_size, 1, 0);
+ event_list.push_back(q.ext_oneapi_copy(src, src_offset, desc_src,
+ dest_host_ptr, dest_offset,
+ dest_extend, copy_extend));
+ }
+ return event_list;
+}
+
+static inline std::vector
+dpct_memcpy(sycl::ext::oneapi::experimental::image_mem_handle src,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc_src,
+ size_t w_offset_src, size_t h_offset_src, void *dest, size_t s,
+ sycl::queue q) {
+ if (dpct::detail::get_pointer_attribute(q, dest) ==
+ dpct::detail::pointer_access_attribute::device_only) {
+ std::vector event_list;
+ dpct::detail::host_buffer buf(s, q, event_list);
+ auto copy_events = dpct_memcpy_to_host(src, desc_src, w_offset_src,
+ h_offset_src, buf.get_ptr(), s, q);
+ event_list.push_back(dpct::detail::dpct_memcpy(
+ q, dest, buf.get_ptr(), s, memcpy_direction::host_to_device,
+ copy_events));
+ return event_list;
+ }
+ return dpct_memcpy_to_host(src, desc_src, w_offset_src, h_offset_src, dest, s,
+ q);
+}
+
+static inline sycl::event
+dpct_memcpy(const void *src,
+ sycl::ext::oneapi::experimental::image_mem_handle dest,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc_dest,
+ size_t w_offset_dest, size_t h_offset_dest, size_t p, size_t w,
+ size_t h, sycl::queue q) {
+ const auto ele_size = get_ele_size(desc_dest);
+ const auto src_offset = sycl::range<3>(0, 0, 0);
+ const auto src_extend = sycl::range<3>(p / ele_size, 0, 0);
+ const auto dest_offset =
+ sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
+ const auto copy_extend = sycl::range<3>(w / ele_size, h, 0);
+ // TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
+ return q.ext_oneapi_copy(const_cast(src), src_offset, src_extend,
+ dest, dest_offset, desc_dest, copy_extend);
+}
+
+static inline std::vector dpct_memcpy_from_host(
+ const void *src_host_ptr,
+ sycl::ext::oneapi::experimental::image_mem_handle dest,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc_dest,
+ size_t w_offset_dest, size_t h_offset_dest, size_t s, sycl::queue q) {
+ std::vector event_list;
+ const auto ele_size = get_ele_size(desc_dest);
+ const auto w = desc_dest.width * ele_size;
+ size_t offset_src = 0;
+ while (s - offset_src >= w - w_offset_dest) {
+ const auto src_offset = sycl::range<3>(offset_src / ele_size, 0, 0);
+ const auto src_extend = sycl::range<3>(0, 0, 0);
+ const auto dest_offset =
+ sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
+ const auto copy_extend =
+ sycl::range<3>((w - w_offset_dest) / ele_size, 1, 0);
+ // TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
+ event_list.push_back(q.ext_oneapi_copy(
+ const_cast(src_host_ptr), src_offset, src_extend, dest,
+ dest_offset, desc_dest, copy_extend));
+ offset_src += w - w_offset_dest;
+ w_offset_dest = 0;
+ ++h_offset_dest;
+ }
+ if (s - offset_src > 0) {
+ const auto src_offset = sycl::range<3>(offset_src / ele_size, 0, 0);
+ const auto src_extend = sycl::range<3>(0, 0, 0);
+ const auto dest_offset =
+ sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
+ const auto copy_extend = sycl::range<3>((s - offset_src) / ele_size, 1, 0);
+ // TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
+ event_list.push_back(q.ext_oneapi_copy(
+ const_cast(src_host_ptr), src_offset, src_extend, dest,
+ dest_offset, desc_dest, copy_extend));
+ }
+ return event_list;
+}
+
+static inline std::vector dpct_memcpy(
+ const void *src, sycl::ext::oneapi::experimental::image_mem_handle dest,
+ const sycl::ext::oneapi::experimental::image_descriptor &desc_dest,
+ size_t w_offset_dest, size_t h_offset_dest, size_t s, sycl::queue q) {
+ if (dpct::detail::get_pointer_attribute(q, src) ==
+ dpct::detail::pointer_access_attribute::device_only) {
+ std::vector event_list;
+ dpct::detail::host_buffer buf(s, q, event_list);
+ event_list.push_back(dpct::detail::dpct_memcpy(
+ q, buf.get_ptr(), src, s, memcpy_direction::device_to_host));
+ auto copy_events = dpct_memcpy_from_host(
+ buf.get_ptr(), dest, desc_dest, w_offset_dest, h_offset_dest, s, q);
+ event_list.insert(event_list.end(), copy_events.begin(), copy_events.end());
+ return event_list;
+ }
+ return dpct_memcpy_from_host(src, dest, desc_dest, w_offset_dest,
+ h_offset_dest, s, q);
+}
+
+static inline sycl::event
+dpct_memcpy(const image_mem_wrapper *src, const sycl::id<3> &src_id,
+ const size_t src_x_offest_byte, pitched_data &dest,
+ const sycl::id<3> &dest_id, const size_t dest_x_offest_byte,
+ const sycl::range<3> &size, const size_t copy_x_size_byte,
+ sycl::queue q) {
+ const auto ele_size = get_ele_size(src->get_desc());
+ const auto src_offset = sycl::range<3>(
+ src_x_offest_byte != 0 ? src_x_offest_byte / ele_size : src_id[0],
+ src_id[1], src_id[2]);
+ const auto dest_offset = sycl::range<3>(
+ dest_x_offest_byte != 0 ? dest_x_offest_byte / ele_size : dest_id[0],
+ dest_id[1], dest_id[2]);
+ const auto dest_extend =
+ sycl::range<3>(dest.get_pitch() / ele_size, dest.get_y(), 1);
+ const auto copy_extend = sycl::range<3>(
+ copy_x_size_byte != 0 ? copy_x_size_byte / ele_size : size[0], size[1],
+ size[2]);
+ return q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
+ dest.get_data_ptr(), dest_offset, dest_extend,
+ copy_extend);
+}
+
+static inline sycl::event
+dpct_memcpy(pitched_data src, const sycl::id<3> &src_id,
+ const size_t src_x_offest_byte, image_mem_wrapper *dest,
+ const sycl::id<3> &dest_id, const size_t dest_x_offest_byte,
+ const sycl::range<3> &size, const size_t copy_x_size_byte,
+ sycl::queue q) {
+ const auto ele_size = get_ele_size(dest->get_desc());
+ const auto src_offset = sycl::range<3>(
+ src_x_offest_byte != 0 ? src_x_offest_byte / ele_size : src_id[0],
+ src_id[1], src_id[2]);
+ const auto src_extend =
+ sycl::range<3>(src.get_pitch() / ele_size, src.get_y(), 1);
+ const auto dest_offset = sycl::range<3>(
+ dest_x_offest_byte != 0 ? dest_x_offest_byte / ele_size : dest_id[0],
+ dest_id[1], dest_id[2]);
+ const auto copy_extend = sycl::range<3>(
+ copy_x_size_byte != 0 ? copy_x_size_byte / ele_size : size[0], size[1],
+ size[2]);
+ return q.ext_oneapi_copy(src.get_data_ptr(), src_offset, src_extend,
+ dest->get_handle(), dest_offset, dest->get_desc(),
+ copy_extend);
+}
+
+static inline sycl::event
+dpct_memcpy(const image_mem_wrapper *src, const sycl::id<3> &src_id,
+ const size_t src_x_offest_byte, image_mem_wrapper *dest,
+ const sycl::id<3> &dest_id, const size_t dest_x_offest_byte,
+ const sycl::range<3> &size, const size_t copy_x_size_byte,
+ sycl::queue q) {
+ // TODO: Need change logic when sycl support image_mem to image_mem copy.
+ const auto from_ele_size = get_ele_size(src->get_desc());
+ const auto to_ele_size = get_ele_size(dest->get_desc());
+ const auto copy_extend = sycl::range<3>(
+ copy_x_size_byte != 0 ? copy_x_size_byte / from_ele_size : size[0],
+ size[1], size[2]);
+ std::vector event_list;
+ dpct::detail::host_buffer buf(
+ copy_extend.size() * std::max(from_ele_size, to_ele_size), q, event_list);
+ auto to = pitched_data(buf.get_ptr(), copy_extend[0] * from_ele_size,
+ copy_extend[0], copy_extend[1]);
+ dpct_memcpy(src, src_id, src_x_offest_byte, to, sycl::id<3>(0, 0, 0), 0,
+ copy_extend, 0, q);
+ auto from = pitched_data(buf.get_ptr(), copy_extend[0] * to_ele_size,
+ copy_extend[0], copy_extend[1]);
+ event_list.push_back(dpct_memcpy(from, sycl::id<3>(0, 0, 0), 0, dest, dest_id,
+ dest_x_offest_byte, copy_extend, 0, q));
+ return event_list.front();
+}
+
+template
+inline bool check_duplicate_entries(int count, T **entries) {
+ std::set uniqueEntries;
+
+ for (size_t i = 0; i < count; ++i) {
+ auto result = uniqueEntries.insert(entries[i]);
+ if (!result.second) {
+ return true;
+ }
+ }
+
+ return false;
+}
+} // namespace detail
+
+/// Get image_mem_wrapper according to sampled image handle.
+/// \param [in] handle The bindless image handle.
+/// \returns The image_mem_wrapper of sampled image.
+inline image_mem_wrapper *&get_img_mem(
+ const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
+ static std::map
+ img_mem_map;
+ return img_mem_map[handle];
+}
+
+/// Get image_mem_wrapper according to unsampled image handle.
+/// \param [in] handle The unsampled bindless image handle.
+/// \returns The image_mem_wrapper of unsampled image.
+inline image_mem_wrapper *&get_img_mem(
+ const sycl::ext::oneapi::experimental::unsampled_image_handle handle) {
+ static std::map
+ img_mem_map;
+ return img_mem_map[handle];
+}
+
+/// Associate an image memory wrapper with a bindless unsampled image handle in
+/// the global registry.
+///
+/// Inserts or updates an entry in the internal image memory map, using the
+/// specified bindless unsampled image handle as the key and the provided image
+/// memory wrapper as the mapped value.
+/// \param [in] handle The unsampled bindless image handle.
+/// \param [in] img_mem The image_mem_wrapper associated with the unsampled handle.
+static inline void set_img_mem(
+ const sycl::ext::oneapi::experimental::unsampled_image_handle handle,
+ image_mem_wrapper *img_mem) {
+ dpct::experimental::get_img_mem(handle) = img_mem;
+}
+
+/// Associate an image memory wrapper with a bindless sampled image handle in
+/// the global registry.
+///
+/// Inserts or updates an entry in the internal image memory map, using the
+/// specified bindless unsampled image handle as the key and the provided image
+/// memory wrapper as the mapped value.
+/// \param [in] handle The sampled bindless image handle.
+/// \param [in] img_mem The image_mem_wrapper associated with the sampled handle.
+static inline void set_img_mem(
+ const sycl::ext::oneapi::experimental::sampled_image_handle handle,
+ image_mem_wrapper *img_mem) {
+ dpct::experimental::get_img_mem(handle) = img_mem;
+}
+
+#ifdef _WIN32
+/// Map the resource memories to mem handles
+/// \param [in] count The count of resources to map.
+/// \param [in] handles The external mem wrappers used to map the resources.
+/// \param [in] q The queue used to map the resource with.
+inline void map_resources(int count, external_mem_wrapper **handles,
+ queue_ptr q_ptr = &get_default_queue()) {
+ if (detail::check_duplicate_entries(count, handles)) {
+ throw std::runtime_error(
+ "Duplicate handle entries found during resource mapping!");
+ }
+
+ for (int i = 0; i < count; i++) {
+ handles[i]->map_resource(*q_ptr);
+ }
+}
+
+/// Unmap the resource memories to mem handles
+/// \param [in] count The count of resources to unmap.
+/// \param [in] handles The external mem wrappers used to unmap the resources.
+/// \param [in] q The queue used to unmap the resource with.
+inline void unmap_resources(int count, external_mem_wrapper **handles,
+ queue_ptr q_ptr = &get_default_queue()) {
+ if (detail::check_duplicate_entries(count, handles)) {
+ throw std::runtime_error(
+ "Duplicate handle entries found during resource unmapping!");
+ }
+
+ for (int i = 0; i < count; i++) {
+ handles[i]->unmap_resource(*q_ptr);
+ }
+}
+#endif // _WIN32
+
+/// Imports resource's external memory into a SYCL external memory object.
+/// \param [in] extMem SYCL external memory object ptr to be initialized.
+/// \param [in] memHandleDesc Pointer to the external memory resource handle
+/// descriptor.
+/// \throws std::runtime_error If no valid external memory resource handle is
+/// provided.
+inline void
+import_external_memory(sycl::ext::oneapi::experimental::external_mem *extMem,
+ external_mem_handle_desc *memHandleDesc) {
+#ifdef _WIN32
+ if (memHandleDesc->get_win32_handle()) {
+ *extMem = sycl::ext::oneapi::experimental::import_external_memory(
+ sycl::ext::oneapi::experimental::external_mem_descriptor<
+ sycl::ext::oneapi::experimental::resource_win32_handle>{
+ {memHandleDesc->get_win32_handle()},
+ memHandleDesc->get_handle_type(),
+ memHandleDesc->get_res_size()},
+ get_default_queue());
+ } else if (memHandleDesc->get_win32_obj_name()) {
+ throw std::runtime_error(
+ "Importing external memory resource using Windows NT name"
+ " type is not supported. Only NT handle type is"
+ " supported");
+ }
+#else
+ if (memHandleDesc->get_fd_handle() != -1) {
+ *extMem = sycl::ext::oneapi::experimental::import_external_memory(
+ sycl::ext::oneapi::experimental::external_mem_descriptor<
+ sycl::ext::oneapi::experimental::resource_fd>{
+ {memHandleDesc->get_fd_handle()},
+ memHandleDesc->get_handle_type(),
+ memHandleDesc->get_res_size()},
+ get_default_queue());
+ }
+#endif // _WIN32
+ else {
+ throw std::runtime_error(
+ "No external memory resource handle is set for importing!");
+ }
+}
+
+/// Imports external semaphore resource into a SYCL external semaphore object.
+/// \param [in] extSem SYCL external semaphore object ptr to be initialized.
+/// \param [in] semHandleDesc Pointer to the external semaphore resource handle
+/// descriptor.
+/// \throws std::runtime_error If no valid external semaphore resource handle is
+/// provided.
+inline void import_external_semaphore(external_sem_wrapper **extSem,
+ external_sem_handle_desc *semHandleDesc) {
+ sycl::ext::oneapi::experimental::external_semaphore extSemObj;
+#ifdef _WIN32
+ if (semHandleDesc->get_win32_handle()) {
+ extSemObj = sycl::ext::oneapi::experimental::import_external_semaphore(
+ sycl::ext::oneapi::experimental::external_semaphore_descriptor<
+ sycl::ext::oneapi::experimental::resource_win32_handle>{
+ {semHandleDesc->get_win32_handle()},
+ semHandleDesc->get_handle_type()},
+ get_default_queue());
+ } else if (semHandleDesc->get_win32_obj_name()) {
+ throw std::runtime_error(
+ "Importing external semaphore resource using Windows NT name"
+ " type is not supported. Only NT handle type is"
+ " supported");
+ }
+#else
+ if (semHandleDesc->get_fd_handle() != -1) {
+ extSemObj = sycl::ext::oneapi::experimental::import_external_semaphore(
+ sycl::ext::oneapi::experimental::external_semaphore_descriptor<
+ sycl::ext::oneapi::experimental::resource_fd>{
+ {semHandleDesc->get_fd_handle()}, semHandleDesc->get_handle_type()},
+ get_default_queue());
+ }
+#endif // _WIN32
+ else {
+ throw std::runtime_error(
+ "No external semaphore resource handle is set for importing!");
+ }
+
+ *extSem =
+ new external_sem_wrapper(extSemObj, semHandleDesc->get_handle_type());
+}
+
+/// Signals the external semaphore resource.
+/// \param [in] extSem ptr to dpct external semaphore resource ptr
+/// \param [in] semSignalParams Pointer to the external semaphore signal
+/// parameters.
+/// \param [in] numExtSems Number of external semaphores to signal.
+/// \param [in] q_ptr The queue used to signal the external semaphore resource.
+inline void signal_external_semaphores(external_sem_wrapper **extSem,
+ external_sem_params *semSignalParams,
+ unsigned int numExtSems,
+ queue_ptr q_ptr = &get_default_queue()) {
+ for (int i = 0; i < numExtSems; i++) {
+ auto extSemType = extSem[i]->get_handle_type();
+ switch (extSemType) {
+#ifdef _WIN32
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ win32_nt_dx12_fence:
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ timeline_win32_nt_handle:
+ q_ptr->ext_oneapi_signal_external_semaphore(
+ extSem[i]->get(), semSignalParams[i].get_value());
+ break;
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ win32_nt_handle:
+ q_ptr->ext_oneapi_signal_external_semaphore(extSem[i]->get());
+ break;
+#else // _WIN32
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ timeline_fd:
+ q_ptr->ext_oneapi_signal_external_semaphore(
+ extSem[i]->get(), semSignalParams[i].get_value());
+ break;
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ opaque_fd:
+ q_ptr->ext_oneapi_signal_external_semaphore(extSem[i]->get());
+ break;
+#endif // _WIN32
+ default:
+ throw std::runtime_error(
+ "Unsupported external semaphore resource handle type!");
+ break;
+ }
+ }
+}
+
+/// Wait for the external semaphore resource.
+/// \param [in] extSem ptr to dpct external semaphore resource ptr
+/// \param [in] semWaitParams Pointer to the external semaphore wait parameters.
+/// \param [in] numExtSems Number of external semaphores to wait.
+/// \param [in] q_ptr The queue used to wait on the external semaphore resource.
+inline void wait_external_semaphores(external_sem_wrapper **extSem,
+ external_sem_params *semWaitParams,
+ unsigned int numExtSems,
+ queue_ptr q_ptr = &get_default_queue()) {
+ for (int i = 0; i < numExtSems; i++) {
+ auto extSemType = extSem[i]->get_handle_type();
+ switch (extSemType) {
+#ifdef _WIN32
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ win32_nt_dx12_fence:
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ timeline_win32_nt_handle:
+ q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get(),
+ semWaitParams[i].get_value());
+ break;
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ win32_nt_handle:
+ q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get());
+ break;
+#else // _WIN32
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ timeline_fd:
+ q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get(),
+ semWaitParams[i].get_value());
+ break;
+ case sycl::ext::oneapi::experimental::external_semaphore_handle_type::
+ opaque_fd:
+ q_ptr->ext_oneapi_wait_external_semaphore(extSem[i]->get());
+ break;
+#endif // _WIN32
+ default:
+ throw std::runtime_error(
+ "Unsupported external semaphore resource handle type!");
+ break;
+ }
+ }
+}
+
+/// Create bindless image according to image data and sampling info.
+/// \param [in] data The image data used to create bindless image.
+/// \param [in] info The image sampling info used to create bindless image.
+/// \param [in] q The queue where the image creation be executed.
+/// \returns The sampled image handle of created bindless image.
+static inline sycl::ext::oneapi::experimental::sampled_image_handle
+create_bindless_image(image_data data, sampling_info info,
+ sycl::queue q = get_default_queue()) {
+ auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
+ info.get_addressing_mode(), info.get_coordinate_normalization_mode(),
+ info.get_filtering_mode(), info.get_mipmap_filtering(),
+ info.get_min_mipmap_level_clamp(), info.get_max_mipmap_level_clamp(),
+ info.get_max_anisotropy());
+
+ switch (data.get_data_type()) {
+ case image_data_type::linear: {
+ // linear memory only use sycl::filtering_mode::nearest.
+ samp.filtering = sycl::filtering_mode::nearest;
+ // TODO: Use pointer to create image when bindless image support.
+ auto mem = new image_mem_wrapper(
+ data.get_channel(), data.get_x() / data.get_channel().get_total_size());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), samp, mem->get_desc(), q);
+ set_img_mem(img, mem);
+ auto ptr = data.get_data_ptr();
+#ifdef DPCT_USM_LEVEL_NONE
+ q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
+ mem->get_handle(), mem->get_desc())
+ .wait();
+#else
+ q.ext_oneapi_copy(ptr, mem->get_handle(), mem->get_desc()).wait();
+#endif
+ detail::get_img_info_map(img) = {data, info};
+ return img;
+ }
+ case image_data_type::pitch: {
+#ifdef DPCT_USM_LEVEL_NONE
+ auto mem =
+ new image_mem_wrapper(data.get_channel(), data.get_x(), data.get_y());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), samp, mem->get_desc(), q);
+ set_img_mem(img, mem);
+ q.ext_oneapi_copy(
+ get_buffer(data.get_data_ptr()).get_host_access().get_pointer(),
+ mem->get_handle(), mem->get_desc())
+ .wait();
+#else
+ auto desc = sycl::ext::oneapi::experimental::image_descriptor(
+ {data.get_x(), data.get_y()}, data.get_channel().get_channel_num(),
+ data.get_channel_type());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ data.get_data_ptr(), data.get_pitch(), samp, desc, q);
+#endif
+ detail::get_img_info_map(img) = {data, info};
+ return img;
+ }
+ case image_data_type::matrix: {
+ const auto mem = static_cast(data.get_data_ptr());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), samp, mem->get_desc(), q);
+ detail::get_img_info_map(img) = {data, info};
+ return img;
+ }
+ default:
+ throw std::runtime_error(
+ "Unsupported image_data_type in create_bindless_image!");
+ break;
+ }
+ // Must not reach here.
+ return sycl::ext::oneapi::experimental::sampled_image_handle();
+}
+
+/// Create bindless image according to image data.
+/// \param [in] data The image data used to create bindless image.
+/// \param [in] q The queue where the image creation be executed.
+/// \returns The sampled image handle of created bindless image.
+static inline sycl::ext::oneapi::experimental::unsampled_image_handle
+create_bindless_image(image_data data, sycl::queue q = get_default_queue()) {
+ switch (data.get_data_type()) {
+ case image_data_type::linear: {
+ // TODO: Use pointer to create image when bindless image support.
+ auto mem = new image_mem_wrapper(
+ data.get_channel(), data.get_x() / data.get_channel().get_total_size());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), mem->get_desc(), q);
+ set_img_mem(img, mem);
+ auto ptr = data.get_data_ptr();
+#ifdef DPCT_USM_LEVEL_NONE
+ q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
+ mem->get_handle(), mem->get_desc())
+ .wait();
+#else
+ q.ext_oneapi_copy(ptr, mem->get_handle(), mem->get_desc()).wait();
+#endif
+ return img;
+ }
+ case image_data_type::pitch: {
+ auto mem =
+ new image_mem_wrapper(data.get_channel(), data.get_x(), data.get_y());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), mem->get_desc(), q);
+ set_img_mem(img, mem);
+#ifdef DPCT_USM_LEVEL_NONE
+ q.ext_oneapi_copy(
+ get_buffer(data.get_data_ptr()).get_host_access().get_pointer(),
+ mem->get_handle(), mem->get_desc())
+ .wait();
+#else
+ q.ext_oneapi_copy(data.get_data_ptr(), mem->get_handle(), mem->get_desc())
+ .wait();
+
+#endif
+ return img;
+ }
+ case image_data_type::matrix: {
+ const auto mem = static_cast(data.get_data_ptr());
+ auto img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), mem->get_desc(), q);
+ return img;
+ }
+ default:
+ throw std::runtime_error(
+ "Unsupported image_data_type in create_bindless_image!");
+ break;
+ }
+ // Must not reach here.
+ return sycl::ext::oneapi::experimental::unsampled_image_handle();
+}
+
+/// Destroy bindless image.
+/// \param [in] handle The bindless image should be destroyed.
+/// \param [in] q The queue where the image destruction be executed.
+template
+static inline void destroy_bindless_image(T handle,
+ sycl::queue q = get_default_queue()) {
+ auto &mem = get_img_mem(handle);
+ if (mem) {
+ delete mem;
+ mem = nullptr;
+ }
+ sycl::ext::oneapi::experimental::destroy_image_handle(handle, q);
+}
+
+/// Get the image data according to sampled image handle.
+/// \param [in] handle The bindless image handle.
+/// \returns The image data of sampled image.
+template static inline image_data get_data(const T handle) {
+ return detail::get_img_info_map(handle).first;
+}
+
+/// Get the sampling info according to sampled image handle.
+/// \param [in] handle The bindless image handle.
+/// \returns The sampling info of sampled image.
+static inline sampling_info get_sampling_info(
+ const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
+ return detail::get_img_info_map(handle).second;
+}
+
+/// The base class of different template specialization bindless_image_wrapper
+/// class.
+class bindless_image_wrapper_base {
+public:
+ bindless_image_wrapper_base() {
+ // Make sure that singleton class dev_mgr will destruct later than this.
+ dev_mgr::instance();
+ }
+
+ /// Destroy bindless image wrapper.
+ ~bindless_image_wrapper_base() {
+ destroy_bindless_image(_img, get_default_queue());
+ }
+
+ /// Attach linear data to bindless image.
+ /// \param [in] data The linear data used to create bindless image.
+ /// \param [in] size The size of linear data used to create bindless image.
+ /// \param [in] channel The image channel used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(void *data, size_t size, const image_channel &channel,
+ sycl::queue q = get_default_queue()) {
+ detach(q);
+ auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
+ _addressing_mode, _coordinate_normalization_mode, _filtering_mode);
+ // TODO: Use pointer to create image when bindless image support.
+ auto mem = new image_mem_wrapper(channel, size);
+ _img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), samp, mem->get_desc(), q);
+ set_img_mem(_img, mem);
+ auto ptr = data;
+#ifdef DPCT_USM_LEVEL_NONE
+ q.ext_oneapi_copy(get_buffer(data).get_host_access().get_pointer(),
+ mem->get_handle(), mem->get_desc())
+ .wait();
+#else
+ q.ext_oneapi_copy(ptr, mem->get_handle(), mem->get_desc()).wait();
+#endif
+ }
+
+ /// Attach linear data to bindless image.
+ /// \param [in] data The linear data used to create bindless image.
+ /// \param [in] size The size of linear data used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(void *data, size_t size, sycl::queue q = get_default_queue()) {
+ attach(data, size, _channel, q);
+ }
+
+ /// Attach device_ptr data to bindless image.
+ /// \param [in] desc The image_descriptor used to create bindless image.
+ /// \param [in] ptr The data pointer used to create bindless image.
+ /// \param [in] pitch The pitch of 2D data used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(const sycl::ext::oneapi::experimental::image_descriptor *desc,
+ device_ptr ptr, size_t pitch,
+ sycl::queue q = get_default_queue()) {
+ detach(q);
+ auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
+ _addressing_mode, _coordinate_normalization_mode, _filtering_mode);
+#ifdef DPCT_USM_LEVEL_NONE
+ auto mem = new image_mem_wrapper(desc);
+ _img = sycl::ext::oneapi::experimental::create_image(mem->get_handle(),
+ samp, *desc, q);
+ set_img_mem(_img, mem);
+ q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
+ mem->get_handle(), mem->get_desc())
+ .wait();
+#else
+ _img = sycl::ext::oneapi::experimental::create_image(ptr, pitch, samp,
+ *desc, q);
+#endif
+ }
+
+ /// Attach 2D data to bindless image.
+ /// \param [in] data The 2D data used to create bindless image.
+ /// \param [in] width The width of 2D data used to create bindless image.
+ /// \param [in] height The height of 2D data used to create bindless image.
+ /// \param [in] pitch The pitch of 2D data used to create bindless image.
+ /// \param [in] channel The image channel used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(void *data, size_t width, size_t height, size_t pitch,
+ const image_channel &channel,
+ sycl::queue q = get_default_queue()) {
+ auto desc = sycl::ext::oneapi::experimental::image_descriptor(
+ {width, height}, channel.get_channel_num(), channel.get_channel_type());
+ attach(&desc, static_cast(data), pitch, q);
+ }
+
+ /// Attach 2D data to bindless image.
+ /// \param [in] data The 2D data used to create bindless image.
+ /// \param [in] width The width of 2D data used to create bindless image.
+ /// \param [in] height The height of 2D data used to create bindless image.
+ /// \param [in] pitch The pitch of 2D data used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(void *data, size_t width, size_t height, size_t pitch,
+ sycl::queue q = get_default_queue()) {
+ attach(data, width, height, pitch, _channel, q);
+ }
+
+ /// Attach image memory to bindless image.
+ /// \param [in] mem The image memory used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(image_mem_wrapper *mem, const image_channel &channel,
+ sycl::queue q = get_default_queue()) {
+ detach(q);
+ auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
+ _addressing_mode, _coordinate_normalization_mode, _filtering_mode);
+ _img = sycl::ext::oneapi::experimental::create_image(
+ mem->get_handle(), samp, mem->get_desc(), q);
+ set_img_mem(_img, mem);
+ }
+
+ /// Attach image memory to bindless image.
+ /// \param [in] mem The image memory used to create bindless image.
+ /// \param [in] q The queue where the image creation be executed.
+ void attach(image_mem_wrapper *mem, sycl::queue q = get_default_queue()) {
+ attach(mem, _channel, q);
+ }
+
+ /// Detach bindless image data.
+ /// \param [in] q The queue where the image destruction be executed.
+ void detach(sycl::queue q = get_default_queue()) {
+ destroy_bindless_image(_img, q);
+ }
+
+ /// Set image channel of bindless image.
+ /// \param [in] channel The image channel used to set.
+ void set_channel(image_channel channel) { _channel = channel; }
+ /// Get image channel of bindless image.
+ /// \returns The image channel of bindless image.
+ image_channel get_channel() { return _channel; }
+
+ /// Set channel size of image channel.
+ /// \param [in] channel_num The channels number to set.
+ /// \param [in] channel_size The size for each channel in bits.
+ void set_channel_size(unsigned channel_num, unsigned channel_size) {
+ return _channel.set_channel_size(channel_num, channel_size);
+ }
+ /// Get channel size of image channel.
+ /// \returns The size for each channel in bits.
+ unsigned get_channel_size() { return _channel.get_channel_size(); }
+
+ /// Set image channel data type of image channel.
+ /// \param [in] type The image channel data type to set.
+ void set_channel_data_type(image_channel_data_type type) {
+ _channel.set_channel_data_type(type);
+ }
+ /// Get image channel data type of image channel.
+ /// \returns The image channel data type of image channel.
+ image_channel_data_type get_channel_data_type() {
+ return _channel.get_channel_data_type();
+ }
+
+ /// Set channel num of bindless image.
+ /// \param [in] filtering_mode The channel num used to set.
+ void set_channel_num(unsigned num) { return _channel.set_channel_num(num); }
+
+ /// Set channel type of bindless image.
+ /// \param [in] filtering_mode The channel type used to set.
+ void set_channel_type(sycl::image_channel_type type) {
+ return _channel.set_channel_type(type);
+ }
+
+ /// Set addressing mode of bindless image.
+ /// \param [in] addressing_mode The addressing mode used to set.
+ void set(sycl::addressing_mode addressing_mode) {
+ _addressing_mode = addressing_mode;
+ }
+ /// Get addressing mode of bindless image.
+ /// \returns The addressing mode of bindless image.
+ sycl::addressing_mode get_addressing_mode() { return _addressing_mode; }
+
+ /// Set coordinate normalization mode of bindless image.
+ /// \param [in] coordinate_normalization_mode The coordinate normalization
+ /// mode used to set.
+ void set(sycl::coordinate_normalization_mode coordinate_normalization_mode) {
+ _coordinate_normalization_mode = coordinate_normalization_mode;
+ }
+ /// Set coordinate normalization mode of bindless image.
+ /// \param [in] is_normalized Determine whether the coordinate should be
+ /// normalized.
+ void set_coordinate_normalization_mode(int is_normalized) {
+ _coordinate_normalization_mode =
+ is_normalized ? sycl::coordinate_normalization_mode::normalized
+ : sycl::coordinate_normalization_mode::unnormalized;
+ }
+ /// Get coordinate normalization mode of bindless image.
+ /// \returns The coordinate normalization mode of bindless image.
+ bool is_coordinate_normalized() {
+ return _coordinate_normalization_mode ==
+ sycl::coordinate_normalization_mode::normalized;
+ }
+
+ /// Set filtering mode of bindless image.
+ /// \param [in] filtering_mode The filtering mode used to set.
+ void set(sycl::filtering_mode filtering_mode) {
+ _filtering_mode = filtering_mode;
+ }
+ /// Get filtering mode of bindless image.
+ /// \returns The filtering mode of bindless image.
+ sycl::filtering_mode get_filtering_mode() { return _filtering_mode; }
+
+ /// Get bindless image handle.
+ /// \returns The sampled image handle of created bindless image.
+ inline sycl::ext::oneapi::experimental::sampled_image_handle get_handle() {
+ return _img;
+ }
+
+ /// Set mipmap sample filtering mode for bindless image handle
+ /// \param [in] filtering_mode The mipmap filtering mode
+ inline void set_mip_filtering_mode(sycl::filtering_mode mode) {
+ auto sampling_info = get_sampling_info(_img);
+ sampling_info.set_mipmap_filtering(mode);
+ }
+
+ /// Get mipmap sample filtering mode for bindless image handle
+ /// \return The mipmap filtering mode
+ inline sycl::filtering_mode get_mip_filtering_mode(void) {
+ auto sampling_info = get_sampling_info(_img);
+ return sampling_info.get_mipmap_filtering();
+ }
+
+ /// Get mipmap sample filtering mode for bindless image handle
+ /// \return The mipmap filtering mode
+ inline void get_mip_level_clamp(float *min_level_clamp,
+ float *max_level_clamp) {
+ auto sampling_info = get_sampling_info(_img);
+ *min_level_clamp = sampling_info.get_min_mipmap_level_clamp();
+ *max_level_clamp = sampling_info.get_max_mipmap_level_clamp();
+ }
+
+ /// Get mipmap memory wrapper attached the bindless image
+ /// \return The mipmap memory wrapper
+ inline image_mem_wrapper *get_attached_mipmap_data(void) {
+ auto mem = get_img_mem(_img);
+
+ if (mem->get_image_type() !=
+ sycl::ext::oneapi::experimental::image_type::mipmap)
+ throw std::runtime_error("Bindless image data is not of mipmap type");
+
+ return mem;
+ }
+
+private:
+ image_channel _channel;
+ sycl::addressing_mode _addressing_mode = sycl::addressing_mode::clamp_to_edge;
+ sycl::coordinate_normalization_mode _coordinate_normalization_mode =
+ sycl::coordinate_normalization_mode::unnormalized;
+ sycl::filtering_mode _filtering_mode = sycl::filtering_mode::nearest;
+ sycl::ext::oneapi::experimental::sampled_image_handle _img{0};
+};
+
+/// The wrapper class of bindless image.
+template
+class bindless_image_wrapper : public bindless_image_wrapper_base {
+public:
+ /// Create bindless image wrapper according to template argument \p T.
+ bindless_image_wrapper() { set_channel(image_channel::create()); }
+};
+
+/// Asynchronously copies from the image memory to memory specified by a
+/// pointer, The return of the function does NOT guarantee the copy is
+/// completed.
+/// \param [in] dest The destination memory address.
+/// \param [in] p The pitch of destination memory.
+/// \param [in] src The source image memory.
+/// \param [in] w_offset_src The x offset of source image memory.
+/// \param [in] h_offset_src The y offset of source image memory.
+/// \param [in] w The width of matrix to be copied.
+/// \param [in] h The height of matrix to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void async_dpct_memcpy(void *dest, size_t p,
+ const image_mem_wrapper *src,
+ size_t w_offset_src, size_t h_offset_src,
+ size_t w, size_t h,
+ sycl::queue q = get_default_queue()) {
+ detail::dpct_memcpy(src->get_handle(), src->get_desc(), w_offset_src,
+ h_offset_src, dest, p, w, h, q);
+}
+
+/// Synchronously copies from the image memory to memory specified by a
+/// pointer, The function will return after the copy is completed.
+/// \param [in] dest The destination memory address.
+/// \param [in] p The pitch of destination memory.
+/// \param [in] src The source image memory.
+/// \param [in] w_offset_src The x offset of source image memory.
+/// \param [in] h_offset_src The y offset of source image memory.
+/// \param [in] w The width of matrix to be copied.
+/// \param [in] h The height of matrix to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void dpct_memcpy(void *dest, size_t p,
+ const image_mem_wrapper *src,
+ size_t w_offset_src, size_t h_offset_src,
+ size_t w, size_t h,
+ sycl::queue q = get_default_queue()) {
+ detail::dpct_memcpy(src->get_handle(), src->get_desc(), w_offset_src,
+ h_offset_src, dest, p, w, h, q)
+ .wait();
+}
+
+/// Asynchronously copies from the image memory to memory specified by a
+/// pointer, The return of the function does NOT guarantee the copy is
+/// completed.
+/// \param [in] dest The destination memory address.
+/// \param [in] src The source image memory.
+/// \param [in] w_offset_src The x offset of source image memory.
+/// \param [in] h_offset_src The y offset of source image memory.
+/// \param [in] s The size to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void async_dpct_memcpy(void *dest, const image_mem_wrapper *src,
+ size_t w_offset_src, size_t h_offset_src,
+ size_t s,
+ sycl::queue q = get_default_queue()) {
+ detail::dpct_memcpy(src->get_handle(), src->get_desc(), w_offset_src,
+ h_offset_src, dest, s, q);
+}
+
+/// Synchronously copies from the image memory to memory specified by a
+/// pointer, The function will return after the copy is completed.
+/// \param [in] dest The destination memory address.
+/// \param [in] src The source image memory.
+/// \param [in] w_offset_src The x offset of source image memory.
+/// \param [in] h_offset_src The y offset of source image memory.
+/// \param [in] s The size to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void dpct_memcpy(void *dest, const image_mem_wrapper *src,
+ size_t w_offset_src, size_t h_offset_src,
+ size_t s, sycl::queue q = get_default_queue()) {
+ sycl::event::wait(detail::dpct_memcpy(src->get_handle(), src->get_desc(),
+ w_offset_src, h_offset_src, dest, s,
+ q));
+}
+
+/// Asynchronously copies from memory specified by a pointer to the image
+/// memory, The return of the function does NOT guarantee the copy is completed.
+/// \param [in] dest The destination image memory.
+/// \param [in] w_offset_dest The x offset of destination image memory.
+/// \param [in] h_offset_dest The y offset of destination image memory.
+/// \param [in] src The source memory address.
+/// \param [in] p The pitch of source memory.
+/// \param [in] w The width of matrix to be copied.
+/// \param [in] h The height of matrix to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void async_dpct_memcpy(image_mem_wrapper *dest,
+ size_t w_offset_dest, size_t h_offset_dest,
+ const void *src, size_t p, size_t w,
+ size_t h,
+ sycl::queue q = get_default_queue()) {
+ detail::dpct_memcpy(src, dest->get_handle(), dest->get_desc(), w_offset_dest,
+ h_offset_dest, p, w, h, q);
+}
+
+/// Synchronously copies from memory specified by a pointer to the image
+/// memory, The function will return after the copy is completed.
+/// \param [in] dest The destination image memory.
+/// \param [in] w_offset_dest The x offset of destination image memory.
+/// \param [in] h_offset_dest The y offset of destination image memory.
+/// \param [in] src The source memory address.
+/// \param [in] p The pitch of source memory.
+/// \param [in] w The width of matrix to be copied.
+/// \param [in] h The height of matrix to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
+ size_t h_offset_dest, const void *src, size_t p,
+ size_t w, size_t h,
+ sycl::queue q = get_default_queue()) {
+ detail::dpct_memcpy(src, dest->get_handle(), dest->get_desc(), w_offset_dest,
+ h_offset_dest, p, w, h, q)
+ .wait();
+}
+
+/// Asynchronously copies from memory specified by a pointer to the image
+/// memory, The return of the function does NOT guarantee the copy is completed.
+/// \param [in] dest The destination image memory.
+/// \param [in] w_offset_dest The x offset of destination image memory.
+/// \param [in] h_offset_dest The y offset of destination image memory.
+/// \param [in] src The source memory address.
+/// \param [in] s The size to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void async_dpct_memcpy(image_mem_wrapper *dest,
+ size_t w_offset_dest, size_t h_offset_dest,
+ const void *src, size_t s,
+ sycl::queue q = get_default_queue()) {
+ detail::dpct_memcpy(src, dest->get_handle(), dest->get_desc(), w_offset_dest,
+ h_offset_dest, s, q);
+}
+
+/// Synchronously copies from memory specified by a pointer to the image
+/// memory, The function will return after the copy is completed.
+/// \param [in] dest The destination image memory.
+/// \param [in] w_offset_dest The x offset of destination image memory.
+/// \param [in] h_offset_dest The y offset of destination image memory.
+/// \param [in] src The source memory address.
+/// \param [in] s The size to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
+ size_t h_offset_dest, const void *src, size_t s,
+ sycl::queue q = get_default_queue()) {
+ sycl::event::wait(detail::dpct_memcpy(src, dest->get_handle(),
+ dest->get_desc(), w_offset_dest,
+ h_offset_dest, s, q));
+}
+
+/// Synchronously copies from image memory to the image memory, The function
+/// will return after the copy is completed.
+/// \param [in] dest The destination image memory.
+/// \param [in] w_offset_dest The x offset of destination image memory.
+/// \param [in] h_offset_dest The y offset of destination image memory.
+/// \param [in] src The source image memory.
+/// \param [in] w_offset_src The x offset of source image memory.
+/// \param [in] h_offset_src The y offset of source image memory.
+/// \param [in] w The width of matrix to be copied.
+/// \param [in] h The height of matrix to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
+ size_t h_offset_dest,
+ const image_mem_wrapper *src,
+ size_t w_offset_src, size_t h_offset_src,
+ size_t w, size_t h,
+ sycl::queue q = get_default_queue()) {
+ auto temp = (void *)sycl::malloc_device(w * h, q);
+ // TODO: Need change logic when sycl support image_mem to image_mem copy.
+ dpct_memcpy(temp, w, src, w_offset_src, h_offset_src, w, h, q);
+ dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, w, w, h, q);
+ sycl::free(temp, q);
+}
+
+/// Synchronously copies from image memory to the image memory, The function
+/// will return after the copy is completed.
+/// \param [in] dest The destination image memory.
+/// \param [in] w_offset_dest The x offset of destination image memory.
+/// \param [in] h_offset_dest The y offset of destination image memory.
+/// \param [in] src The source image memory.
+/// \param [in] w_offset_src The x offset of source image memory.
+/// \param [in] h_offset_src The y offset of source image memory.
+/// \param [in] s The size to be copied.
+/// \param [in] q The queue to execute the copy task.
+static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
+ size_t h_offset_dest,
+ const image_mem_wrapper *src,
+ size_t w_offset_src, size_t h_offset_src,
+ size_t s, sycl::queue q = get_default_queue()) {
+ auto temp = (void *)sycl::malloc_device(s, q);
+ // TODO: Need change logic when sycl support image_mem to image_mem copy.
+ dpct_memcpy(temp, src, w_offset_src, h_offset_src, s, q);
+ dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, s, q);
+ sycl::free(temp, q);
+}
+// A wrapper for sycl fetch_image function for the byte addressing image.
+template
+DataT fetch_image_by_byte(
+ const sycl::ext::oneapi::experimental::unsampled_image_handle &imageHandle,
+ CoordT &&coords) {
+ if constexpr (std::is_scalar_v) {
+ return sycl::ext::oneapi::experimental::fetch_image(
+ imageHandle, coords / sizeof(DataT));
+ } else {
+ coords[0] = coords[0] / sizeof(DataT);
+ return sycl::ext::oneapi::experimental::fetch_image(
+ imageHandle, coords);
+ }
+}
+
+using image_mem_wrapper_ptr = image_mem_wrapper *;
+#ifdef _WIN32
+using external_mem_wrapper_ptr = external_mem_wrapper *;
+#else
+using external_mem_wrapper_ptr = external_mem_wrapper_base *;
+#endif // _WIN32
+using external_sem_wrapper_ptr = external_sem_wrapper *;
+using bindless_image_wrapper_base_p = bindless_image_wrapper_base *;
+
+#endif
+
+} // namespace experimental
+} // namespace dpct
+
+#endif // !__DPCT_BINDLESS_IMAGE_HPP__
diff --git a/sycl/include/dpct/blas_gemm_utils.hpp b/sycl/include/dpct/blas_gemm_utils.hpp
new file mode 100644
index 00000000000..f41fe0737b4
--- /dev/null
+++ b/sycl/include/dpct/blas_gemm_utils.hpp
@@ -0,0 +1,896 @@
+//==---- blas_gemm_utils.hpp ----------------------*- C++ -*----------------==//
+//
+// Copyright (C) Intel Corporation
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// See https://llvm.org/LICENSE.txt for license information.
+//
+//===----------------------------------------------------------------------===//
+// This file contains the implementation of GEneral Matrix Multiplication by
+// using oneDNN. More datatype combinations and the epilogue can be supported by
+// oneDNN, which is not available in blas_utils.hpp using oneMKL.
+//===----------------------------------------------------------------------===//
+
+#ifndef __DPCT_BLAS_GEMM_UTILS_HPP__
+#define __DPCT_BLAS_GEMM_UTILS_HPP__
+
+#include "compat_service.hpp"
+#include "blas_utils.hpp"
+#include "dnnl_utils.hpp"
+
+namespace dpct::blas_gemm::experimental {
+enum class order_t : std::uint8_t {
+ col,
+ row,
+ col32,
+ col4_4r2_8c,
+ col32_2r_4r4
+};
+} // namespace dpct::blas_gemm::experimental
+
+#include "detail/blas_gemm_utils_detail.hpp"
+
+namespace dpct::blas_gemm::experimental {
+enum class pointer_mode_t {
+ host,
+ device,
+ device_vector,
+ alpha_device_vector_beta_zero,
+ alpha_device_vector_beta_host
+};
+enum class epilogue_t {
+ nop = 1,
+ relu,
+ bias,
+ gelu,
+ gelu_bias,
+ gelu_aux,
+ gelu_aux_bias,
+ dgelu,
+ bgradb
+};
+
+class descriptor;
+using descriptor_ptr = descriptor *;
+class matrix_layout_t;
+using matrix_layout_ptr = matrix_layout_t *;
+class matmul_desc_t;
+using matmul_desc_ptr = matmul_desc_t *;
+class transform_desc_t;
+using transform_desc_ptr = transform_desc_t *;
+
+class descriptor {
+public:
+ descriptor() {}
+ void init(sycl::queue *q_ptr) {
+ _engine = ::dnnl::sycl_interop::make_engine(q_ptr->get_device(),
+ q_ptr->get_context());
+ _engine_stream = ::dnnl::sycl_interop::make_stream(_engine, *q_ptr);
+ }
+ ::dnnl::engine get_engine() const noexcept { return _engine; }
+ ::dnnl::stream get_engine_stream() const noexcept { return _engine_stream; };
+
+private:
+ ::dnnl::engine _engine;
+ ::dnnl::stream _engine_stream;
+};
+
+class matrix_layout_t {
+public:
+ enum class attribute {
+ type,
+ order,
+ rows,
+ cols,
+ ld,
+ batch_count,
+ strided_batch_offset
+ };
+
+ matrix_layout_t(library_data_t type, std::uint64_t rows, std::uint64_t cols,
+ std::int64_t ld)
+ : _type(type), _rows(rows), _cols(cols), _ld(ld) {}
+
+ void set_attribute(attribute attr, const void *mem) {
+ get_set_attr(attr, const_cast(mem));
+ }
+ void get_attribute(attribute attr, void *mem) {
+ get_set_attr(attr, mem);
+ }
+
+private:
+ template void get_set_attr(attribute attr, void *mem) {
+#define CASE(tag) \
+ case attribute::tag: \
+ if constexpr (is_set) { \
+ _##tag = *static_cast(mem); \
+ } else { \
+ *static_cast(mem) = _##tag; \
+ } \
+ break;
+ switch (attr) {
+ CASE(type)
+ CASE(order)
+ CASE(rows)
+ CASE(cols)
+ CASE(ld)
+ CASE(batch_count)
+ CASE(strided_batch_offset)
+ }
+#undef CASE
+ }
+
+ library_data_t _type;
+ order_t _order = order_t::col;
+ std::uint64_t _rows;
+ std::uint64_t _cols;
+ std::int64_t _ld;
+ std::uint64_t _batch_count = 1;
+ std::uint64_t _strided_batch_offset = 0;
+
+ friend sycl::event matmul(descriptor_ptr handle, matmul_desc_ptr computeDesc,
+ const void *alpha, const void *a,
+ matrix_layout_ptr a_desc, const void *b,
+ matrix_layout_ptr b_desc, const void *beta,
+ const void *c, matrix_layout_ptr c_desc, void *d,
+ matrix_layout_ptr d_desc,
+ ::dpct::cs::queue_ptr q_ptr);
+ friend sycl::event
+ matrix_transform(transform_desc_ptr transform_desc, const void *alpha,
+ const void *a, matrix_layout_ptr a_desc, const void *beta,
+ const void *b, matrix_layout_ptr b_desc, void *c,
+ matrix_layout_ptr c_desc, ::dpct::cs::queue_ptr q_ptr);
+};
+
+class matmul_desc_t {
+public:
+ enum class attribute {
+ compute_type,
+ scale_type,
+ bias_data_type,
+ bias_pointer,
+ pointer_mode,
+ trans_a,
+ trans_b,
+ trans_c,
+ epilogue,
+ epilogue_aux_ld,
+ epilogue_aux_pointer,
+ epilogue_aux_data_type,
+ a_scale_pointer,
+ b_scale_pointer,
+ d_scale_pointer,
+ absmax_d_pointer,
+ unsupport
+ };
+
+ matmul_desc_t(compute_type compute_type, library_data_t scale_type)
+ : _compute_type(compute_type), _scale_type(scale_type) {}
+
+ void set_attribute(attribute attr, const void *mem) {
+ if (attr != attribute::unsupport)
+ get_set_attr(attr, const_cast(mem));
+ }
+ void get_attribute(attribute attr, void *mem) {
+ if (attr != attribute::unsupport)
+ get_set_attr(attr, mem);
+ }
+
+private:
+ template void get_set_attr(attribute attr, void *mem) {
+#define CASE(tag) \
+ case attribute::tag: \
+ if constexpr (is_set) { \
+ _##tag = *static_cast(mem); \
+ } else { \
+ *static_cast(mem) = _##tag; \
+ } \
+ break;
+ switch (attr) {
+ CASE(compute_type)
+ CASE(scale_type)
+ CASE(bias_data_type)
+ CASE(bias_pointer)
+ CASE(pointer_mode)
+ CASE(trans_a)
+ CASE(trans_b)
+ CASE(trans_c)
+ CASE(epilogue)
+ CASE(a_scale_pointer)
+ CASE(b_scale_pointer)
+ CASE(d_scale_pointer)
+ CASE(absmax_d_pointer)
+ CASE(epilogue_aux_ld)
+ CASE(epilogue_aux_pointer)
+ CASE(epilogue_aux_data_type)
+ default:
+ break;
+ }
+#undef CASE
+ }
+
+ compute_type _compute_type;
+ library_data_t _scale_type;
+ library_data_t _bias_data_type = library_data_t::real_float;
+ pointer_mode_t _pointer_mode = pointer_mode_t::host;
+ oneapi::mkl::transpose _trans_a = oneapi::mkl::transpose::nontrans;
+ oneapi::mkl::transpose _trans_b = oneapi::mkl::transpose::nontrans;
+ oneapi::mkl::transpose _trans_c = oneapi::mkl::transpose::nontrans;
+ epilogue_t _epilogue = epilogue_t::nop;
+ library_data_t _epilogue_aux_data_type = library_data_t::real_float;
+ size_t _epilogue_aux_ld = 0;
+ void *_a_scale_pointer = nullptr;
+ void *_b_scale_pointer = nullptr;
+ void *_d_scale_pointer = nullptr;
+ void *_absmax_d_pointer = nullptr;
+ void *_bias_pointer = nullptr;
+ void *_epilogue_aux_pointer = nullptr;
+
+ friend sycl::event matmul(descriptor_ptr handle, matmul_desc_ptr computeDesc,
+ const void *alpha, const void *a,
+ matrix_layout_ptr a_desc, const void *b,
+ matrix_layout_ptr b_desc, const void *beta,
+ const void *c, matrix_layout_ptr c_desc, void *d,
+ matrix_layout_ptr d_desc,
+ ::dpct::cs::queue_ptr q_ptr);
+};
+
+/// This function does the following operations:
+/// (1) D_temp = epilogue(alpha * scale_a * op_a(A) * scale_b * op_b(B) + beta * C)
+/// (2) Amax = absmax(D_temp) when matmul_desc_t::attribute::absmax_d_pointer is specified
+/// (3) D = scale_d * D_temp
+/// "op_a" is specified by the matmul_desc_t::attribute::trans_a
+/// (default is nontrans)
+/// "op_b" is specified by the matmul_desc_t::attribute::trans_b
+/// (default is nontrans)
+/// "scale_a" is specified by the matmul_desc_t::attribute::a_scale_pointer
+/// (default is nullptr, which behaves as 1.0)
+/// "scale_b" is specified by the matmul_desc_t::attribute::b_scale_pointer
+/// (default is nullptr, which behaves as 1.0)
+/// "scale_d" is specified by the matmul_desc_t::attribute::d_scale_pointer
+/// (default is nullptr, which behaves as 1.0)
+/// "alpha" can be a scalar value or a vector, which is specified by the
+/// matmul_desc_tattribute::::pointer_mode
+/// "epilogue" is specified by the matmul_desc_t::attribute::epilogue
+/// Currently, this function only supports the following type combinations:
+/// scale_type==int32 && a_type==int8 && b_type==int8 && c_type==int32;
+/// scale_type==float && a_type==int8 && b_type==int8 && c_type==int8;
+/// scale_type==float && a_type==int8 && b_type==int8 && c_type==int32;
+/// scale_type==float && a_type==float && b_type==float && c_type==float.
+/// Currently, this function only supports beta==0 or beta==1.
+/// Currently, this function only supports the relu, bias, gelu, gelu_bias,
+/// gelu_aux, gelu_aux_bias, dgelu and bgradb epilogue.
+/// NOTE: Non-col-major matrix will be converted to col-major matrix before.
+/// TODO: Impl row-major matmul without layout conversion.
+/// multiplication and converted back after multiplication.
+/// \param [in] handle A handle containing context info.
+/// \param [in] compute_desc Describe the computation.
+/// \param [in] alpha Scaling factor alpha.
+/// \param [in] a Input matrix A.
+/// \param [in] a_desc Describe the matrix A.
+/// \param [in] b Input matrix B.
+/// \param [in] b_desc Describe the matrix B.
+/// \param [in] beta Scaling factor beta.
+/// \param [in] c Input matrix C.
+/// \param [in] c_desc Describe the matrix C.
+/// \param [out] d Output matrix D.
+/// \param [in] d_desc Describe the matrix D.
+/// \param [in] q_ptr The queue where the routine should be executed.
+inline sycl::event matmul(descriptor_ptr handle, matmul_desc_ptr compute_desc,
+ const void *alpha, const void *a,
+ matrix_layout_ptr a_desc, const void *b,
+ matrix_layout_ptr b_desc, const void *beta,
+ const void *c, matrix_layout_ptr c_desc, void *d,
+ matrix_layout_ptr d_desc,
+ ::dpct::cs::queue_ptr q_ptr) {
+ const size_t m = compute_desc->_trans_a == oneapi::mkl::transpose::nontrans
+ ? a_desc->_rows
+ : a_desc->_cols;
+ const size_t n = d_desc->_cols;
+ const size_t k = compute_desc->_trans_b == oneapi::mkl::transpose::nontrans
+ ? b_desc->_rows
+ : b_desc->_cols;
+ const library_data_t a_type = a_desc->_type;
+ const library_data_t b_type = b_desc->_type;
+ const library_data_t c_type = c_desc->_type;
+ const library_data_t d_type = d_desc->_type;
+ const library_data_t scale_type = compute_desc->_scale_type;
+
+ if (!q_ptr)
+ q_ptr = &::dpct::cs::get_default_queue();
+ handle->init(q_ptr);
+ bool vector_alpha = false;
+ bool device_alpha = false;
+ if (compute_desc->_pointer_mode == pointer_mode_t::device_vector ||
+ compute_desc->_pointer_mode ==
+ pointer_mode_t::alpha_device_vector_beta_zero ||
+ compute_desc->_pointer_mode ==
+ pointer_mode_t::alpha_device_vector_beta_host) {
+ vector_alpha = true;
+ device_alpha = true;
+ } else if (compute_desc->_pointer_mode == pointer_mode_t::device) {
+ device_alpha = true;
+ }
+
+ bool beta_is_zero = true;
+ if (beta != nullptr) {
+ int beta_value = detail::type_dispatch(
+ compute_desc->_scale_type, beta, q_ptr);
+ if (beta_value != 0) {
+ beta_is_zero = false;
+ if (beta_value != 1)
+ throw std::runtime_error(
+ "dpct::blas_gemm::experimental::matmul() does "
+ "not support non-zero and non-one beta currently.");
+ }
+ }
+
+ if (compute_desc->_epilogue != epilogue_t::nop &&
+ compute_desc->_epilogue != epilogue_t::relu &&
+ compute_desc->_epilogue != epilogue_t::bias &&
+ compute_desc->_epilogue != epilogue_t::gelu &&
+ compute_desc->_epilogue != epilogue_t::gelu_bias &&
+ compute_desc->_epilogue != epilogue_t::gelu_aux &&
+ compute_desc->_epilogue != epilogue_t::gelu_aux_bias &&
+ compute_desc->_epilogue != epilogue_t::dgelu &&
+ compute_desc->_epilogue != epilogue_t::bgradb) {
+ throw std::runtime_error(
+ "dpct::blas_gemm::experimental::matmul() only "
+ "supports relu, bias, gelu, gelu_bias, gelu_aux, "
+ "gelu_aux_bias, dgelu and bgradb epilogue currently.");
+ }
+
+ if (!(compute_desc->_scale_type == library_data_t::real_int32 &&
+ a_desc->_type == library_data_t::real_int8 &&
+ b_desc->_type == library_data_t::real_int8 &&
+ c_desc->_type == library_data_t::real_int32) &&
+ !(compute_desc->_scale_type == library_data_t::real_float &&
+ a_desc->_type == library_data_t::real_int8 &&
+ b_desc->_type == library_data_t::real_int8 &&
+ c_desc->_type == library_data_t::real_int8) &&
+ !(compute_desc->_scale_type == library_data_t::real_float &&
+ a_desc->_type == library_data_t::real_int8 &&
+ b_desc->_type == library_data_t::real_int8 &&
+ c_desc->_type == library_data_t::real_int32) &&
+ !(compute_desc->_scale_type == library_data_t::real_float &&
+ a_desc->_type == library_data_t::real_float &&
+ b_desc->_type == library_data_t::real_float &&
+ c_desc->_type == library_data_t::real_float)) {
+ throw std::runtime_error(
+ "dpct::blas_gemm::experimental::matmul() only supports data type "
+ "combinataions:\n scale_type==int32 && a_type==int8 && b_type==int8"
+ "&& c_type==int32,\n scale_type==float && a_type==int8 && "
+ "b_type==int8 && c_type==int8,\n scale_type==float && a_type==int8"
+ "&& b_type==int8 && c_type==int32 or\n scale_type==float && "
+ "a_type==float"
+ "&& b_type==float && c_type==float.");
+ }
+
+ if (a_desc->_batch_count != b_desc->_batch_count ||
+ a_desc->_batch_count != c_desc->_batch_count ||
+ a_desc->_batch_count != d_desc->_batch_count) {
+ throw std::runtime_error("dpct::blas_gemm::experimental::matmul() does not "
+ "support different batch count.");
+ }
+
+#ifdef DPCT_USM_LEVEL_NONE
+ if (a_desc->_batch_count != 1) {
+ throw std::runtime_error(
+ "dpct::blas_gemm::experimental::matmul() only supports "
+ "batch mode with USM enabled.");
+ }
+#endif
+
+ size_t a_elm_size = dpct::detail::library_data_size[static_cast(
+ a_desc->_type)] /
+ 8;
+ size_t b_elm_size = dpct::detail::library_data_size[static_cast(
+ b_desc->_type)] /
+ 8;
+ size_t c_elm_size = dpct::detail::library_data_size[static_cast(
+ c_desc->_type)] /
+ 8;
+ size_t d_elm_size = dpct::detail::library_data_size[static_cast(
+ d_desc->_type)] /
+ 8;
+
+ static const auto matmul_single =
+ [](descriptor_ptr handle, matmul_desc_ptr compute_desc, const void *alpha,
+ const void *a, matrix_layout_ptr a_desc, const void *b,
+ matrix_layout_ptr b_desc, const void *beta, const void *c,
+ matrix_layout_ptr c_desc, void *d, matrix_layout_ptr d_desc,
+ ::dpct::cs::queue_ptr q_ptr, const size_t m, const size_t n,
+ const size_t k, const library_data_t a_type,
+ const library_data_t b_type, const library_data_t c_type,
+ const library_data_t d_type, const library_data_t scale_type,
+ bool vector_alpha, bool device_alpha, bool beta_is_zero,
+ size_t a_elm_size, size_t b_elm_size, size_t c_elm_size,
+ size_t d_elm_size) -> sycl::event {
+ // For non-col_major matrix, convert it to col_major.
+ const void *new_a = a;
+ const void *new_b = b;
+ const void *new_c = c;
+ void *new_d = d;
+ bool new_b_allocated = false;
+ bool new_c_allocated = false;
+ bool new_d_allocated = false;
+ size_t new_lda = a_desc->_ld, new_ldb = b_desc->_ld, new_ldc = c_desc->_ld,
+ new_ldd = d_desc->_ld;
+ std::vector transform_events;
+
+ if (a_desc->_order != order_t::col)
+ new_lda = a_desc->_rows;
+ new_a = ::dpct::cs::malloc(a_elm_size * a_desc->_cols * new_lda, *q_ptr);
+ sycl::event e_init;
+ if (a_desc->_order != order_t::col)
+ e_init = detail::type_dispatch(
+ a_desc->_type, q_ptr, a_desc->_rows, a_desc->_cols, a_desc->_ld,
+ a_desc->_order, (const std::int8_t *)a, new_lda, order_t::col,
+ (std::int8_t *)new_a, std::vector{});
+ else
+ e_init = ::dpct::cs::memcpy(
+ *q_ptr, (void *)new_a, a, a_elm_size * a_desc->_cols * new_lda,
+ ::dpct::cs::memcpy_direction::device_to_device);
+
+ // alpha = alpha * scale_a * scale_b
+ sycl::event e_scale_new_a = detail::scale_new_a(
+ q_ptr, m, k, (void *)new_a, a_type, alpha, scale_type, vector_alpha,
+ device_alpha, compute_desc->_a_scale_pointer,
+ compute_desc->_b_scale_pointer, {e_init});
+
+ transform_events.push_back(e_scale_new_a);
+
+ if (b_desc->_order != order_t::col) {
+ new_ldb = b_desc->_rows;
+ new_b = ::dpct::cs::malloc(b_elm_size * b_desc->_cols * new_ldb, *q_ptr);
+ new_b_allocated = true;
+ sycl::event e = detail::type_dispatch(
+ b_desc->_type, q_ptr, b_desc->_rows, b_desc->_cols, b_desc->_ld,
+ b_desc->_order, b, new_ldb, order_t::col, const_cast(new_b),
+ std::vector{});
+ transform_events.push_back(e);
+ }
+
+ if (!beta_is_zero && c_desc->_order != order_t::col) {
+ new_ldc = c_desc->_rows;
+ new_c = ::dpct::cs::malloc(c_elm_size * c_desc->_cols * new_ldc, *q_ptr);
+ new_c_allocated = true;
+ sycl::event e = detail::type_dispatch(
+ c_desc->_type, q_ptr, c_desc->_rows, c_desc->_cols, c_desc->_ld,
+ c_desc->_order, c, new_ldc, order_t::col, const_cast(new_c),
+ std::vector{});
+ transform_events.push_back(e);
+ }
+
+ if (d_desc->_order != order_t::col) {
+ new_ldd = d_desc->_rows;
+ new_d = ::dpct::cs::malloc(d_elm_size * d_desc->_cols * new_ldd, *q_ptr);
+ new_d_allocated = true;
+ }
+
+ // start to call oneDNN matmul primitive
+ // a,d are col_major, b is row_major
+ const ::dnnl::memory::dim M = m;
+ const ::dnnl::memory::dim N = n;
+ const ::dnnl::memory::dim K = k;
+
+ ::dnnl::memory::dims src_dims = {M, K};
+ ::dnnl::memory::dims weights_dims = {K, N};
+ ::dnnl::memory::dims bias_dims = {M, N};
+ ::dnnl::memory::dims dst_dims = {M, N};
+
+ const ::dnnl::memory::dims src_strides =
+ compute_desc->_trans_a == oneapi::mkl::transpose::nontrans
+ ? ::dnnl::memory::dims{1, static_cast(new_lda)}
+ : ::dnnl::memory::dims{static_cast(new_lda), 1};
+ const ::dnnl::memory::dims weights_strides =
+ compute_desc->_trans_b == oneapi::mkl::transpose::nontrans
+ ? ::dnnl::memory::dims{1, static_cast(new_ldb)}
+ : ::dnnl::memory::dims{static_cast(new_ldb), 1};
+ const ::dnnl::memory::dims bias_strides =
+ ::dnnl::memory::dims{1, static_cast(new_ldc)};
+ const ::dnnl::memory::dims dst_strides =
+ ::dnnl::memory::dims{1, static_cast(new_ldd)};
+
+ auto src_md = ::dnnl::memory::desc(
+ src_dims, dpct::dnnl::memory_desc_ext::to_dnnl_data_type(a_type),
+ src_strides);
+ auto weights_md = ::dnnl::memory::desc(
+ weights_dims, dpct::dnnl::memory_desc_ext::to_dnnl_data_type(b_type),
+ weights_strides);
+ auto bias_md = ::dnnl::memory::desc(
+ bias_dims, dpct::dnnl::memory_desc_ext::to_dnnl_data_type(c_type),
+ bias_strides);
+ auto dst_md = ::dnnl::memory::desc(
+ dst_dims, dpct::dnnl::memory_desc_ext::to_dnnl_data_type(d_type),
+ dst_strides);
+
+ auto *src_mem =
+ new ::dnnl::memory(src_md, handle->get_engine(), DNNL_MEMORY_NONE);
+ auto *weights_mem =
+ new ::dnnl::memory(weights_md, handle->get_engine(), DNNL_MEMORY_NONE);
+ auto *bias_mem =
+ new ::dnnl::memory(bias_md, handle->get_engine(), DNNL_MEMORY_NONE);
+ auto *dst_mem =
+ new ::dnnl::memory(dst_md, handle->get_engine(), DNNL_MEMORY_NONE);
+
+#ifdef DPCT_USM_LEVEL_NONE
+ detail::type_dispatch(a_type, src_mem, new_a);
+ detail::type_dispatch(b_type, weights_mem, new_b);
+ if (!beta_is_zero)
+ detail::type_dispatch(c_type, bias_mem, new_c);
+ detail::type_dispatch(d_type, dst_mem, new_d);
+#else
+ src_mem->set_data_handle(const_cast(new_a));
+ weights_mem->set_data_handle(const_cast(new_b));
+ if (!beta_is_zero)
+ bias_mem->set_data_handle(const_cast(new_c));
+ dst_mem->set_data_handle(new_d);
+#endif
+
+ std::unordered_map matmul_args;
+ matmul_args.insert({DNNL_ARG_SRC, *src_mem});
+ matmul_args.insert({DNNL_ARG_WEIGHTS, *weights_mem});
+ matmul_args.insert({DNNL_ARG_DST, *dst_mem});
+ ::dnnl::primitive_attr matmul_attr;
+
+ ::dnnl::post_ops matmul_ops;
+ if (!beta_is_zero) {
+ matmul_ops.append_binary(::dnnl::algorithm::binary_add, bias_md);
+ matmul_args.insert({DNNL_ARG_ATTR_MULTIPLE_POST_OP(matmul_ops.len() - 1) |
+ DNNL_ARG_SRC_1,
+ *bias_mem});
+ }
+
+ ::dnnl::memory *po_bias_mem = nullptr;
+ auto po_bias_md =
+ ::dnnl::memory::desc(::dnnl::memory::dims{M, 1},
+ dpct::dnnl::memory_desc_ext::to_dnnl_data_type(
+ compute_desc->_bias_data_type),
+ ::dnnl::memory::dims{1, M});
+ if (compute_desc->_epilogue == epilogue_t::bias ||
+ compute_desc->_epilogue == epilogue_t::gelu_bias ||
+ compute_desc->_epilogue == epilogue_t::gelu_aux_bias) {
+ po_bias_mem = new ::dnnl::memory(po_bias_md, handle->get_engine(),
+ DNNL_MEMORY_NONE);
+#ifdef DPCT_USM_LEVEL_NONE
+ detail::type_dispatch(
+ compute_desc->_bias_data_type, po_bias_mem,
+ compute_desc->_bias_pointer);
+#else
+ po_bias_mem->set_data_handle(compute_desc->_bias_pointer);
+#endif
+ }
+
+ ::dnnl::memory *po_bias_bgradb_mem = nullptr;
+ auto po_bias_bgradb_md = ::dnnl::memory::desc(
+ compute_desc->_trans_b == oneapi::mkl::transpose::nontrans
+ ? ::dnnl::memory::dims{N, 1}
+ : ::dnnl::memory::dims{1, N},
+ dpct::dnnl::memory_desc_ext::to_dnnl_data_type(
+ compute_desc->_bias_data_type),
+ compute_desc->_trans_b == oneapi::mkl::transpose::nontrans
+ ? ::dnnl::memory::dims{1, N}
+ : ::dnnl::memory::dims{N, 1});
+ if (compute_desc->_epilogue == epilogue_t::bgradb) {
+ po_bias_bgradb_mem = new ::dnnl::memory(
+ po_bias_bgradb_md, handle->get_engine(), DNNL_MEMORY_NONE);
+#ifdef DPCT_USM_LEVEL_NONE
+ detail::type_dispatch(
+ compute_desc->_bias_data_type, po_bias_bgradb_mem,
+ compute_desc->_bias_pointer);
+#else
+ po_bias_bgradb_mem->set_data_handle(compute_desc->_bias_pointer);
+#endif
+ }
+
+ ::dnnl::memory *po_aux_mem = nullptr;
+ auto po_aux_md = ::dnnl::memory::desc(
+ ::dnnl::memory::dims{M, N},
+ dpct::dnnl::memory_desc_ext::to_dnnl_data_type(
+ compute_desc->_epilogue_aux_data_type),
+ ::dnnl::memory::dims{
+ 1, static_cast(compute_desc->_epilogue_aux_ld)});
+ if (compute_desc->_epilogue == epilogue_t::dgelu) {
+ po_aux_mem =
+ new ::dnnl::memory(po_aux_md, handle->get_engine(), DNNL_MEMORY_NONE);
+#ifdef DPCT_USM_LEVEL_NONE
+ detail::type_dispatch(
+ compute_desc->_epilogue_aux_data_type, po_aux_mem,
+ compute_desc->_epilogue_aux_pointer);
+#else
+ po_aux_mem->set_data_handle(compute_desc->_epilogue_aux_pointer);
+#endif
+ }
+
+ switch (compute_desc->_epilogue) {
+ case epilogue_t::relu:
+ matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_relu, 0.f, 0.f);
+ break;
+ case epilogue_t::bias:
+ case epilogue_t::gelu_bias: {
+ matmul_ops.append_binary(::dnnl::algorithm::binary_add, po_bias_md);
+ matmul_args.insert({DNNL_ARG_ATTR_MULTIPLE_POST_OP(matmul_ops.len() - 1) |
+ DNNL_ARG_SRC_1,
+ *po_bias_mem});
+ if (compute_desc->_epilogue == epilogue_t::gelu_bias)
+ matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_tanh, 0.f,
+ 0.f);
+ break;
+ }
+ case epilogue_t::gelu:
+ matmul_ops.append_eltwise(::dnnl::algorithm::eltwise_gelu_tanh, 0.f, 0.f);
+ break;
+ default:
+ break;
+ }
+
+ matmul_attr.set_post_ops(matmul_ops);
+
+ auto matmul_pd = ::dnnl::matmul::primitive_desc(
+ handle->get_engine(), src_md, weights_md, dst_md, matmul_attr);
+ auto matmul_prim = ::dnnl::matmul(matmul_pd);
+ sycl::event matmul_prim_event =
+ ::dnnl::sycl_interop::execute(matmul_prim, handle->get_engine_stream(),
+ matmul_args, transform_events);
+
+ // post-op implemented by separate primitives
+ sycl::event post_op_prim_event;
+ if (compute_desc->_epilogue == epilogue_t::gelu_aux ||
+ compute_desc->_epilogue == epilogue_t::gelu_aux_bias) {
+ sycl::event prev_event = matmul_prim_event;
+ if (compute_desc->_epilogue == epilogue_t::gelu_aux_bias) {
+ auto po_bias_pd = ::dnnl::binary::primitive_desc(
+ handle->get_engine(), ::dnnl::algorithm::binary_add, dst_md,
+ po_bias_md, dst_md);
+ auto po_bias_prim = ::dnnl::binary(po_bias_pd);
+ std::unordered_map po_bias_args;
+ po_bias_args.insert({DNNL_ARG_SRC_0, *dst_mem});
+ po_bias_args.insert({DNNL_ARG_SRC_1, *po_bias_mem});
+ po_bias_args.insert({DNNL_ARG_DST, *dst_mem});
+ sycl::event prev_event = ::dnnl::sycl_interop::execute(
+ po_bias_prim, handle->get_engine_stream(), po_bias_args,
+ {matmul_prim_event});
+ }
+ size_t size_of_element =
+ dpct::detail::library_data_size[static_cast(
+ d_desc->_type)] /
+ 8;
+ sycl::event copy_e = dpct::blas::matrix_mem_copy_async(
+ compute_desc->_epilogue_aux_pointer, new_d,
+ compute_desc->_epilogue_aux_ld, d_desc->_ld, m, n, size_of_element,
+ ::dpct::cs::memcpy_direction::automatic, *q_ptr, {prev_event});
+
+ auto gelu_pd = ::dnnl::eltwise_forward::primitive_desc(
+ handle->get_engine(), ::dnnl::prop_kind::forward_training,
+ ::dnnl::algorithm::eltwise_gelu_tanh, dst_md, dst_md);
+ auto gelu_prim = ::dnnl::eltwise_forward(gelu_pd);
+ std::unordered_map gelu_args;
+ gelu_args.insert({DNNL_ARG_SRC, *dst_mem});
+ gelu_args.insert({DNNL_ARG_DST, *dst_mem});
+ post_op_prim_event = ::dnnl::sycl_interop::execute(
+ gelu_prim, handle->get_engine_stream(), gelu_args, {copy_e});
+ } else if (compute_desc->_epilogue == epilogue_t::dgelu) {
+ auto gelu_pd = ::dnnl::eltwise_forward::primitive_desc(
+ handle->get_engine(), ::dnnl::prop_kind::forward_training,
+ ::dnnl::algorithm::eltwise_gelu_tanh, po_aux_md, po_aux_md);
+ auto dgelu_pd = ::dnnl::eltwise_backward::primitive_desc(
+ handle->get_engine(), ::dnnl::algorithm::eltwise_gelu_tanh, dst_md,
+ dst_md, po_aux_md, gelu_pd);
+ auto dgelu_prim = ::dnnl::eltwise_backward(dgelu_pd);
+ std::unordered_map dgelu_args;
+ dgelu_args.insert({DNNL_ARG_SRC, *po_aux_mem});
+ dgelu_args.insert({DNNL_ARG_DIFF_DST, *dst_mem});
+ dgelu_args.insert({DNNL_ARG_DIFF_SRC, *dst_mem});
+ post_op_prim_event =
+ ::dnnl::sycl_interop::execute(dgelu_prim, handle->get_engine_stream(),
+ dgelu_args, {matmul_prim_event});
+ } else if (compute_desc->_epilogue == epilogue_t::bgradb) {
+ auto reduction_pd = ::dnnl::reduction::primitive_desc(
+ handle->get_engine(), ::dnnl::algorithm::reduction_sum, weights_md,
+ po_bias_bgradb_md, 0.f, 0.f);
+ auto reduction_prim = ::dnnl::reduction(reduction_pd);
+ std::unordered_map reduction_args;
+ reduction_args.insert({DNNL_ARG_SRC, *weights_mem});
+ reduction_args.insert({DNNL_ARG_DST, *po_bias_bgradb_mem});
+ post_op_prim_event = ::dnnl::sycl_interop::execute(
+ reduction_prim, handle->get_engine_stream(), reduction_args,
+ {matmul_prim_event});
+ }
+
+#ifndef DPCT_USM_LEVEL_NONE
+ // WA to avoid incorrect result on CPU with USM
+ if (q_ptr->get_device().is_cpu()) {
+ matmul_prim_event.wait();
+ post_op_prim_event.wait();
+ }
+#endif
+ // end of calling oneDNN
+
+ sycl::event absmax_d_event;
+ if (auto absmax_ptr = compute_desc->_absmax_d_pointer) {
+ absmax_d_event = detail::type_dispatch(
+ d_desc->_type, absmax_ptr, new_d, new_ldd, d_desc->_rows,
+ d_desc->_cols, q_ptr,
+ std::vector{matmul_prim_event, post_op_prim_event});
+ }
+
+ sycl::event scale_d_event;
+ if (auto d_scale_ptr = compute_desc->_d_scale_pointer) {
+ scale_d_event = detail::type_dispatch(
+ d_desc->_type, d_scale_ptr, new_d, new_ldd, d_desc->_rows,
+ d_desc->_cols, q_ptr, compute_desc->_scale_type,
+ std::vector{matmul_prim_event, absmax_d_event,
+ post_op_prim_event});
+ }
+
+ sycl::event transform_d_event;
+ if (d_desc->_order != order_t::col) {
+ transform_d_event = detail::type_dispatch(
+ d_desc->_type, q_ptr, d_desc->_rows, d_desc->_cols, new_ldd,
+ order_t::col, new_d, d_desc->_ld, d_desc->_order, d,
+ std::vector{matmul_prim_event, absmax_d_event,
+ post_op_prim_event});
+ }
+
+ sycl::event free_event = q_ptr->submit([&](sycl::handler &cgh) {
+ cgh.depends_on({transform_d_event, matmul_prim_event, absmax_d_event,
+ post_op_prim_event});
+ cgh.host_task([=] {
+ delete src_mem;
+ delete weights_mem;
+ delete bias_mem;
+ delete dst_mem;
+ if (po_bias_mem)
+ delete po_bias_mem;
+ if (po_bias_bgradb_mem)
+ delete po_bias_bgradb_mem;
+ if (po_aux_mem)
+ delete po_aux_mem;
+ ::dpct::cs::free((void *)new_a, *q_ptr);
+ if (new_b_allocated)
+ ::dpct::cs::free((void *)new_b, *q_ptr);
+ if (new_c_allocated)
+ ::dpct::cs::free((void *)new_c, *q_ptr);
+ if (new_d_allocated)
+ ::dpct::cs::free((void *)new_d, *q_ptr);
+ });
+ });
+ return free_event;
+ };
+
+ std::vector events;
+ const std::byte *offsetted_a = static_cast(a);
+ const std::byte *offsetted_b = static_cast(b);
+ const std::byte *offsetted_c = static_cast(c);
+ std::byte *offsetted_d = static_cast(d);
+ for (std::uint64_t i = 0; i < a_desc->_batch_count; i++) {
+ sycl::event e = matmul_single(
+ handle, compute_desc, alpha, offsetted_a, a_desc, offsetted_b, b_desc,
+ beta, offsetted_c, c_desc, offsetted_d, d_desc, q_ptr, m, n, k, a_type,
+ b_type, c_type, d_type, scale_type, vector_alpha, device_alpha,
+ beta_is_zero, a_elm_size, b_elm_size, c_elm_size, d_elm_size);
+ events.push_back(e);
+
+ offsetted_a += a_elm_size * a_desc->_strided_batch_offset;
+ offsetted_b += b_elm_size * b_desc->_strided_batch_offset;
+ offsetted_c += c_elm_size * c_desc->_strided_batch_offset;
+ offsetted_d += d_elm_size * d_desc->_strided_batch_offset;
+ }
+
+ return q_ptr->single_task(events, [] {});
+}
+
+class transform_desc_t {
+public:
+ enum class attribute { scale_type, pointer_mode, trans_a, trans_b };
+
+ transform_desc_t(library_data_t scale_type) : _scale_type(scale_type) {}
+ void set_attribute(attribute attr, const void *mem) {
+ get_set_attr(attr, const_cast(mem));
+ }
+ void get_attribute(attribute attr, void *mem) {
+ get_set_attr(attr, mem);
+ }
+
+private:
+ template