From 4148c19e022bd0809fe3c19ff4b1ba0a37a908c3 Mon Sep 17 00:00:00 2001 From: vchamarthi Date: Wed, 9 Sep 2026 09:24:35 -0500 Subject: [PATCH 1/3] add asv benchmarks --- benchmarks/README.md | 48 +++++++ benchmarks/asv.conf.json | 25 ++++ benchmarks/benchmarks/__init__.py | 23 +++ benchmarks/benchmarks/_utils.py | 159 +++++++++++++++++++++ benchmarks/benchmarks/bench_compile.py | 133 +++++++++++++++++ benchmarks/benchmarks/bench_construct.py | 119 +++++++++++++++ benchmarks/benchmarks/bench_copy.py | 110 ++++++++++++++ benchmarks/benchmarks/bench_enumerate.py | 75 ++++++++++ benchmarks/benchmarks/bench_queue_cache.py | 55 +++++++ benchmarks/benchmarks/bench_submit.py | 88 ++++++++++++ benchmarks/benchmarks/bench_usm.py | 96 +++++++++++++ benchmarks/requirements.txt | 1 + pyproject.toml | 1 + 13 files changed, 933 insertions(+) create mode 100644 benchmarks/README.md create mode 100644 benchmarks/asv.conf.json create mode 100644 benchmarks/benchmarks/__init__.py create mode 100644 benchmarks/benchmarks/_utils.py create mode 100644 benchmarks/benchmarks/bench_compile.py create mode 100644 benchmarks/benchmarks/bench_construct.py create mode 100644 benchmarks/benchmarks/bench_copy.py create mode 100644 benchmarks/benchmarks/bench_enumerate.py create mode 100644 benchmarks/benchmarks/bench_queue_cache.py create mode 100644 benchmarks/benchmarks/bench_submit.py create mode 100644 benchmarks/benchmarks/bench_usm.py create mode 100644 benchmarks/requirements.txt diff --git a/benchmarks/README.md b/benchmarks/README.md new file mode 100644 index 0000000000..1d6170adcb --- /dev/null +++ b/benchmarks/README.md @@ -0,0 +1,48 @@ +# dpctl ASV Benchmarks + +Runtime-overhead benchmarks for [dpctl](https://github.com/IntelPython/dpctl) +using [ASV](https://asv.readthedocs.io/en/stable/): object construction, +queue caching, USM allocation, device enumeration, kernel bundle +compilation, data movement, kernel submission. No compute throughput. + +## Coverage + +| File | API | +|------|-----| +| `bench_construct.py` | `SyclDevice`/`SyclContext`/`SyclQueue`/`SyclPlatform` construction, device/queue attribute reads | +| `bench_queue_cache.py` | `get_device_cached_queue` for each key kind, vs. an uncached baseline | +| `bench_usm.py` | `MemoryUSM{Device,Host,Shared}` alloc/free, aligned alloc, first touch, USM pointer queries | +| `bench_enumerate.py` | `get_devices`, `get_num_devices`, `get_platforms`, `select_*_device`, `has_*_devices`, `select_device_with_aspects` | +| `bench_compile.py` | Kernel bundles from SPIR-V / OpenCL C source / SYCL source, kernel lookup, availability probes | +| `bench_copy.py` | `SyclQueue.memcpy`/`memcpy_async`/`fill`/`memset`, `_Memory.copy_to_host`/`copy_from_host`/`copy_from_device` | +| `bench_submit.py` | `submit`, `submit_async`, batched submission, `submit_barrier`, idle `wait` | + +## Device axis + +Benchmarks parameterize over the `cpu`/`gpu` filter selectors and skip when a +selector has no matching device, so the same benchmark names run on any +node. Sizes over 25% of a device's `global_mem_size` skip the same way. + +## Compilation caching + +`benchmarks/__init__.py` disables the persistent JIT cache +(`SYCL_CACHE_PERSISTENT=0`). `time_bundle_from_source_cold` mints a unique +kernel name per call to defeat the in-memory cache too; `_warm` reuses one +name to measure the cache-hit path. + +`create_kernel_bundle_from_source` runs on the OpenCL backend only. +`create_kernel_bundle_from_sycl_source` needs a device where +`can_compile("sycl")` is true. + +## Running + +```bash +pip install ".[benchmark]" +cd benchmarks && asv machine --yes && asv run --python=same --quick HEAD^! +``` + +One module: `asv run --python=same --quick --bench bench_compile HEAD^!` + +Compare commits: `asv continuous --python=same HEAD~1 HEAD` + +View results: `asv publish && asv preview` diff --git a/benchmarks/asv.conf.json b/benchmarks/asv.conf.json new file mode 100644 index 0000000000..e57e1542ee --- /dev/null +++ b/benchmarks/asv.conf.json @@ -0,0 +1,25 @@ +{ + "version": 1, + "project": "dpctl", + "project_url": "https://github.com/IntelPython/dpctl", + "show_commit_url": "https://github.com/IntelPython/dpctl/commit/", + "repo": "..", + "branches": [ + "master", + "dev-milestone" + ], + "environment_type": "conda", + "conda_channels": [ + "https://software.repos.intel.com/python/conda/", + "conda-forge" + ], + "benchmark_dir": "benchmarks", + "env_dir": ".asv/env", + "results_dir": ".asv/results", + "html_dir": ".asv/html", + "build_cache_size": 2, + "default_benchmark_timeout": 900, + "regressions_thresholds": { + ".*": 0.2 + } +} diff --git a/benchmarks/benchmarks/__init__.py b/benchmarks/benchmarks/__init__.py new file mode 100644 index 0000000000..903ce277fb --- /dev/null +++ b/benchmarks/benchmarks/__init__.py @@ -0,0 +1,23 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""ASV benchmarks for dpctl.""" + +import os + +# Disable the persistent JIT cache before dpctl is imported, so cold-compile +# benchmarks in bench_compile.py measure a real compile, not a cache hit. +os.environ.setdefault("SYCL_CACHE_PERSISTENT", "0") diff --git a/benchmarks/benchmarks/_utils.py b/benchmarks/benchmarks/_utils.py new file mode 100644 index 0000000000..8efcea5691 --- /dev/null +++ b/benchmarks/benchmarks/_utils.py @@ -0,0 +1,159 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Shared utilities for dpctl benchmarks. + +Every benchmark that needs a device goes through :func:`queue_for` or +:func:`device_for` so that benchmark names stay identical on every node in +the pool: a node without a given device reports the parameter as skipped +rather than failing the whole suite. +""" + +import os + +from asv_runner.benchmarks.mark import SkipNotImplemented + +import dpctl +import dpctl.compiler as dpc + +# Device axis. Filter-selector strings accepted by SyclQueue/SyclDevice. +_SELECTORS = ["cpu", "gpu"] + +# Allocation and transfer size sweep, in bytes: 4 KiB, 1 MiB, 16 MiB, 256 MiB. +_SIZES = [4 * 1024, 1024**2, 16 * 1024**2, 256 * 1024**2] + +# USM kinds, mapped to their constructors. +_USM_TYPES = ["device", "host", "shared"] + +# Fraction of device global memory a single benchmark allocation may claim. +_MEM_BUDGET = 0.25 + +_queues = {} +_devices = {} +_spirv = None + + +def queue_for(selector): + """Return a memoized queue for *selector*, or skip when unavailable.""" + if selector not in _queues: + try: + _queues[selector] = dpctl.SyclQueue(selector) + except dpctl.SyclQueueCreationError: + _queues[selector] = None + q = _queues[selector] + if q is None: + raise SkipNotImplemented(f"no {selector} device available") + return q + + +def device_for(selector): + """Return a memoized device for *selector*, or skip when unavailable.""" + if selector not in _devices: + try: + _devices[selector] = dpctl.SyclDevice(selector) + except dpctl.SyclDeviceCreationError: + _devices[selector] = None + d = _devices[selector] + if d is None: + raise SkipNotImplemented(f"no {selector} device available") + return d + + +def usm_class(usm_type): + """Return the dpctl.memory class allocating *usm_type* memory.""" + import dpctl.memory as dpm + + return { + "device": dpm.MemoryUSMDevice, + "host": dpm.MemoryUSMHost, + "shared": dpm.MemoryUSMShared, + }[usm_type] + + +def skip_unless_fits(queue, nbytes): + """Skip when *nbytes* exceeds this device's allocation budget.""" + budget = _MEM_BUDGET * queue.sycl_device.global_mem_size + if nbytes > budget: + raise SkipNotImplemented( + f"{nbytes} bytes exceeds the device memory budget" + ) + + +def opencl_queue_or_skip(): + """Return a memoized OpenCL queue, or skip. + + ``create_kernel_bundle_from_source`` only supports the OpenCL backend. + """ + return queue_for("opencl") + + +def sycl_source_queue_or_skip(selector): + """Return a queue whose device can compile SYCL source, or skip.""" + q = queue_for(selector) + if not dpc.is_sycl_source_compilation_available(): + raise SkipNotImplemented("SYCL source compilation extension absent") + if not q.sycl_device.can_compile("sycl"): + raise SkipNotImplemented("device cannot compile SYCL source") + return q + + +def spirv_bytes(): + """Return the SPIR-V module shipped with the installed dpctl, or skip. + + Defines ``add(int*, int*, int*)`` and ``axpy(int*, int*, int*, int)``. + """ + global _spirv + if _spirv is None: + path = os.path.join( + os.path.dirname(os.path.abspath(dpctl.__file__)), + "tests", + "input_files", + "multi_kernel.spv", + ) + if not os.path.exists(path): + raise SkipNotImplemented(f"SPIR-V module not found at {path}") + with open(path, "rb") as fh: + _spirv = fh.read() + return _spirv + + +def ocl_axpy_source(kernel_name="axpy"): + """Return OpenCL C source for an axpy kernel called *kernel_name*.""" + return ( + f"kernel void {kernel_name}(" + " global int *a, global int *b, global int *c, int d) {" + " size_t index = get_global_id(0);" + " c[index] = d * a[index] + b[index];" + "}" + ) + + +def sycl_axpy_source(kernel_name="axpy"): + """Return SYCL source for an axpy kernel called *kernel_name*.""" + return f""" + #include + + namespace syclext = sycl::ext::oneapi::experimental; + + extern "C" SYCL_EXTERNAL + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclext::nd_range_kernel<1>)) + void {kernel_name}(int* a, int* b, int* c, int d) {{ + sycl::nd_item<1> item = + sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + size_t i = item.get_global_linear_id(); + c[i] = d * a[i] + b[i]; + }} + """ diff --git a/benchmarks/benchmarks/bench_compile.py b/benchmarks/benchmarks/bench_compile.py new file mode 100644 index 0000000000..200dbb5a56 --- /dev/null +++ b/benchmarks/benchmarks/bench_compile.py @@ -0,0 +1,133 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for kernel bundle creation. + +Cold and warm compilation are separate benchmarks. Cold gives every call a +kernel name no compiler has seen, so neither the in-memory nor the +persistent cache (disabled in ``benchmarks/__init__.py``) can serve it; warm +re-submits identical source to measure the cache-hit path instead. +""" + +import itertools + +import dpctl.compiler as dpc + +from ._utils import ( + _SELECTORS, + ocl_axpy_source, + opencl_queue_or_skip, + queue_for, + spirv_bytes, + sycl_axpy_source, + sycl_source_queue_or_skip, +) + + +class BundleFromSPIRV: + """Kernel bundle from a pre-compiled SPIR-V module.""" + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.queue = queue_for(selector) + self.spirv = spirv_bytes() + self.bundle = dpc.create_kernel_bundle_from_spirv( + self.queue, self.spirv + ) + + def time_bundle_from_spirv(self, selector): + dpc.create_kernel_bundle_from_spirv(self.queue, self.spirv) + + def time_get_sycl_kernel(self, selector): + self.bundle.get_sycl_kernel("axpy") + + def time_has_sycl_kernel(self, selector): + self.bundle.has_sycl_kernel("axpy") + + +class BundleFromOpenCLSource: + """Kernel bundle built from OpenCL C source (OpenCL backend only).""" + + timeout = 300 + number = 1 + repeat = 3 + warmup_time = 0 + + def setup(self): + self.queue = opencl_queue_or_skip() + self.counter = itertools.count() + self.warm_source = ocl_axpy_source() + dpc.create_kernel_bundle_from_source(self.queue, self.warm_source) + + def time_bundle_from_source_cold(self): + name = f"axpy_{next(self.counter)}" + dpc.create_kernel_bundle_from_source(self.queue, ocl_axpy_source(name)) + + def time_bundle_from_source_warm(self): + dpc.create_kernel_bundle_from_source(self.queue, self.warm_source) + + +class BundleFromSYCLSource: + """Kernel bundle built from SYCL source via the kernel_compiler + extension. + + Skipped unless the extension is present and the device reports it can + compile SYCL source. + """ + + params = [_SELECTORS] + param_names = ["selector"] + timeout = 600 + number = 1 + repeat = 2 + warmup_time = 0 + + def setup(self, selector): + self.queue = sycl_source_queue_or_skip(selector) + self.counter = itertools.count() + self.warm_source = sycl_axpy_source() + dpc.create_kernel_bundle_from_sycl_source(self.queue, self.warm_source) + + def time_bundle_from_sycl_source_cold(self, selector): + name = f"axpy_{next(self.counter)}" + dpc.create_kernel_bundle_from_sycl_source( + self.queue, sycl_axpy_source(name) + ) + + def time_bundle_from_sycl_source_warm(self, selector): + dpc.create_kernel_bundle_from_sycl_source(self.queue, self.warm_source) + + +class SourceCompilationProbe: + """Cost of the availability probes themselves. + + Called by consumers before every compilation attempt. + """ + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.device = queue_for(selector).sycl_device + dpc.is_sycl_source_compilation_available() + + def time_is_sycl_source_compilation_available(self, selector): + dpc.is_sycl_source_compilation_available() + + def time_can_compile(self, selector): + self.device.can_compile("sycl") diff --git a/benchmarks/benchmarks/bench_construct.py b/benchmarks/benchmarks/bench_construct.py new file mode 100644 index 0000000000..cd66cb667a --- /dev/null +++ b/benchmarks/benchmarks/bench_construct.py @@ -0,0 +1,119 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for construction of the core dpctl runtime objects.""" + +import dpctl + +from ._utils import _SELECTORS, device_for, queue_for + + +class Construct: + """Construction cost of SyclDevice, SyclContext, SyclQueue. + + ``SyclQueue(selector)`` also builds a context, while + ``SyclQueue(context, device)`` reuses one. Both are measured so that a + regression can be attributed to the queue or to the context. + """ + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.device = device_for(selector) + self.context = dpctl.SyclContext(self.device) + + def time_sycl_device(self, selector): + dpctl.SyclDevice(selector) + + def time_sycl_context_from_device(self, selector): + dpctl.SyclContext(self.device) + + def time_sycl_queue_from_selector(self, selector): + dpctl.SyclQueue(selector) + + def time_sycl_queue_from_device(self, selector): + dpctl.SyclQueue(self.device) + + def time_sycl_queue_from_context_device(self, selector): + dpctl.SyclQueue(self.context, self.device) + + def time_sycl_queue_in_order(self, selector): + dpctl.SyclQueue(self.device, property="in_order") + + +class ConstructPlatform: + """Construction and enumeration cost of SyclPlatform.""" + + def time_sycl_platform_default(self): + dpctl.SyclPlatform() + + def time_get_platforms(self): + dpctl.get_platforms() + + +class DeviceProperties: + """Hot device attribute reads. + + Catches an attribute that stops being cached and starts round-tripping + to the SYCL runtime on every access. + """ + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.device = device_for(selector) + # touch every attribute once so a first-access cost is not charged + # to the first measured iteration + self.time_device_property_reads(selector) + + def time_device_property_reads(self, selector): + d = self.device + d.name + d.driver_version + d.max_compute_units + d.max_work_group_size + d.global_mem_size + d.has_aspect_fp64 + + def time_device_filter_string(self, selector): + self.device.filter_string + + +class QueueAccessors: + """Accessor cost on an existing queue.""" + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.queue = queue_for(selector) + + def time_sycl_device_attr(self, selector): + self.queue.sycl_device + + def time_get_sycl_device(self, selector): + self.queue.get_sycl_device() + + def time_sycl_context_attr(self, selector): + self.queue.sycl_context + + def time_queue_hash(self, selector): + hash(self.queue) + + def time_queue_eq(self, selector): + self.queue == self.queue diff --git a/benchmarks/benchmarks/bench_copy.py b/benchmarks/benchmarks/bench_copy.py new file mode 100644 index 0000000000..a110cfb854 --- /dev/null +++ b/benchmarks/benchmarks/bench_copy.py @@ -0,0 +1,110 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for data movement through a queue and through _Memory. + +Every asynchronous benchmark waits before returning, so no benchmark leaves +work queued for the next iteration to absorb. +""" + +import numpy as np + +import dpctl.memory as dpm + +from ._utils import _SELECTORS, _SIZES, queue_for, skip_unless_fits + + +class HostDeviceCopy: + """Blocking and asynchronous host/device transfers.""" + + params = [_SELECTORS, _SIZES] + param_names = ["selector", "nbytes"] + + def setup(self, selector, nbytes): + self.queue = queue_for(selector) + skip_unless_fits(self.queue, 2 * nbytes) + self.nbytes = nbytes + self.host = np.zeros(nbytes, dtype="u1") + self.dev = dpm.MemoryUSMDevice(nbytes, queue=self.queue) + # page in both sides before measuring + self.queue.memcpy(self.dev, self.host, nbytes) + self.queue.memcpy(self.host, self.dev, nbytes) + + def time_memcpy_h2d(self, selector, nbytes): + self.queue.memcpy(self.dev, self.host, self.nbytes) + + def time_memcpy_d2h(self, selector, nbytes): + self.queue.memcpy(self.host, self.dev, self.nbytes) + + def time_memcpy_async_h2d_wait(self, selector, nbytes): + self.queue.memcpy_async(self.dev, self.host, self.nbytes).wait() + + def time_copy_from_host(self, selector, nbytes): + self.dev.copy_from_host(self.host) + + def time_copy_to_host(self, selector, nbytes): + self.dev.copy_to_host(self.host) + + +class DeviceDeviceCopy: + """Device-to-device transfers.""" + + params = [_SELECTORS, _SIZES] + param_names = ["selector", "nbytes"] + + def setup(self, selector, nbytes): + self.queue = queue_for(selector) + skip_unless_fits(self.queue, 2 * nbytes) + self.nbytes = nbytes + self.src = dpm.MemoryUSMDevice(nbytes, queue=self.queue) + self.dst = dpm.MemoryUSMDevice(nbytes, queue=self.queue) + self.src.memset() + self.queue.memcpy(self.dst, self.src, nbytes) + + def time_memcpy_d2d(self, selector, nbytes): + self.queue.memcpy(self.dst, self.src, self.nbytes) + + def time_copy_from_device(self, selector, nbytes): + self.dst.copy_from_device(self.src) + + +class Fill: + """Byte-wise and typed fills.""" + + params = [_SELECTORS, _SIZES] + param_names = ["selector", "nbytes"] + + def setup(self, selector, nbytes): + self.queue = queue_for(selector) + skip_unless_fits(self.queue, nbytes) + self.nbytes = nbytes + self.dev = dpm.MemoryUSMDevice(nbytes, queue=self.queue) + self.dev.memset() + + def time_memset(self, selector, nbytes): + self.queue.memset(self.dev, 0, self.nbytes) + + def time_memory_memset(self, selector, nbytes): + self.dev.memset() + + def time_fill_u1(self, selector, nbytes): + self.queue.fill(self.dev, 0, self.nbytes, "u1") + + def time_fill_f4(self, selector, nbytes): + self.queue.fill(self.dev, 0.0, self.nbytes // 4, "f4") + + def time_memset_async_wait(self, selector, nbytes): + self.queue.memset_async(self.dev, 0, self.nbytes).wait() diff --git a/benchmarks/benchmarks/bench_enumerate.py b/benchmarks/benchmarks/bench_enumerate.py new file mode 100644 index 0000000000..7aceb42936 --- /dev/null +++ b/benchmarks/benchmarks/bench_enumerate.py @@ -0,0 +1,75 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for device enumeration and device selection.""" + +from asv_runner.benchmarks.mark import SkipNotImplemented + +import dpctl + +_DEVICE_TYPES = ["all", "cpu", "gpu"] +_BACKENDS = ["all", "opencl", "level_zero"] + + +class Enumerate: + """get_devices / get_num_devices across device_type and backend.""" + + params = [_DEVICE_TYPES, _BACKENDS] + param_names = ["device_type", "backend"] + + def setup(self, device_type, backend): + # first enumeration initializes the SYCL platform list + dpctl.get_devices(backend=backend, device_type=device_type) + + def time_get_devices(self, device_type, backend): + dpctl.get_devices(backend=backend, device_type=device_type) + + def time_get_num_devices(self, device_type, backend): + dpctl.get_num_devices(backend=backend, device_type=device_type) + + +class Select: + """Device selectors and availability predicates.""" + + def setup(self): + dpctl.select_default_device() + try: + dpctl.select_device_with_aspects("fp64") + self.has_fp64 = True + except dpctl.SyclDeviceCreationError: + self.has_fp64 = False + + def time_select_default_device(self): + dpctl.select_default_device() + + def time_select_cpu_device(self): + if dpctl.has_cpu_devices(): + dpctl.select_cpu_device() + + def time_select_gpu_device(self): + if dpctl.has_gpu_devices(): + dpctl.select_gpu_device() + + def time_has_cpu_devices(self): + dpctl.has_cpu_devices() + + def time_has_gpu_devices(self): + dpctl.has_gpu_devices() + + def time_select_device_with_aspects(self): + if not self.has_fp64: + raise SkipNotImplemented("no device with the fp64 aspect") + dpctl.select_device_with_aspects("fp64") diff --git a/benchmarks/benchmarks/bench_queue_cache.py b/benchmarks/benchmarks/bench_queue_cache.py new file mode 100644 index 0000000000..962f996a32 --- /dev/null +++ b/benchmarks/benchmarks/bench_queue_cache.py @@ -0,0 +1,55 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for the cached-queue lookup in dpctl._sycl_queue_manager. + +``get_device_cached_queue`` accepts three key kinds and they do not cost the +same: only the ``(SyclContext, SyclDevice)`` tuple key reaches the map +without building a queue first, while the device key and the filter-string +key each construct a ``SyclQueue`` before the lookup. All four paths are +tracked so the asymmetry stays visible. +""" + +import dpctl + +from ._utils import _SELECTORS, device_for + + +class QueueCache: + """Cached queue lookup, one benchmark per key kind.""" + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.selector = selector + self.device = device_for(selector) + # populate the cache so every measurement below is a hit + q = dpctl.get_device_cached_queue(self.device) + self.context = q.sycl_context + self.ctx_dev = (self.context, self.device) + + def time_cached_queue_ctx_dev_key(self, selector): + dpctl.get_device_cached_queue(self.ctx_dev) + + def time_cached_queue_device_key(self, selector): + dpctl.get_device_cached_queue(self.device) + + def time_cached_queue_str_key(self, selector): + dpctl.get_device_cached_queue(selector) + + def time_uncached_queue(self, selector): + dpctl.SyclQueue(selector) diff --git a/benchmarks/benchmarks/bench_submit.py b/benchmarks/benchmarks/bench_submit.py new file mode 100644 index 0000000000..c1b02c21c3 --- /dev/null +++ b/benchmarks/benchmarks/bench_submit.py @@ -0,0 +1,88 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for kernel submission and synchronization overhead. + +The kernel is deliberately trivial and the global range is tiny, so what is +measured is dpctl's submission path and the runtime round trip, not compute. +""" + +import dpctl +import dpctl.compiler as dpc +import dpctl.memory as dpm + +from ._utils import _SELECTORS, queue_for, spirv_bytes + +_BATCH = 100 + + +class Submit: + """Submission overhead on a one-work-item range.""" + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.queue = queue_for(selector) + bundle = dpc.create_kernel_bundle_from_spirv( + self.queue, spirv_bytes() + ) + self.kernel = bundle.get_sycl_kernel("add") + nbytes = 4 * 1024 + self.args = [ + dpm.MemoryUSMDevice(nbytes, queue=self.queue) for _ in range(3) + ] + for m in self.args: + m.memset() + self.range = [1] + # first launch on this queue builds runtime state; do not charge it + # to the first measured iteration + self.queue.submit(self.kernel, self.args, self.range) + + def time_submit_roundtrip(self, selector): + self.queue.submit(self.kernel, self.args, self.range) + + def time_submit_async_wait(self, selector): + self.queue.submit_async(self.kernel, self.args, self.range).wait() + + def time_submit_async_batch(self, selector): + q = self.queue + for _ in range(_BATCH): + ev = q.submit_async(self.kernel, self.args, self.range) + ev.wait() + + def time_submit_with_local_range(self, selector): + self.queue.submit(self.kernel, self.args, [1], [1]) + + +class Synchronize: + """Barrier, queue wait, and event handling costs on an idle queue.""" + + params = [_SELECTORS] + param_names = ["selector"] + + def setup(self, selector): + self.queue = queue_for(selector) + self.queue.submit_barrier().wait() + + def time_submit_barrier_wait(self, selector): + self.queue.submit_barrier().wait() + + def time_queue_wait_idle(self, selector): + self.queue.wait() + + def time_default_event_wait(self, selector): + dpctl.SyclEvent().wait() diff --git a/benchmarks/benchmarks/bench_usm.py b/benchmarks/benchmarks/bench_usm.py new file mode 100644 index 0000000000..2336d1b3d4 --- /dev/null +++ b/benchmarks/benchmarks/bench_usm.py @@ -0,0 +1,96 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Benchmarks for USM allocation and USM pointer queries. + +Each timed call allocates and releases in the same call, so ASV iterations +cannot accumulate device memory. Sizes that do not fit the device budget are +skipped rather than left to fail the run. +""" + +from ._utils import ( + _SELECTORS, + _SIZES, + _USM_TYPES, + queue_for, + skip_unless_fits, + usm_class, +) + + +class USMAllocation: + """Allocate and free one USM block per timed call.""" + + params = [_SELECTORS, _USM_TYPES, _SIZES] + param_names = ["selector", "usm_type", "nbytes"] + + def setup(self, selector, usm_type, nbytes): + self.queue = queue_for(selector) + skip_unless_fits(self.queue, nbytes) + self.cls = usm_class(usm_type) + # first allocation of a given kind may initialize a runtime pool + self.cls(nbytes, queue=self.queue) + + def time_alloc_free(self, selector, usm_type, nbytes): + self.cls(nbytes, queue=self.queue) + + def time_alloc_free_aligned(self, selector, usm_type, nbytes): + self.cls(nbytes, alignment=4096, queue=self.queue) + + +class USMFirstTouch: + """Allocate, write one byte, free. + + Separates lazy allocation from the page-in that the first write pays. + """ + + params = [_SELECTORS, _USM_TYPES, [1024**2, 16 * 1024**2]] + param_names = ["selector", "usm_type", "nbytes"] + + def setup(self, selector, usm_type, nbytes): + self.queue = queue_for(selector) + skip_unless_fits(self.queue, nbytes) + self.cls = usm_class(usm_type) + m = self.cls(nbytes, queue=self.queue) + m.memset() + + def time_alloc_touch_free(self, selector, usm_type, nbytes): + m = self.cls(nbytes, queue=self.queue) + m.memset() + + +class USMQueries: + """Pointer and interface queries on an existing allocation.""" + + params = [_SELECTORS, _USM_TYPES] + param_names = ["selector", "usm_type"] + + def setup(self, selector, usm_type): + self.queue = queue_for(selector) + self.mem = usm_class(usm_type)(1024**2, queue=self.queue) + self.mem.get_usm_type() + + def time_get_usm_type(self, selector, usm_type): + self.mem.get_usm_type() + + def time_get_usm_type_enum(self, selector, usm_type): + self.mem.get_usm_type_enum() + + def time_sycl_usm_array_interface(self, selector, usm_type): + self.mem.__sycl_usm_array_interface__ + + def time_sycl_queue_attr(self, selector, usm_type): + self.mem.sycl_queue diff --git a/benchmarks/requirements.txt b/benchmarks/requirements.txt new file mode 100644 index 0000000000..24ce15ab7e --- /dev/null +++ b/benchmarks/requirements.txt @@ -0,0 +1 @@ +numpy diff --git a/pyproject.toml b/pyproject.toml index b1d88ddcfc..f7b2d68480 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -61,6 +61,7 @@ readme = {file = "README.md", content-type = "text/markdown"} requires-python = ">=3.10" [project.optional-dependencies] +benchmark = ["asv>=0.6", "numpy"] coverage = ["Cython>=3.1.0", "pytest", "coverage", "tomli"] docs = [ "Cython>=3.1.0", From 9a82d4b574fd8bcdc2b5fdf852c48b66816f32ac Mon Sep 17 00:00:00 2001 From: vchamarthi Date: Wed, 9 Sep 2026 18:15:04 -0500 Subject: [PATCH 2/3] failsafe guarding over unimplemented methods --- benchmarks/benchmarks/_utils.py | 11 ++++++++--- benchmarks/benchmarks/bench_compile.py | 16 ++++++++++++++-- benchmarks/benchmarks/bench_copy.py | 9 +++++++++ benchmarks/benchmarks/bench_submit.py | 10 +++++++++- 4 files changed, 40 insertions(+), 6 deletions(-) diff --git a/benchmarks/benchmarks/_utils.py b/benchmarks/benchmarks/_utils.py index 8efcea5691..3fcf5b3f65 100644 --- a/benchmarks/benchmarks/_utils.py +++ b/benchmarks/benchmarks/_utils.py @@ -27,10 +27,11 @@ from asv_runner.benchmarks.mark import SkipNotImplemented import dpctl -import dpctl.compiler as dpc -# Device axis. Filter-selector strings accepted by SyclQueue/SyclDevice. -_SELECTORS = ["cpu", "gpu"] +# Device axis. Fully-qualified filter-selector strings, so a node exposing +# a device through more than one backend still resolves each benchmark to a +# fixed, unambiguous device. +_SELECTORS = ["opencl:cpu", "level_zero:gpu"] # Allocation and transfer size sweep, in bytes: 4 KiB, 1 MiB, 16 MiB, 256 MiB. _SIZES = [4 * 1024, 1024**2, 16 * 1024**2, 256 * 1024**2] @@ -102,6 +103,10 @@ def opencl_queue_or_skip(): def sycl_source_queue_or_skip(selector): """Return a queue whose device can compile SYCL source, or skip.""" + try: + import dpctl.compiler as dpc + except ImportError: + raise SkipNotImplemented("dpctl.compiler is not available") q = queue_for(selector) if not dpc.is_sycl_source_compilation_available(): raise SkipNotImplemented("SYCL source compilation extension absent") diff --git a/benchmarks/benchmarks/bench_compile.py b/benchmarks/benchmarks/bench_compile.py index 200dbb5a56..9cafc1f8be 100644 --- a/benchmarks/benchmarks/bench_compile.py +++ b/benchmarks/benchmarks/bench_compile.py @@ -24,7 +24,12 @@ import itertools -import dpctl.compiler as dpc +from asv_runner.benchmarks.mark import SkipNotImplemented + +try: + import dpctl.compiler as dpc +except ImportError: + dpc = None from ._utils import ( _SELECTORS, @@ -44,6 +49,8 @@ class BundleFromSPIRV: param_names = ["selector"] def setup(self, selector): + if dpc is None: + raise SkipNotImplemented("dpctl.compiler is not available") self.queue = queue_for(selector) self.spirv = spirv_bytes() self.bundle = dpc.create_kernel_bundle_from_spirv( @@ -69,6 +76,8 @@ class BundleFromOpenCLSource: warmup_time = 0 def setup(self): + if dpc is None: + raise SkipNotImplemented("dpctl.compiler is not available") self.queue = opencl_queue_or_skip() self.counter = itertools.count() self.warm_source = ocl_axpy_source() @@ -98,6 +107,8 @@ class BundleFromSYCLSource: warmup_time = 0 def setup(self, selector): + if dpc is None: + raise SkipNotImplemented("dpctl.compiler is not available") self.queue = sycl_source_queue_or_skip(selector) self.counter = itertools.count() self.warm_source = sycl_axpy_source() @@ -124,9 +135,10 @@ class SourceCompilationProbe: def setup(self, selector): self.device = queue_for(selector).sycl_device - dpc.is_sycl_source_compilation_available() def time_is_sycl_source_compilation_available(self, selector): + if dpc is None: + raise SkipNotImplemented("dpctl.compiler is not available") dpc.is_sycl_source_compilation_available() def time_can_compile(self, selector): diff --git a/benchmarks/benchmarks/bench_copy.py b/benchmarks/benchmarks/bench_copy.py index a110cfb854..0515aece08 100644 --- a/benchmarks/benchmarks/bench_copy.py +++ b/benchmarks/benchmarks/bench_copy.py @@ -21,6 +21,7 @@ """ import numpy as np +from asv_runner.benchmarks.mark import SkipNotImplemented import dpctl.memory as dpm @@ -95,16 +96,24 @@ def setup(self, selector, nbytes): self.dev.memset() def time_memset(self, selector, nbytes): + if not hasattr(self.queue, "memset"): + raise SkipNotImplemented("SyclQueue.memset not available") self.queue.memset(self.dev, 0, self.nbytes) def time_memory_memset(self, selector, nbytes): self.dev.memset() def time_fill_u1(self, selector, nbytes): + if not hasattr(self.queue, "fill"): + raise SkipNotImplemented("SyclQueue.fill not available") self.queue.fill(self.dev, 0, self.nbytes, "u1") def time_fill_f4(self, selector, nbytes): + if not hasattr(self.queue, "fill"): + raise SkipNotImplemented("SyclQueue.fill not available") self.queue.fill(self.dev, 0.0, self.nbytes // 4, "f4") def time_memset_async_wait(self, selector, nbytes): + if not hasattr(self.queue, "memset_async"): + raise SkipNotImplemented("SyclQueue.memset_async not available") self.queue.memset_async(self.dev, 0, self.nbytes).wait() diff --git a/benchmarks/benchmarks/bench_submit.py b/benchmarks/benchmarks/bench_submit.py index c1b02c21c3..eced0c251c 100644 --- a/benchmarks/benchmarks/bench_submit.py +++ b/benchmarks/benchmarks/bench_submit.py @@ -20,10 +20,16 @@ measured is dpctl's submission path and the runtime round trip, not compute. """ +from asv_runner.benchmarks.mark import SkipNotImplemented + import dpctl -import dpctl.compiler as dpc import dpctl.memory as dpm +try: + import dpctl.compiler as dpc +except ImportError: + dpc = None + from ._utils import _SELECTORS, queue_for, spirv_bytes _BATCH = 100 @@ -36,6 +42,8 @@ class Submit: param_names = ["selector"] def setup(self, selector): + if dpc is None: + raise SkipNotImplemented("dpctl.compiler is not available") self.queue = queue_for(selector) bundle = dpc.create_kernel_bundle_from_spirv( self.queue, spirv_bytes() From 53ac512b6be267fd15857137c14572e0f79a8c39 Mon Sep 17 00:00:00 2001 From: vchamarthi Date: Wed, 9 Sep 2026 18:22:00 -0500 Subject: [PATCH 3/3] remove dead weight benchmarks --- benchmarks/benchmarks/bench_compile.py | 5 ----- benchmarks/benchmarks/bench_construct.py | 17 +---------------- benchmarks/benchmarks/bench_submit.py | 3 --- benchmarks/benchmarks/bench_usm.py | 6 ------ 4 files changed, 1 insertion(+), 30 deletions(-) diff --git a/benchmarks/benchmarks/bench_compile.py b/benchmarks/benchmarks/bench_compile.py index 9cafc1f8be..ba5c16bad6 100644 --- a/benchmarks/benchmarks/bench_compile.py +++ b/benchmarks/benchmarks/bench_compile.py @@ -136,10 +136,5 @@ class SourceCompilationProbe: def setup(self, selector): self.device = queue_for(selector).sycl_device - def time_is_sycl_source_compilation_available(self, selector): - if dpc is None: - raise SkipNotImplemented("dpctl.compiler is not available") - dpc.is_sycl_source_compilation_available() - def time_can_compile(self, selector): self.device.can_compile("sycl") diff --git a/benchmarks/benchmarks/bench_construct.py b/benchmarks/benchmarks/bench_construct.py index cd66cb667a..ebf1f22176 100644 --- a/benchmarks/benchmarks/bench_construct.py +++ b/benchmarks/benchmarks/bench_construct.py @@ -56,14 +56,11 @@ def time_sycl_queue_in_order(self, selector): class ConstructPlatform: - """Construction and enumeration cost of SyclPlatform.""" + """Construction cost of SyclPlatform.""" def time_sycl_platform_default(self): dpctl.SyclPlatform() - def time_get_platforms(self): - dpctl.get_platforms() - class DeviceProperties: """Hot device attribute reads. @@ -90,9 +87,6 @@ def time_device_property_reads(self, selector): d.global_mem_size d.has_aspect_fp64 - def time_device_filter_string(self, selector): - self.device.filter_string - class QueueAccessors: """Accessor cost on an existing queue.""" @@ -106,14 +100,5 @@ def setup(self, selector): def time_sycl_device_attr(self, selector): self.queue.sycl_device - def time_get_sycl_device(self, selector): - self.queue.get_sycl_device() - def time_sycl_context_attr(self, selector): self.queue.sycl_context - - def time_queue_hash(self, selector): - hash(self.queue) - - def time_queue_eq(self, selector): - self.queue == self.queue diff --git a/benchmarks/benchmarks/bench_submit.py b/benchmarks/benchmarks/bench_submit.py index eced0c251c..5c597529d5 100644 --- a/benchmarks/benchmarks/bench_submit.py +++ b/benchmarks/benchmarks/bench_submit.py @@ -72,9 +72,6 @@ def time_submit_async_batch(self, selector): ev = q.submit_async(self.kernel, self.args, self.range) ev.wait() - def time_submit_with_local_range(self, selector): - self.queue.submit(self.kernel, self.args, [1], [1]) - class Synchronize: """Barrier, queue wait, and event handling costs on an idle queue.""" diff --git a/benchmarks/benchmarks/bench_usm.py b/benchmarks/benchmarks/bench_usm.py index 2336d1b3d4..eb7c4cd66a 100644 --- a/benchmarks/benchmarks/bench_usm.py +++ b/benchmarks/benchmarks/bench_usm.py @@ -86,11 +86,5 @@ def setup(self, selector, usm_type): def time_get_usm_type(self, selector, usm_type): self.mem.get_usm_type() - def time_get_usm_type_enum(self, selector, usm_type): - self.mem.get_usm_type_enum() - def time_sycl_usm_array_interface(self, selector, usm_type): self.mem.__sycl_usm_array_interface__ - - def time_sycl_queue_attr(self, selector, usm_type): - self.mem.sycl_queue