|
| 1 | +# Data Parallel Control (dpctl) |
| 2 | +# |
| 3 | +# Copyright 2026 Intel Corporation |
| 4 | +# |
| 5 | +# Licensed under the Apache License, Version 2.0 (the "License"); |
| 6 | +# you may not use this file except in compliance with the License. |
| 7 | +# You may obtain a copy of the License at |
| 8 | +# |
| 9 | +# http://www.apache.org/licenses/LICENSE-2.0 |
| 10 | +# |
| 11 | +# Unless required by applicable law or agreed to in writing, software |
| 12 | +# distributed under the License is distributed on an "AS IS" BASIS, |
| 13 | +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| 14 | +# See the License for the specific language governing permissions and |
| 15 | +# limitations under the License. |
| 16 | + |
| 17 | +"""Shared utilities for dpctl benchmarks. |
| 18 | +
|
| 19 | +Every benchmark that needs a device goes through :func:`queue_for` or |
| 20 | +:func:`device_for` so that benchmark names stay identical on every node in |
| 21 | +the pool: a node without a given device reports the parameter as skipped |
| 22 | +rather than failing the whole suite. |
| 23 | +""" |
| 24 | + |
| 25 | +import os |
| 26 | + |
| 27 | +from asv_runner.benchmarks.mark import SkipNotImplemented |
| 28 | + |
| 29 | +import dpctl |
| 30 | + |
| 31 | +# Device selectors |
| 32 | +_SELECTORS = ["opencl:cpu", "level_zero:gpu"] |
| 33 | + |
| 34 | +# Allocation and transfer sizes, in bytes. |
| 35 | +_SIZES = [4 * 1024, 1024**2, 16 * 1024**2, 256 * 1024**2] |
| 36 | + |
| 37 | +# USM kinds. |
| 38 | +_USM_TYPES = ["device", "host", "shared"] |
| 39 | + |
| 40 | +# Max fraction of device memory a single allocation may claim. |
| 41 | +_MEM_BUDGET = 0.25 |
| 42 | + |
| 43 | +_queues = {} |
| 44 | +_devices = {} |
| 45 | +_spirv = None |
| 46 | + |
| 47 | + |
| 48 | +def queue_for(selector): |
| 49 | + """Return a memoized queue for *selector*, or skip when unavailable.""" |
| 50 | + if selector not in _queues: |
| 51 | + try: |
| 52 | + _queues[selector] = dpctl.SyclQueue(selector) |
| 53 | + except dpctl.SyclQueueCreationError: |
| 54 | + _queues[selector] = None |
| 55 | + q = _queues[selector] |
| 56 | + if q is None: |
| 57 | + raise SkipNotImplemented(f"no {selector} device available") |
| 58 | + return q |
| 59 | + |
| 60 | + |
| 61 | +def device_for(selector): |
| 62 | + """Return a memoized device for *selector*, or skip when unavailable.""" |
| 63 | + if selector not in _devices: |
| 64 | + try: |
| 65 | + _devices[selector] = dpctl.SyclDevice(selector) |
| 66 | + except dpctl.SyclDeviceCreationError: |
| 67 | + _devices[selector] = None |
| 68 | + d = _devices[selector] |
| 69 | + if d is None: |
| 70 | + raise SkipNotImplemented(f"no {selector} device available") |
| 71 | + return d |
| 72 | + |
| 73 | + |
| 74 | +def usm_class(usm_type): |
| 75 | + """Return the dpctl.memory class allocating *usm_type* memory.""" |
| 76 | + import dpctl.memory as dpm |
| 77 | + |
| 78 | + return { |
| 79 | + "device": dpm.MemoryUSMDevice, |
| 80 | + "host": dpm.MemoryUSMHost, |
| 81 | + "shared": dpm.MemoryUSMShared, |
| 82 | + }[usm_type] |
| 83 | + |
| 84 | + |
| 85 | +def skip_unless_fits(queue, nbytes): |
| 86 | + """Skip when *nbytes* exceeds this device's allocation budget.""" |
| 87 | + budget = _MEM_BUDGET * queue.sycl_device.global_mem_size |
| 88 | + if nbytes > budget: |
| 89 | + raise SkipNotImplemented( |
| 90 | + f"{nbytes} bytes exceeds the device memory budget" |
| 91 | + ) |
| 92 | + |
| 93 | + |
| 94 | +def opencl_queue_or_skip(): |
| 95 | + """Return a memoized OpenCL queue, or skip. |
| 96 | +
|
| 97 | + ``create_kernel_bundle_from_source`` only supports the OpenCL backend. |
| 98 | + """ |
| 99 | + return queue_for("opencl") |
| 100 | + |
| 101 | + |
| 102 | +def sycl_source_queue_or_skip(selector): |
| 103 | + """Return a queue whose device can compile SYCL source, or skip.""" |
| 104 | + try: |
| 105 | + import dpctl.compiler as dpc |
| 106 | + except ImportError: |
| 107 | + raise SkipNotImplemented("dpctl.compiler is not available") |
| 108 | + q = queue_for(selector) |
| 109 | + if not dpc.is_sycl_source_compilation_available(): |
| 110 | + raise SkipNotImplemented("SYCL source compilation extension absent") |
| 111 | + if not q.sycl_device.can_compile("sycl"): |
| 112 | + raise SkipNotImplemented("device cannot compile SYCL source") |
| 113 | + return q |
| 114 | + |
| 115 | + |
| 116 | +def spirv_bytes(): |
| 117 | + """Return the SPIR-V module shipped with the installed dpctl, or skip. |
| 118 | +
|
| 119 | + Defines ``add(int*, int*, int*)`` and ``axpy(int*, int*, int*, int)``. |
| 120 | + """ |
| 121 | + global _spirv |
| 122 | + if _spirv is None: |
| 123 | + path = os.path.join( |
| 124 | + os.path.dirname(os.path.abspath(dpctl.__file__)), |
| 125 | + "tests", |
| 126 | + "input_files", |
| 127 | + "multi_kernel.spv", |
| 128 | + ) |
| 129 | + if not os.path.exists(path): |
| 130 | + raise SkipNotImplemented(f"SPIR-V module not found at {path}") |
| 131 | + with open(path, "rb") as fh: |
| 132 | + _spirv = fh.read() |
| 133 | + return _spirv |
| 134 | + |
| 135 | + |
| 136 | +def ocl_axpy_source(kernel_name="axpy"): |
| 137 | + """Return OpenCL C source for an axpy kernel called *kernel_name*.""" |
| 138 | + return ( |
| 139 | + f"kernel void {kernel_name}(" |
| 140 | + " global int *a, global int *b, global int *c, int d) {" |
| 141 | + " size_t index = get_global_id(0);" |
| 142 | + " c[index] = d * a[index] + b[index];" |
| 143 | + "}" |
| 144 | + ) |
| 145 | + |
| 146 | + |
| 147 | +def sycl_axpy_source(kernel_name="axpy"): |
| 148 | + """Return SYCL source for an axpy function called *kernel_name*.""" |
| 149 | + return f""" |
| 150 | + #include <sycl/sycl.hpp> |
| 151 | +
|
| 152 | + extern "C" SYCL_EXTERNAL |
| 153 | + void {kernel_name}(int* a, int* b, int* c, int d, size_t i) {{ |
| 154 | + c[i] = d * a[i] + b[i]; |
| 155 | + }} |
| 156 | + """ |
0 commit comments