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..3fcf5b3f65
--- /dev/null
+++ b/benchmarks/benchmarks/_utils.py
@@ -0,0 +1,164 @@
+#                      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
+
+# 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]
+
+# 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."""
+    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")
+    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 <sycl/sycl.hpp>
+
+    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..ba5c16bad6
--- /dev/null
+++ b/benchmarks/benchmarks/bench_compile.py
@@ -0,0 +1,140 @@
+#                      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
+
+from asv_runner.benchmarks.mark import SkipNotImplemented
+
+try:
+    import dpctl.compiler as dpc
+except ImportError:
+    dpc = None
+
+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):
+        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(
+            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):
+        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()
+        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):
+        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()
+        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
+
+    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..ebf1f22176
--- /dev/null
+++ b/benchmarks/benchmarks/bench_construct.py
@@ -0,0 +1,104 @@
+#                      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 cost of SyclPlatform."""
+
+    def time_sycl_platform_default(self):
+        dpctl.SyclPlatform()
+
+
+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
+
+
+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_sycl_context_attr(self, selector):
+        self.queue.sycl_context
diff --git a/benchmarks/benchmarks/bench_copy.py b/benchmarks/benchmarks/bench_copy.py
new file mode 100644
index 0000000000..0515aece08
--- /dev/null
+++ b/benchmarks/benchmarks/bench_copy.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 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
+from asv_runner.benchmarks.mark import SkipNotImplemented
+
+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):
+        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_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..5c597529d5
--- /dev/null
+++ b/benchmarks/benchmarks/bench_submit.py
@@ -0,0 +1,93 @@
+#                      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.
+"""
+
+from asv_runner.benchmarks.mark import SkipNotImplemented
+
+import dpctl
+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
+
+
+class Submit:
+    """Submission overhead on a one-work-item range."""
+
+    params = [_SELECTORS]
+    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()
+        )
+        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()
+
+
+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..eb7c4cd66a
--- /dev/null
+++ b/benchmarks/benchmarks/bench_usm.py
@@ -0,0 +1,90 @@
+#                      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_sycl_usm_array_interface(self, selector, usm_type):
+        self.mem.__sycl_usm_array_interface__
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",