diff --git a/.github/workflows/test-python-package.yml b/.github/workflows/test-python-package.yml index f86240c7f..d1a151032 100644 --- a/.github/workflows/test-python-package.yml +++ b/.github/workflows/test-python-package.yml @@ -24,16 +24,26 @@ jobs: os: [ubuntu-latest, macos-latest] steps: - - uses: actions/checkout@v4 + - uses: actions/checkout@v7 - name: Setup Nox - uses: fjwillemsen/setup-nox2@v3.0.0 + uses: fjwillemsen/setup-nox2@v4 - name: Setup Poetry uses: Gr1N/setup-poetry@v9 - run: poetry self add poetry-plugin-export + - uses: julia-actions/setup-julia@v3 + with: + version: '1.11' # when changed, also see `require_julia` in noxfile.py and the Julia version in Project.toml - name: Run tests with Nox run: | + rm -rf .nox pip install nox-poetry - nox -- skip-gpu github-action + nox -- skip-gpu skip-julia github-action + - name: Run Julia tests with Nox + run: | + rm -rf .nox + pip install nox-poetry + nox --session tests-3.14 -- skip-gpu github-action + # [ -d "~" ] && mkdir -p ~/.julia/registries # - name: Upload Coverage report to CodeCov # uses: codecov/codecov-action@v3 # with: diff --git a/.vscode/extensions.json b/.vscode/extensions.json index 035235be3..2fce92093 100644 --- a/.vscode/extensions.json +++ b/.vscode/extensions.json @@ -1,14 +1,15 @@ { - // See https://go.microsoft.com/fwlink/?LinkId=827846 to learn about workspace recommendations. - // Extension identifier format: ${publisher}.${name}. Example: vscode.csharp - // List of extensions which should be recommended for users of this workspace. - "recommendations": [ - "ms-python.python", - "ms-python.black-formatter", - "charliermarsh.ruff", - "bungcip.better-toml", - "njpwerner.autodocstring", - ], - // List of extensions recommended by VS Code that should not be recommended for users of this workspace. - "unwantedRecommendations": [] + // See https://go.microsoft.com/fwlink/?LinkId=827846 to learn about workspace recommendations. + // Extension identifier format: ${publisher}.${name}. Example: vscode.csharp + // List of extensions which should be recommended for users of this workspace. + "recommendations": [ + "ms-python.python", + "ms-python.black-formatter", + "charliermarsh.ruff", + "bungcip.better-toml", + "njpwerner.autodocstring", + "julialang.language-julia", + ], + // List of extensions recommended by VS Code that should not be recommended for users of this workspace. + "unwantedRecommendations": [] } \ No newline at end of file diff --git a/.vscode/settings.json b/.vscode/settings.json index 3089f374a..f5337c24f 100755 --- a/.vscode/settings.json +++ b/.vscode/settings.json @@ -17,9 +17,7 @@ "black-formatter.args": [ "--config=pyproject.toml" ], - "ruff.args": [ - "--config=pyproject.toml" - ], + "ruff.configuration": "pyproject.toml", "autoDocstring.docstringFormat": "google-notypes", "esbonio.sphinx.confDir": "", "python.testing.pytestArgs": [ @@ -27,4 +25,8 @@ ], "python.testing.unittestEnabled": false, "python.testing.pytestEnabled": true, -} + "sonarlint.connectedMode.project": { + "projectKey": "KernelTuner_kernel_tuner", + "connectionId": "kerneltuner", + } +} \ No newline at end of file diff --git a/CONTRIBUTING.rst b/CONTRIBUTING.rst index 7b8a46dc3..47fef18a5 100644 --- a/CONTRIBUTING.rst +++ b/CONTRIBUTING.rst @@ -26,7 +26,7 @@ Before creating a pull request please ensure the following: * You are working in an up-to-date development environment * You are a human developer. We are not interested in purely AI generated code contributions. -* You have written unit tests to test your additions and all unit tests pass (run :bash:`nox`). If you do not have the required hardware, you can run :bash:`nox -- skip-gpu`, or :bash:`skip-cuda`, :bash:`skip-hip`, :bash:`skip-opencl`. +* You have written unit tests to test your additions and all unit tests pass (run :bash:`nox`). If you do not have the required hardware, you can run :bash:`nox -- skip-gpu`, or :bash:`skip-cuda`, :bash:`skip-hip`, :bash:`skip-opencl`, :bash:`skip-julia`. * The examples still work and produce the same (or better) results * An entry about the change or addition is created in :bash:`CHANGELOG.md` diff --git a/doc/source/dev-environment.rst b/doc/source/dev-environment.rst index 0adb3c83e..275f40039 100644 --- a/doc/source/dev-environment.rst +++ b/doc/source/dev-environment.rst @@ -40,7 +40,7 @@ Steps with :bash:`sudo` access (e.g. on a local device): * Activate the environment with :bash:`pyenv activate kerneltuner`. * Make sure :bash:`which python` and :bash:`which pip` point to the expected Python location and version. * Update Pip with :bash:`pip install --upgrade pip`. -#. Install the project, dependencies and extras: :bash:`poetry install --with test,docs -E cuda -E opencl -E hip`, leaving out :bash:`-E cuda`, :bash:`-E opencl` or :bash:`-E hip` if this does not apply on your system. To go all-out, use :bash:`--all-extras` +#. Install the project, dependencies and extras: :bash:`poetry install --with test,docs -E cuda -E opencl -E hip -E julia`, leaving out :bash:`-E cuda`, :bash:`-E opencl` etc. if this does not apply on your system. To go all-out, use :bash:`--all-extras` * Depending on the environment, it may be necessary or convenient to install extra packages such as :bash:`cupy-cuda11x` / :bash:`cupy-cuda12x`, and :bash:`cuda-python`. These are currently not defined as dependencies for kernel-tuner, but can be part of tests. * Do not forget to make sure the paths are set correctly. If you're using CUDA, the desired CUDA version should be in :bash:`$PATH`, :bash:`$LD_LIBARY_PATH` and :bash:`$CPATH`. * Re-open the shell for changes to take effect. @@ -72,8 +72,9 @@ Steps without :bash:`sudo` access (e.g. on a cluster): #. `Install Poetry `__. * Use :bash:`curl -sSL https://install.python-poetry.org | python3 -` to install Poetry. * Add the poetry export plugin with :bash:`poetry self add poetry-plugin-export`. -#. Install the project, dependencies and extras: :bash:`poetry install --with test,docs -E cuda -E opencl -E hip`, leaving out :bash:`-E cuda`, :bash:`-E opencl` or :bash:`-E hip` if this does not apply on your system. To go all-out, use :bash:`--all-extras`. +#. Install the project, dependencies and extras: :bash:`poetry install --with test,docs -E cuda -E opencl -E hip -E julia`, leaving out :bash:`-E cuda`, :bash:`-E opencl` etc. if this does not apply on your system. To go all-out, use :bash:`--all-extras`. * If you run into "keyring" or other seemingly weird issues, this is a known issue with Poetry on some systems. Do: :bash:`pip install keyring`, :bash:`python3 -m keyring --disable`. + * Kernel Tuner has a dependency on Python-Constraint, which provides binaries for most systems. On some older systems, these binaries may not be compatible (on Linux usually due to an outdated LDD version, check with :bash:`ldd --version`, must be >=2.35). Without binaries, your system will try to build it yourself, so make sure build tools are available (e.g. with :bash:`module load gcc/13.2.0`). * Depending on the environment, it may be necessary or convenient to install extra packages such as :bash:`cupy-cuda11x` / :bash:`cupy-cuda12x`, and :bash:`cuda-python`. These are currently not defined as dependencies for kernel-tuner, but can be part of tests. * Verify that your development environment has no missing installs or updates with :bash:`poetry install --sync --dry-run --with test`. #. Check if the environment is setup correctly by running :bash:`pytest`. All tests should pass, except if you're not on a GPU node, or one or more extras has been left out in the previous step, then these tests will skip gracefully. @@ -102,9 +103,10 @@ If you do not have fully compatible hardware or environment, you can use the fol * :bash:`nox -- skip-cuda` to skip tests involving CUDA. * :bash:`nox -- skip-hip` to skip tests involving HIP. * :bash:`nox -- skip-opencl` to skip tests involving OpenCL. -* :bash:`nox -- skip-gpu` to skip all tests on the GPU (the same as :bash:`nox -- skip-cuda skip-hip skip-opencl`), especially helpful if you don't have a GPU locally. +* :bash:`nox -- skip-julia` to skip tests involving Julia. +* :bash:`nox -- skip-gpu` to skip all tests on the GPU (the same as :bash:`nox -- skip-cuda skip-hip skip-opencl`), especially helpful if you don't have a GPU in your system. -Contributions you make to the Kernel Tuner should not break any of the tests even if you cannot run them locally! +Contributions you make to Kernel Tuner should not break any of the tests even if you cannot run them locally! Running with :bash:`pytest` will test against your local Python version and PIP packages. In this case, tests that require PyCuda and/or a CUDA capable GPU will be skipped automatically if these are not installed/present. diff --git a/kernel_tuner/accuracy.py b/kernel_tuner/accuracy.py index a04b0a81b..74d3851a1 100644 --- a/kernel_tuner/accuracy.py +++ b/kernel_tuner/accuracy.py @@ -1,19 +1,25 @@ +"""Module for accuracy measurement and tunable parameters.""" + +import logging +import re from collections import UserDict from typing import Dict + import numpy as np -import logging -import re from kernel_tuner.observers import OutputObserver class Tunable(UserDict): + """``Tunable`` can be used as a parameter value dependent input argument when tuning kernels.""" + def __init__(self, param_key: str, arrays: Dict): - """The ``Tunable`` object can be used as an input argument when tuning - kernels. It is a container that holds several arrays internally and + """The ``Tunable`` object can be used as an input argument when tuning kernels. + + It is a container that holds several arrays internally and selects one array during benchmarking based on the value of a tunable parameter. - Example + Example: ------- Consider this example:: @@ -37,6 +43,7 @@ def __init__(self, param_key: str, arrays: Dict): self.param_key = param_key def select_for_configuration(self, params): + """Select the array for the given configuration.""" if callable(self.param_key): option = self.param_key(params) elif self.param_key in params: @@ -46,13 +53,16 @@ def select_for_configuration(self, params): if option not in self.data: list = ", ".join(map(str, self.data.keys())) - raise KeyError( - f"'{option}' is not a valid parameter value, should be one of: {list}" - ) + raise KeyError(f"'{option}' is not a valid parameter value, should be one of: {list}") - return self.data[option] + # continue recursively until we find a non-Tunable + el = self.data[option] + if isinstance(el, Tunable): + return el.select_for_configuration(params) + else: + return el - def __call__(self, params): + def __call__(self, params): # noqa: D102 return self.select_for_configuration(params) @@ -70,6 +80,7 @@ def _find_bfloat16_if_available(): if dtype is None: try: from ml_dtypes import bfloat16 + dtype = bfloat16 except ImportError: pass @@ -78,6 +89,7 @@ def _find_bfloat16_if_available(): if dtype is None: try: from jax.numpy import bfloat16 + dtype = bfloat16 except ImportError: pass @@ -86,6 +98,7 @@ def _find_bfloat16_if_available(): if dtype is None: try: from tensorflow import bfloat16 + dtype = bfloat16.as_numpy_dtype except ImportError: pass @@ -100,9 +113,9 @@ def _find_bfloat16_if_available(): def _to_float_dtype(x: str) -> np.dtype: - """Convert a string to a numpy data type (``dtype``). This function recognizes - common names (such as ``f16`` or ``kfloat``), and uses ``np.dtype(x)`` as a - fallback. + """Convert a string to a numpy data type (``dtype``). + + This function recognizes common names (such as ``f16`` or ``kfloat``), and uses ``np.dtype(x)`` as a fallback. """ if isinstance(x, str): x = x.lower() @@ -123,16 +136,17 @@ def _to_float_dtype(x: str) -> np.dtype: class TunablePrecision(Tunable): - def __init__( - self, param_key: str, array: np.ndarray, dtypes: Dict[str, np.dtype] = None - ): - """The ``Tunable`` object can be used as an input argument when tuning - kernels. It is a container that internally holds several arrays + """``TunablePrecision`` can be used as a precision-level dependent input argument when tuning kernels.""" + + def __init__(self, param_key: str, array: np.ndarray, dtypes: Dict[str, np.dtype] = None): + """The ``Tunable`` object can be used as an input argument when tuning kernels. + + It is a container that internally holds several arrays containing the same data, but stored in using different levels of precision. During benchamrking, one array is selected based on the value of the tunable parameter called ``param_key``. - Example + Example: ------- Consider this example:: @@ -156,7 +170,6 @@ def __init__( if bfloat16 is not None: dtypes["bfloat16"] = bfloat16 - # If dtype is a list, convert it to a dictionary if isinstance(dtypes, (list, tuple)): dtypes = dict((name, _to_float_dtype(name)) for name in dtypes) @@ -197,7 +210,6 @@ def error_metric_from_name(user_key, EPS=1e-8): The value of `EPS` is used for relative errors to prevent division by zero. `` """ - # Prepocess the provided name: # - convert to lowercase # - remove the word "error" @@ -278,17 +290,13 @@ def metric(a, b): raise ValueError(f"invalid error metric provided: {user_key}") # cast both arguments to f64 before passing them to the metric - return lambda a, b: metric( - a.astype(np.float64, copy=False), b.astype(np.float64, copy=False) - ) + return lambda a, b: metric(a.astype(np.float64, copy=False), b.astype(np.float64, copy=False)) class AccuracyObserver(OutputObserver): - """``AccuracyObserver`` measures the error on the output produced by a kernel - by comparing the output against a reference output. + """``AccuracyObserver`` measures the error on the output produced by a kernel by comparing to a reference output. - By default, it uses the root mean-squared error (RMSE) and uses the - metric name ``"error"``. + By default, it uses the root mean-squared error (RMSE) and uses the metric name ``"error"``. """ def __init__(self, metric=None, key="error", *, atol=1e-8): @@ -303,7 +311,6 @@ def __init__(self, metric=None, key="error", *, atol=1e-8): :param atol: The tolerance used in relative metrics to prevent division by zero. It is ignored by absolute error metrics. """ - # Default metric is RMSE if not metric: metric = "rmse" @@ -317,6 +324,7 @@ def __init__(self, metric=None, key="error", *, atol=1e-8): self.result = None def process_output(self, answers, outputs): + """Process the output produced by the kernel and compare it to the reference answers.""" errors = [] for answer, output in zip(answers, outputs): @@ -326,4 +334,5 @@ def process_output(self, answers, outputs): self.result = max(errors) def get_results(self): + """Get the results produced by this observer.""" return dict([(self.key, self.result)]) diff --git a/kernel_tuner/backends/compiler.py b/kernel_tuner/backends/compiler.py index 06402ff7c..ec2b2da6f 100644 --- a/kernel_tuner/backends/compiler.py +++ b/kernel_tuner/backends/compiler.py @@ -1,13 +1,13 @@ -""" This module contains the functionality for running and compiling C functions """ +"""This module contains the functionality for running and compiling C functions""" -from collections import namedtuple -import subprocess -import platform +import _ctypes +import ctypes as C import errno -import re import logging -import ctypes as C -import _ctypes +import platform +import re +import subprocess +from collections import namedtuple import numpy as np import numpy.ctypeslib @@ -15,10 +15,9 @@ from kernel_tuner.backends.backend import CompilerBackend from kernel_tuner.observers.compiler import CompilerRuntimeObserver from kernel_tuner.util import ( - get_temp_filename, delete_temp_file, + get_temp_filename, write_file, - SkippableFailure, ) try: @@ -88,7 +87,7 @@ class CompilerFunctions(CompilerBackend): """Class that groups the code for running and compiling C functions""" def __init__(self, iterations=7, compiler_options=None, compiler=None, observers=None): - """instantiate CFunctions object used for interacting with C code + """Instantiate CFunctions object used for interacting with C code :param iterations: Number of iterations used while benchmarking a kernel, 7 by default. :type iterations: int @@ -146,7 +145,7 @@ def __init__(self, iterations=7, compiler_options=None, compiler=None, observers self.name = platform.processor() def ready_argument_list(self, arguments): - """ready argument list to be passed to the C function + """Ready argument list to be passed to the C function :param arguments: List of arguments to be passed to the C function. The order should match the argument list on the C function. @@ -181,7 +180,7 @@ def ready_argument_list(self, arguments): return ctype_args def compile(self, kernel_instance): - """call the C compiler to compile the kernel, return the function + """Call the C compiler to compile the kernel, return the function :param kernel_instance: An object representing the specific instance of the tunable kernel in the parameter space. @@ -311,29 +310,33 @@ def compile(self, kernel_instance): def start_event(self): """Records the event that marks the start of a measurement - C backend does not use events""" + C backend does not use events + """ pass def stop_event(self): """Records the event that marks the end of a measurement - C backend does not use events""" + C backend does not use events + """ pass def kernel_finished(self): """Returns True if the kernel has finished, False otherwise - C backend does not support asynchronous launches""" + C backend does not support asynchronous launches + """ return True def synchronize(self): """Halts execution until device has finished its tasks - C backend does not support asynchronous launches""" + C backend does not support asynchronous launches + """ pass def run_kernel(self, func, c_args, threads, grid, stream=None): - """runs the kernel once, returns whatever the kernel returns + """Runs the kernel once, returns whatever the kernel returns :param func: A C function compiled for this specific configuration :type func: ctypes._FuncPtr @@ -367,7 +370,7 @@ def run_kernel(self, func, c_args, threads, grid, stream=None): return time def memset(self, allocation, value, size): - """set the memory in allocation to the value in value + """Set the memory in allocation to the value in value :param allocation: An Argument for some memory allocation unit :type allocation: Argument @@ -419,7 +422,7 @@ def refresh_memory(self, _, arguments, should_sync): self.memcpy_dtoh(arg, self.allocations[i]) def cleanup_lib(self): - """unload the previously loaded shared library""" + """Unload the previously loaded shared library""" if self.lib is None: return diff --git a/kernel_tuner/backends/hip/hip.py b/kernel_tuner/backends/hip/hip.py index 9c416614d..848b4d46b 100644 --- a/kernel_tuner/backends/hip/hip.py +++ b/kernel_tuner/backends/hip/hip.py @@ -9,8 +9,8 @@ import numpy as np from kernel_tuner.backends.backend import GPUBackend -from kernel_tuner.observers.hip import HipRuntimeObserver from kernel_tuner.backends.hip.util import hip_check +from kernel_tuner.observers.hip import HipRuntimeObserver try: from hip import hip, hiprtc diff --git a/kernel_tuner/backends/hypertuner.py b/kernel_tuner/backends/hypertuner.py index d6f23475a..1e7566d59 100644 --- a/kernel_tuner/backends/hypertuner.py +++ b/kernel_tuner/backends/hypertuner.py @@ -27,12 +27,14 @@ def after_finish(self): self.scores.append(self.dev.last_score) def get_results(self): - results = {'score': mean(self.scores), 'scores': self.scores.copy()} + results = {"score": mean(self.scores), "scores": self.scores.copy()} self.scores = [] return results + class HypertunerFunctions(Backend): """Class for executing hyperparameter tuning.""" + units = {} def __init__(self, iterations, compiler_options=None): @@ -51,30 +53,30 @@ def __init__(self, iterations, compiler_options=None): "name": "dedispersion_milo", "folder": folder, "input_file": "dedispersion_milo.json", - "objective_performance_keys": ["time"] + "objective_performance_keys": ["time"], }, { "name": "hotspot_milo", "folder": folder, "input_file": "hotspot_milo.json", - "objective_performance_keys": ["GFLOP/s"] + "objective_performance_keys": ["GFLOP/s"], }, { "name": "convolution_milo", "folder": folder, "input_file": "convolution_milo.json", - "objective_performance_keys": ["time"] + "objective_performance_keys": ["time"], }, { "name": "gemm_milo", "folder": folder, "input_file": "gemm_milo.json", - "objective_performance_keys": ["time"] - } + "objective_performance_keys": ["time"], + }, ] # any additional settings - self.override = { - "experimental_groups_defaults": { + self.override = { + "experimental_groups_defaults": { "repeats": 25, "samples": self.iterations, "minimum_fraction_of_budget_valid": 0.1, @@ -84,10 +86,8 @@ def __init__(self, iterations, compiler_options=None): "cutoff_percentile": 0.95, "cutoff_percentile_start": 0.01, "cutoff_type": "time", - "objective_time_keys": [ - "all" - ] - } + "objective_time_keys": ["all"], + }, } # override the defaults with compiler options if provided @@ -113,7 +113,7 @@ def ready_argument_list(self, arguments): if arglist is None: arglist = [] return arglist - + def compile(self, kernel_instance): super().compile(kernel_instance) path = Path(__file__).parent.parent.parent / "hyperparamtuning" @@ -121,37 +121,47 @@ def compile(self, kernel_instance): # strategy settings strategy: str = kernel_instance.arguments[0] - hyperparams = [{'name': k, 'value': v} for k, v in kernel_instance.params.items()] + hyperparams = [{"name": k, "value": v} for k, v in kernel_instance.params.items()] hyperparams_string = "_".join(f"{k}={str(v)}" for k, v in kernel_instance.params.items()) - searchspace_strategies = [{ - "autotuner": "KernelTuner", - "name": f"{strategy.lower()}_{hyperparams_string}", - "display_name": strategy.replace('_', ' ').capitalize(), - "search_method": strategy.lower(), - 'search_method_hyperparameters': hyperparams - }] + searchspace_strategies = [ + { + "autotuner": "KernelTuner", + "name": f"{strategy.lower()}_{hyperparams_string}", + "display_name": strategy.replace("_", " ").capitalize(), + "search_method": strategy.lower(), + "search_method_hyperparameters": hyperparams, + } + ] name = kernel_instance.name if len(kernel_instance.name) > 0 else kernel_instance.kernel_source.kernel_name - experiments_filepath = generate_experiment_file(name, path, searchspace_strategies, self.applications, self.gpus, - override=self.override, generate_unique_file=True, overwrite_existing_file=True) + experiments_filepath = generate_experiment_file( + name, + path, + searchspace_strategies, + self.applications, + self.gpus, + override=self.override, + generate_unique_file=True, + overwrite_existing_file=True, + ) return str(experiments_filepath) - + def start_event(self): return super().start_event() - + def stop_event(self): return super().stop_event() - + def kernel_finished(self): super().kernel_finished() return True - + def synchronize(self): return super().synchronize() - + def run_kernel(self, func, gpu_args=None, threads=None, grid=None, stream=None): # from cProfile import Profile - + # # generate the experiments file # experiments_filepath = Path(func) @@ -161,23 +171,23 @@ def run_kernel(self, func, gpu_args=None, threads=None, grid=None, stream=None): # pr.dump_stats('diff_evo_hypertune_hotspot.prof') # self.last_score = scores[list(scores.keys())[0]]['score'] # raise ValueError(scores) - + # generate the experiments file experiments_filepath = Path(func) # run the methodology to get a fitness score for this configuration scores = get_strategy_scores(str(experiments_filepath), full_validate_on_load=False) - self.last_score = scores[list(scores.keys())[0]]['score'] + self.last_score = scores[list(scores.keys())[0]]["score"] # remove the experiments file experiments_filepath.unlink() - + def memset(self, allocation, value, size): return super().memset(allocation, value, size) - + def memcpy_dtoh(self, dest, src): return super().memcpy_dtoh(dest, src) - + def memcpy_htod(self, dest, src): return super().memcpy_htod(dest, src) diff --git a/kernel_tuner/backends/julia.py b/kernel_tuner/backends/julia.py new file mode 100644 index 000000000..16fe88723 --- /dev/null +++ b/kernel_tuner/backends/julia.py @@ -0,0 +1,511 @@ +"""Kernel Tuner backend for running Julia CUDA.jl kernels via JuliaCall. + +This backend allows Julia kernels to be compiled, launched, and observed from Python using Kernel Tuner. + +Requirements: + pip install juliacall + and in Julia: ] add CUDA / ROCBackend / oneAPI / Metal (will be automatically installed if not present) + +Notes: +- The kernel string should contain a valid Julia GPU kernel function definition. +- The kernel name must match the Julia function to be launched. +- Currently supports CuArray and scalar arguments; constant and texture memory are not implemented. +""" + +from pathlib import Path +from warnings import warn +from dataclasses import dataclass + +import numpy as np + +from kernel_tuner.backends.backend import GPUBackend +from kernel_tuner.observers.julia import JuliaRuntimeObserver +from kernel_tuner.util import SkippableFailure + +from .julia_helper import backend_map, detect_julia_gpu_backends + +try: + from juliacall import JuliaError + from juliacall import Main as jl +except ImportError: + jl = None + + +@dataclass(frozen=True) +class JuliaKernel: + """Immutable Julia kernel representation, particularly useful for caching and parallel tuning.""" + function: object + params: tuple + + +class JuliaFunctions(GPUBackend): + """Backend for running Julia kernels (CUDA.jl) through JuliaCall.""" + + units = {"time": "ms"} + last_selected_device = None + + def __init__(self, device=0, iterations=7, compiler_options=None, observers=None): + """Initialize Julia backend using JuliaCall.""" + if jl is None: + raise ImportError("JuliaCall not installed. Please run `pip install juliacall`.") + + # process passed options and backends + self.process_compiler_options(compiler_options) + self.available_backends = detect_julia_gpu_backends() + backend_name = self.verify_backends_with_options(compiler_options) + + # Initialize backend attributes + self.device = device + self.iterations = iterations + self.compiler_options = compiler_options or [] + self.allocations = [] + self.current_kernel = None + self.smem_size = 0 + + # Initialize Julia backend + self.backend = None + self.start_evt = None + self.end_evt = None + self.host_time = None + self.initialize_backend(device, backend_name=backend_name) + + # setup observers + self.observers = observers or [] + self.observers.append( + # TODO this single stateful default observer currently prevents parallel tuning + JuliaRuntimeObserver( + jl.Main.KernelAbstractions, + self, + self.backend, + self.backend_mod, + self.backend_mod_name, + stream=self.stream, + start_event=self.start_evt, + end_event=self.end_evt, + ) + ) + for observer in self.observers: + observer.register_device(self) + + # setup Julia module for kernel launch + if self.backend_mod_name == "CPU": + jl.seval( + f""" + global dest_tmp, src_tmp # for memcpy_htod + module KernelTunerHelper + using KernelAbstractions + const kt_julia_backend = CPU() + const GPUArrayType = {self.GPUArrayType} + include("{str(Path(__file__).parent / "julia_helper.jl")}") + end + """ + ) + else: + jl.seval( + f""" + global dest_tmp, src_tmp # for memcpy_htod + module KernelTunerHelper + using {self.backend_mod_name} + const kt_julia_backend = {self.backend_mod_instname}() + const GPUArrayType = {self.GPUArrayType} + include("{str(Path(__file__).parent / "julia_helper.jl")}") + end + """ + ) + + self.to_gpuarray = jl.KernelTunerHelper.to_gpuarray + self.launch_kernel = jl.KernelTunerHelper.launch_kernel + + # env info + self.env = { + "device_name": self.name, + "compute_capability": self.cc, + "iterations": iterations, + "compiler_options": self.compiler_options, + } + + def initialize_backend(self, device, backend_name): + """Initialize for a choice of Julia backends by backend_name, one of 'cuda', 'amd', 'intel', 'metal'.""" + backend_name = backend_name.upper() + if backend_name not in backend_map: + raise ValueError(f"Unknown backend: {backend_name}") + info = backend_map[backend_name] + backend_pkg = info["pkg"] + + if backend_pkg is not None: + # Ensure the package is installed + self.check_package_and_install(backend_pkg) + + # # Set debug level if needed + # if backend_name == "cuda": + # jl.seval("ENV[\"JULIA_CUDA_DEBUG\"] = \"2\"") + + # Bring module into Python + self.backend_mod_name = info["module"] + self.backend_mod_instname = info["module_backend"] + # jl.seval(f"using KernelAbstractions, {info['module']}") + jl.seval("using KernelAbstractions") + if backend_pkg is not None: + jl.seval(f"using {info['module']}") + backend_mod = getattr(jl.Main, self.backend_mod_name) + self.backend_mod = backend_mod + jl.seval(f"tmp_arr = {info['GPUArrayType']}(Float32.(zeros(2)))") + self.backend = jl.seval("KernelAbstractions.get_backend(tmp_arr)") + self.GPUArrayType = info["GPUArrayType"] + jl.seval("tmp_arr = nothing; GC.gc()") # free temporary array + + # Select device + try: + if int(device) == 0 and backend_pkg != "CUDA": + # Julia uses 1-based indexing, but the CUDA backend uses 0-based so we skip that + device = 1 + jl.seval(info["device_select"](int(device))) + self.last_selected_device = device + except Exception as e: + raise RuntimeError(f"Failed to select Julia {info['module']} device {device}: {e}") from e + + # Query device name + try: + self.name = str(jl.seval(info["name"])) + except JuliaError: + self.name = f"{backend_name}-device-{device}" + + # Query capability if available + if info["capability"] is not None: + try: + cc_tuple = jl.seval(info["capability"]) + # CUDA returns structs with major/minor fields + self.cc = f"{cc_tuple.major}{cc_tuple.minor}" + except JuliaError: + self.cc = None + else: + self.cc = None + + # Query max threads + try: + self.max_threads = int(jl.seval(info["max_threads"])) + except JuliaError: + self.max_threads = None + + # Get the device and context + try: + self.backend_device = self.backend_mod.device() + except Exception: + self.backend_device = None + if backend_name == "CUDA": + self.contextqueue = self.backend_mod.context + # elif backend_name == "AMD": + # self.contextqueue = self.backend_mod.queue + # elif backend_name == "INTEL": + # self.contextqueue = jl.seval( + # f"ZeCommandQueue(ZeContext(first(drivers())), devices(first(drivers()))[{int(device) + 1}]))" + # ) + elif backend_name == "METAL": + self.contextqueue = self.backend_mod.MTLCommandQueue(self.backend_device) + + self.setup_streams(backend_name) + + def __del__(self): + # drop GPUArray references to let Julia GC handle them + try: + for a in self.allocations: + del a + except Exception: + pass + jl.seval("GC.gc()") + + # ------------------------- + # Memory and argument setup + # ------------------------- + + def ready_argument_list(self, arguments): + """Convert arrays to GPU Array in Julia.""" + gpu_args = [] + for arg in arguments: + try: + arr = self.to_gpuarray(arg) + gpu_args.append(arr) + self.allocations.append(arr) + except Exception as e: + raise RuntimeError(f"Failed to move array to GPU: {e}") + return gpu_args + + # ------------------------- + # Compilation + # ------------------------- + + def compile(self, kernel_instance): + """Define Julia kernel function from kernel_instance.kernel_string.""" + kernel_code = kernel_instance.kernel_string + kernel_name = kernel_instance.name + self.kernel_source = kernel_instance.kernel_source + self.host_time = None # reset host time for this kernel instance + + # Extract all 'using' statements and check for required packages + uses = [] + for line in kernel_code.splitlines(): + stripped = line.strip() + # iterate over multiple using/import statements + if stripped.startswith("using ") or stripped.startswith("import "): + for part in stripped.split(","): + uses.append(part.replace("import ", "").replace("using ", "").strip()) + for package in uses: + self.check_package_and_install(package) + + # Wrap in a module to avoid name conflicts + if self.backend_mod_name == "CPU": + module_code = f""" +module KernelTunerUserKernel + {kernel_code} +end + """ + else: + module_code = f""" +module KernelTunerUserKernel + using {self.backend_mod_name} + {kernel_code} +end + """ + try: + jl.seval(module_code) + function = jl.seval(f"KernelTunerUserKernel.{kernel_name}") + return JuliaKernel(function, tuple(kernel_instance.params.values())) # important: the order of params must match the order in the kernel definition + except Exception as e: + raise SkippableFailure(f"Failed to compile Julia kernel: {e} \n{module_code}") + + # ------------------------- + # Kernel launch and timing + # ------------------------- + + def run_kernel(self, func, gpu_args, threads, grid, stream=None): + """Launch a compiled Julia kernel.""" + if func is None or not isinstance(func, JuliaKernel): + raise RuntimeError("No Julia kernel compiled or provided.") + + args_tuple = tuple(gpu_args) + julia_func = func.function + params = func.params + + remove_trailing_ones = lambda tup: tup[ + : len(tup) - next((int(i) for i, x in enumerate(reversed(tup)) if x != 1), len(tup)) + ] + + # Kernel Tuner's grid is number of workgroups; KernelAbstractions' ndrange is number of global work items + if len(grid) != len(threads): + raise ValueError(f"grid and threads must have equal rank: {grid=}, {threads=}") + global_size = tuple(int(g) * int(t) for g, t in zip(grid, threads)) + ndrange = remove_trailing_ones(global_size) + + # prepare launch parameters + ndrange = (1,) if len(ndrange) == 0 else ndrange + workgroupsize = remove_trailing_ones(threads) + workgroupsize = (1,) if len(workgroupsize) == 0 else workgroupsize + + # run the kernel + try: + self.host_time = self.launch_kernel( + julia_func, + args_tuple, + params, + ndrange, + workgroupsize, + int(self.smem_size), + self.observers[-1].start, + self.observers[-1].end, + self.observers[-1].stream, + ) + except JuliaError as e: + if self.raise_errors: + raise e + else: + raise SkippableFailure(f"Julia kernel launch failed for {params=}: {e}") + + def start_event(self): + """Records the event that marks the start of a measurement.""" + if self.backend_mod_name == "CUDA": + evt = self.start_evt() + self.backend_mod.record(evt, self.stream) + self.backend_mod.synchronize(evt) + return evt + elif self.backend_mod_name == "AMDGPU": + evt = self.start_evt(self.stream, do_record=False, timing=True) + self.backend_mod.HIP.record(evt) + self.backend_mod.HIP.synchronize(evt) + return evt + elif self.backend_mod_name == "Metal": + # Because our kernel launch happens via Kernel Abstractions, we wrap our kernel between two command buffers. + # Normally you would just use one command buffer for the actual kernel and take GPUEndTime - GPUStartTime. + jl.start_buf = self.create_metal_buffer() + jl.seval("Metal.commit!(start_buf)") + self.backend_mod.wait_completed(jl.start_buf) + return float(jl.start_buf.GPUEndTime) + + def stop_event(self): + """Records the event that marks the end of a measurement.""" + if self.backend_mod_name == "CUDA": + evt = self.end_evt() + self.backend_mod.record(evt, self.stream) + self.backend_mod.synchronize(evt) + return evt + elif self.backend_mod_name == "AMDGPU": + evt = self.end_evt(self.stream, do_record=False, timing=True) + self.backend_mod.HIP.record(evt) + self.backend_mod.HIP.synchronize(evt) + return evt + elif self.backend_mod_name == "Metal": + jl.end_buf = self.create_metal_buffer() + jl.seval("Metal.commit!(end_buf)") + self.backend_mod.wait_completed(jl.end_buf) + return float(jl.end_buf.GPUStartTime) + + def kernel_finished(self): + """Returns True if the kernel has finished, False otherwise.""" + return True # JuliaCall synchronizes on record + + # @staticmethod + def synchronize(self): + try: + jl.Main.KernelAbstractions.synchronize(self.backend) + except JuliaError as e: + raise RuntimeError(f"Julia synchronize failed: {e}") + + # ------------------------- + # Memory utilities + # ------------------------- + + @staticmethod + def memset(allocation, value, size): + raise NotImplementedError("memset not yet implemented for Julia backend.") + # try: + # jl.allocation_tmp = allocation + # jl.seval(f"CUDA.fill!(allocation_tmp, {int(value)})") + # del jl.allocation_tmp + # except JuliaError as e: + # raise RuntimeError(f"Julia memset failed: {e}") + + @staticmethod + def memcpy_dtoh(dest, src): + """Perform a device to host memory copy.""" + try: + np.copyto(dest, src) + except JuliaError as e: + raise RuntimeError(f"Julia memcpy_dtoh failed: {e}") + + @staticmethod + def memcpy_htod(dest, src): + """Perform a host to device memory copy.""" + dest_ptr = repr(jl.UInt64(jl.pointer_from_objref(dest))) + jl.dest_tmp = dest + jl.src_tmp = src + jl.seval("copyto!(dest_tmp, Array(src_tmp))") + assert dest_ptr == repr(jl.UInt64(jl.pointer_from_objref(jl.dest_tmp))) + + def copy_constant_memory_args(self, cmem_args): + raise NotImplementedError( + "Constant memory not yet supported in Julia backend. Submit a feature request if needed." + ) + + def copy_shared_memory_args(self, smem_args): + raise NotImplementedError( + "Shared memory not yet supported in Julia backend. Submit a feature request if needed." + ) + # self.smem_size = int(smem_args.get("size", 0)) + + def copy_texture_memory_args(self, texmem_args): + raise NotImplementedError( + "Texture memory not yet supported in Julia backend. Submit a feature request if needed." + ) + + # ------------------------- + # Helper functions + # ------------------------- + + def check_package_and_install(self, package): + """Checks if the Julia package is available, and installs it if not.""" + try: + jl.seval(f"import {package}") + except Exception: + try: + warn(f"{package}.jl not found, attempting to install it directly.") + import juliapkg + juliapkg.add(package) + juliapkg.resolve() + jl.seval(f"import {package}") + except Exception as e: + raise ImportError( + f'{package}.jl not found in your Julia environment. Run `using Pkg; Pkg.add("{package}")` in Julia.' + ) from e + + def create_metal_buffer(self): + """Create a Metal buffer in the command queue.""" + try: + buf = self.contextqueue.commandBuffer() + except Exception: + buf = self.backend_mod.MTLCommandBuffer(self.contextqueue) + return buf + + def setup_streams(self, backend_name: str): + """Set up stream and event attributes for observers.""" + # Optional: common KernelAbstractions stream abstraction + try: + self.stream = self.backend_mod.get_default_stream() + except Exception: + self.stream = None + + # Set up stream and event attributes for observers + if backend_name == "CUDA": + self.stream = self.backend_mod.stream() + self.start_evt = self.backend_mod.CuEvent + self.end_evt = self.backend_mod.CuEvent + elif backend_name == "AMD": + self.stream = self.backend_mod.stream() + self.start_evt = self.backend_mod.HIP.HIPEvent + self.end_evt = self.backend_mod.HIP.HIPEvent + elif backend_name == "INTEL": + # OneAPI: no events available + self.start_evt = None + self.end_evt = None + elif backend_name == "METAL": + self.start_evt = self.start_event + self.end_evt = self.stop_event + elif backend_name == "CPU": + # CPU, use host-side timing + self.start_evt = None + self.end_evt = None + else: + raise NotImplementedError(f"Backend {backend_name} not supported in Julia backend.") + + def process_compiler_options(self, compiler_options=None): + """Process the given compiler options.""" + self.raise_errors = False + for c in compiler_options or []: + if c.lower().startswith("raise_errors="): + raise_errors_str = c.split("=", 1)[1].strip().lower() + if raise_errors_str in ("true", "1", "yes"): + self.raise_errors = True + elif raise_errors_str in ("false", "0", "no"): + self.raise_errors = False + else: + raise ValueError(f"Invalid value for raise_errors: {raise_errors_str}. Use true/false.") + compiler_options.remove(c) + + def verify_backends_with_options(self, compiler_options=None): + """Verify that the requested backend is available.""" + if compiler_options is not None and len(compiler_options) == 1: + if compiler_options[0].upper() not in self.available_backends: + raise ValueError( + f"Requested Julia backend '{compiler_options[0]}' not available. " + f"Available backends: {self.available_backends}" + ) + backend_name = compiler_options[0].upper() + else: + if len(self.available_backends) != 1: + if "CPU" in self.available_backends and len(self.available_backends) == 2: + self.available_backends.remove("CPU") + else: + raise ValueError( + f"Multiple or no Julia backends detected: {self.available_backends}. " + "Please specify exactly one backend in compiler_options." + ) + backend_name = self.available_backends[0] + return backend_name diff --git a/kernel_tuner/backends/julia_helper.jl b/kernel_tuner/backends/julia_helper.jl new file mode 100644 index 000000000..852421cf7 --- /dev/null +++ b/kernel_tuner/backends/julia_helper.jl @@ -0,0 +1,88 @@ +export to_gpuarray, launch_kernel + +function to_gpuarray(a) + a = deepcopy(a) # ensure we have a separate copy of the array to avoid unintended side effects + if isa(a, AbstractArray) + a = GPUArrayType(a) + end + return a +end + +function launch_kernel(kernel, args::Tuple, params::Tuple, ndrange::Tuple, workgroupsize::Tuple, shmem::Int, start_evt::Any, end_evt::Any, stream::Any) + t = Inf + # Check if this is a KernelAbstractions kernel + if isdefined(Main, :KernelAbstractions) && kt_julia_backend !== nothing && applicable(kernel, kt_julia_backend, workgroupsize) + configured_kernel = kernel(kt_julia_backend, workgroupsize) + # Launch kernel + mktemp() do tmppath, _ + open(tmppath, "w") do tmpio + # kernel errors are printed to stdout, capture them + redirect_stdout(tmpio) do + try + val_params = Val.(params) # convert parameters to Val types for kernel invocation + start_buff = nothing + end_buff = nothing + start = time_ns() # simple host-side timing as fallback in case of issues with GPU timing + if start_evt !== nothing + if isdefined(Main, :CUDA) && isa(start_evt, CuEvent) + Main.CUDA.record(start_evt, stream) + elseif isdefined(Main, :AMDGPU) && isa(start_evt, AMDGPU.HIP.HIPEvent) + Main.AMDGPU.HIP.record(start_evt) + elseif isdefined(Main, :Metal) + # prepare the next command buffers for timing as they can only be used once + start_buff = create_metal_buffer(Metal.device()) + end_buff = create_metal_buffer(Metal.device()) + Metal.commit!(start_buff) + else + error("Unsupported event type for timing: $(typeof(start_evt))") + end + end + configured_kernel(args..., val_params...; ndrange=ndrange) # launch the kernel + Main.KernelAbstractions.synchronize(kt_julia_backend) # synchronize to ensure kernel completion + if end_evt !== nothing + if isdefined(Main, :CUDA) && isa(end_evt, CuEvent) + Main.CUDA.record(end_evt, stream) + Main.CUDA.synchronize(end_evt) # ensure the event is recorded before we read it + elseif isdefined(Main, :AMDGPU) && isa(end_evt, AMDGPU.HIP.HIPEvent) + Main.AMDGPU.HIP.record(end_evt) + Main.AMDGPU.HIP.synchronize(end_evt) # ensure the event is recorded before we read it + elseif isdefined(Main, :Metal) + Metal.commit!(end_buff) + Metal.wait_completed(end_buff) # ensure the command buffer is completed before we read the time + t = (float(end_buff.GPUStartTime) - float(start_buff.GPUEndTime)) * 1000 + else + error("Unsupported event type for timing: $(typeof(end_evt))") + end + else + # host-side timing fallback if events are not available + t = float((time_ns() - start) / 1e6) # convert to milliseconds + end + catch e + redirect_stdout(stdout) # restore stdout + close(tmpio) + stdout_output = read(tmppath, String) + print("Kernel stdout during exception:\n", stdout_output) + # Rethrow the exception to be caught outside + throw(stdout_output * "\n" * sprint(showerror, e, catch_backtrace())) + end + end + end + # print any stdout output from the kernel + print(read(tmppath, String)) + end + else + error("Only KernelAbstractions kernels are supported.") + end + return t +end + +function create_metal_buffer(device) + # Create a Metal buffer in the command queue for timing + if isdefined(Main, :Metal) + contextqueue = Main.Metal.MTLCommandQueue(device) + return Metal.MTLCommandBuffer(contextqueue) + # return contextqueue.commandBuffer() + else + error("Metal backend is not available.") + end +end diff --git a/kernel_tuner/backends/julia_helper.py b/kernel_tuner/backends/julia_helper.py new file mode 100644 index 000000000..0b30bc7fd --- /dev/null +++ b/kernel_tuner/backends/julia_helper.py @@ -0,0 +1,134 @@ +"""Helper functions for Julia backend detection and interaction. + +We might want to consider moving this to a utility module or Julia package as it can be useful. +""" + +import subprocess +from json import JSONDecodeError +from json import loads as json_loads +from re import search as regex_search +from warnings import warn + +# Map name → Julia module and device-selection calls +backend_map = { + "CUDA": { + "pkg": "CUDA", + "module": "CUDA", + "module_backend": "CUDABackend", + "device_select": lambda d: f"CUDA.device!({d})", + "name": "CUDA.name(CUDA.device())", + "max_threads": "CUDA.attribute(CUDA.device(), CUDA.DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK)", + "capability": "CUDA.capability(CUDA.device())", + "GPUArrayType": "CuArray", + }, + "AMD": { + "pkg": "AMDGPU", + "module": "AMDGPU", + "module_backend": "ROCBackend", + "device_select": lambda d: f"AMDGPU.device!(AMDGPU.devices()[{d}])", + "name": "AMDGPU.HIP.name(AMDGPU.HIP.device())", + "max_threads": "AMDGPU.HIP.attribute(AMDGPU.HIP.device(), AMDGPU.HIP.hipDeviceAttributeMaxThreadsPerBlock)", + "capability": None, + "GPUArrayType": "ROCArray", + }, + "INTEL": { + "pkg": "oneAPI", + "module": "oneAPI", + "module_backend": "oneAPIBackend", + "device_select": lambda d: f"device!(devices(first(drivers()))[{d}])", + "name": "oneAPI.name(oneAPI.device())", + "max_threads": "oneAPI.compute_properties(oneAPI.device()).maxTotalGroupSize", + "capability": None, + "GPUArrayType": "oneArray", + }, + "METAL": { + "pkg": "Metal", + "module": "Metal", + "module_backend": "MetalBackend", + "device_select": lambda d: "Metal.device!(Metal.device())", # only single device support in Metal.jl + "name": "Metal.device().name", + "max_threads": "Int(Metal.device().maxThreadsPerThreadgroup.width)", + "capability": None, + "GPUArrayType": "MtlArray", + }, + "CPU": { + "pkg": None, + "module": "CPU", + "module_backend": "CPU", + "device_select": lambda d: "nothing", + "name": "CPU", + "max_threads": "1024", # arbitrary as CPU doesn't have a max threads per block concept + "capability": None, + "GPUArrayType": "Array", + }, +} + + +def detect_julia_gpu_backends(): + """Detect the Julia backends available.""" + available_backends = [] + if julia_backend_available_cuda(): + available_backends.append("CUDA") + if julia_backend_available_amd(): + available_backends.append("AMD") + if julia_backend_available_metal(): + available_backends.append("METAL") + if len(available_backends) == 0: + # this can give false positives for other backends too, so skip if we've already detected another backend + if julia_backend_available_intel(): + available_backends.append("INTEL") + + available_backends.append("CPU") # always add CPU backend last + return available_backends + + +def julia_backend_available_cuda(): + """Check if CUDA backend is available.""" + try: + subprocess.check_output("nvidia-smi") + return True + except (FileNotFoundError, subprocess.CalledProcessError): + return False + + +def julia_backend_available_amd(): + """Check if AMD backend is available.""" + try: + subprocess.check_output("rocm-smi") + return True + except (FileNotFoundError, subprocess.CalledProcessError): + return False + + +def julia_backend_available_metal(): + """Check if Metal backend is available.""" + try: + output = subprocess.check_output("system_profiler -json SPDisplaysDataType".split()) + json_output = json_loads(output)["SPDisplaysDataType"] + except (FileNotFoundError, subprocess.CalledProcessError, JSONDecodeError): + return False + for gpu in json_output: + if "spdisplays_mtlgpufamilysupport" in gpu: + supported = gpu["spdisplays_mtlgpufamilysupport"].lower() + if "metal" in supported: + version = regex_search(r".*metal([\d.]+)", supported).group(1) + if float(version) < 3: + warn( + f"Metal backend detected, but {supported} < 3. " + "Metal.jl requires Metal version 3 or higher." + ) + else: + return True + return False + + +def julia_backend_available_intel(): + """Check if Intel backend is available. May give false positives if other backends are present.""" + try: + # not a perfect check but should work in most cases + subprocess.check_output( + "ls /dev/dri/by-path/".split() + ) + return True + except (FileNotFoundError, subprocess.CalledProcessError): + return False diff --git a/kernel_tuner/backends/nvcuda.py b/kernel_tuner/backends/nvcuda.py index f55933c44..f9f28def3 100644 --- a/kernel_tuner/backends/nvcuda.py +++ b/kernel_tuner/backends/nvcuda.py @@ -1,5 +1,4 @@ """This module contains all NVIDIA cuda-python specific kernel_tuner functions.""" -from warnings import warn import numpy as np import uuid @@ -13,11 +12,13 @@ # embedded in try block to be able to generate documentation # and run tests without cuda-python installed try: - from cuda.bindings import driver, runtime, nvrtc + from cuda.bindings import driver, nvrtc, runtime except ImportError: try: # backward compatibility hack for older cuda-python versions - from cuda import cuda as driver, cudart as runtime, nvrtc as nvrtc + from cuda import cuda as driver + from cuda import cudart as runtime + from cuda import nvrtc as nvrtc except ImportError: driver = None diff --git a/kernel_tuner/backends/opencl.py b/kernel_tuner/backends/opencl.py index af3be1c00..5f7ea83f5 100644 --- a/kernel_tuner/backends/opencl.py +++ b/kernel_tuner/backends/opencl.py @@ -1,4 +1,5 @@ """This module contains all OpenCL specific kernel_tuner functions.""" + from __future__ import print_function import numpy as np @@ -16,9 +17,7 @@ class OpenCLFunctions(GPUBackend): """Class that groups the OpenCL functions on maintains some state about the device.""" - def __init__( - self, device=0, platform=0, iterations=7, compiler_options=None, observers=None - ): + def __init__(self, device=0, platform=0, iterations=7, compiler_options=None, observers=None): """Creates OpenCL device context and reads device properties. :param device: The ID of the OpenCL device to use for benchmarking @@ -37,14 +36,10 @@ def __init__( platforms = cl.get_platforms() self.ctx = cl.Context(devices=[platforms[platform].get_devices()[device]]) - self.queue = cl.CommandQueue( - self.ctx, properties=cl.command_queue_properties.PROFILING_ENABLE - ) + self.queue = cl.CommandQueue(self.ctx, properties=cl.command_queue_properties.PROFILING_ENABLE) self.mf = cl.mem_flags # inspect device properties - self.max_threads = self.ctx.devices[0].get_info( - cl.device_info.MAX_WORK_GROUP_SIZE - ) + self.max_threads = self.ctx.devices[0].get_info(cl.device_info.MAX_WORK_GROUP_SIZE) self.compiler_options = compiler_options or [] # observer stuff @@ -108,9 +103,7 @@ def compile(self, kernel_instance): :returns: An OpenCL kernel that can be called directly. :rtype: pyopencl.Kernel """ - prg = cl.Program(self.ctx, kernel_instance.kernel_string).build( - options=self.compiler_options - ) + prg = cl.Program(self.ctx, kernel_instance.kernel_string).build(options=self.compiler_options) func = getattr(prg, kernel_instance.name) return func @@ -136,7 +129,7 @@ def synchronize(self): """Halts execution until device has finished its tasks.""" self.queue.finish() - def run_kernel(self, func, gpu_args, threads, grid): + def run_kernel(self, func, gpu_args, threads, grid, stream=None): """Runs the OpenCL kernel passed as 'func'. :param func: An OpenCL Kernel diff --git a/kernel_tuner/backends/pycuda.py b/kernel_tuner/backends/pycuda.py index 9e3eb0d68..1fc8c6ab1 100644 --- a/kernel_tuner/backends/pycuda.py +++ b/kernel_tuner/backends/pycuda.py @@ -1,4 +1,5 @@ """This module contains all CUDA specific kernel_tuner functions.""" + from __future__ import print_function import logging @@ -97,13 +98,9 @@ def _finish_up(): PyCudaFunctions.last_selected_context = self.context # inspect device properties - devprops = { - str(k): v for (k, v) in self.context.get_device().get_attributes().items() - } + devprops = {str(k): v for (k, v) in self.context.get_device().get_attributes().items()} self.max_threads = devprops["MAX_THREADS_PER_BLOCK"] - cc = str(devprops.get("COMPUTE_CAPABILITY_MAJOR", "0")) + str( - devprops.get("COMPUTE_CAPABILITY_MINOR", "0") - ) + cc = str(devprops.get("COMPUTE_CAPABILITY_MAJOR", "0")) + str(devprops.get("COMPUTE_CAPABILITY_MINOR", "0")) if cc == "00": cc = self.context.get_device().compute_capability() self.cc = str(cc) @@ -351,14 +348,7 @@ def run_kernel(self, func, gpu_args, threads, grid, stream=None): """ if stream is None: stream = self.stream - func( - *gpu_args, - block=threads, - grid=grid, - stream=stream, - shared=self.smem_size, - texrefs=self.texrefs - ) + func(*gpu_args, block=threads, grid=grid, stream=stream, shared=self.smem_size, texrefs=self.texrefs) def memset(self, allocation, value, size): """Set the memory in allocation to the value in value. diff --git a/kernel_tuner/core.py b/kernel_tuner/core.py index b85e28068..4e30194ef 100644 --- a/kernel_tuner/core.py +++ b/kernel_tuner/core.py @@ -4,6 +4,7 @@ import re import time from collections import namedtuple +from warnings import warn import numpy as np @@ -28,7 +29,7 @@ def _get_cupy(): try: from hip._util.types import DeviceArray except ImportError: - DeviceArray = Exception # using Exception here as a type that will never be among kernel arguments + DeviceArray = Exception # using Exception here as a type that will never be among kernel arguments _KernelInstance = namedtuple( @@ -105,15 +106,13 @@ def get_kernel_string(self, index=0, params=None): """ logging.debug("get_kernel_string called") - if hasattr(self, 'lang') and self.lang.upper() == "HYPERTUNER": + if hasattr(self, "lang") and self.lang.upper() == "HYPERTUNER": return "" kernel_source = self.kernel_sources[index] return util.get_kernel_string(kernel_source, params) - def prepare_list_of_files( - self, kernel_name, params, grid, threads, block_size_names - ): + def prepare_list_of_files(self, kernel_name, params, grid, threads, block_size_names): """Prepare the kernel string along with any additional files. The first file in the list is allowed to include or read in the others @@ -201,7 +200,7 @@ def get_suffix(self, index=0): if suffix is not None: return suffix - _suffixes = {"CUDA": ".cu", "OpenCL": ".cl", "C": ".c"} + _suffixes = {"CUDA": ".cu", "OpenCL": ".cl", "C": ".c", "JULIA": ".jl"} try: return _suffixes[self.lang] except KeyError: @@ -214,10 +213,29 @@ def check_argument_lists(self, kernel_name, arguments): """ for i, f in enumerate(self.kernel_sources): if not callable(f): - util.check_argument_list(kernel_name, self.get_kernel_string(i), arguments) + util.check_argument_list(kernel_name, self.get_kernel_string(i), arguments, lang=self.lang) else: logging.debug("Checking of arguments list not supported yet for code generators.") + def infer_julia_backend(self): + """Infer the Julia backend from the kernel source.""" + backend = None + if self.lang.upper() != "JULIA": + return backend + + kernel_string = self.get_kernel_string(0) + if kernel_string.find("using CUDA") != -1: + backend = "cuda" + elif kernel_string.find("using ROCBackend") != -1: + backend = "amd" + elif kernel_string.find("using oneAPI") != -1: + backend = "intel" + elif kernel_string.find("using Metal") != -1: + backend = "metal" + else: + raise ValueError("Could not infer Julia backend from kernel source, provide it as a `compiler_option`") + return backend + def instantiate_observer(observer, args): """Instantiate or build an observer from a class/factory/instance.""" @@ -327,18 +345,21 @@ def __init__( elif lang.upper() == "HIP": from kernel_tuner.backends.hip import HipFunctions backend = HipFunctions + elif lang.upper() == "JULIA": + from kernel_tuner.backends.julia import JuliaFunctions + backend = JuliaFunctions + elif lang.upper() == "HYPERTUNER": + from kernel_tuner.backends.hypertuner import HypertunerFunctions + backend = HypertunerFunctions + self.requires_warmup = False elif lang.upper() in ["C", "FORTRAN"]: from kernel_tuner.backends.compiler import CompilerFunctions backend = CompilerFunctions backend_options["compiler"] = compiler backend_options["observers"] = observers - elif lang.upper() == "HYPERTUNER": - from kernel_tuner.backends.hypertuner import HypertunerFunctions - backend = HypertunerFunctions - self.requires_warmup = False else: raise NotImplementedError( - "Sorry, support for languages other than CUDA, OpenCL, HIP, C, and Fortran is not implemented yet" + "Sorry, support for languages other than CUDA, OpenCL, HIP, C, Julia, Fortran is not implemented yet" ) if issubclass(backend, GPUBackend): @@ -371,6 +392,12 @@ def __init__( if isinstance(obs, PrologueObserver): self.prologue_observers.append(obs) + # for JULIA, add the JIT warmup prologue observer + if lang.upper() == "JULIA": + from kernel_tuner.observers.julia import JuliaJITWarmup + self.prologue_observers.append(JuliaJITWarmup(self.dev.backend)) + self.prologue_observers.append(JuliaJITWarmup(self.dev.backend)) + # Take list of observers from self.dev because Backends tend to add their own observer self.benchmark_observers = [ obs for obs in self.dev.observers if not isinstance(obs, (ContinuousObserver, PrologueObserver)) @@ -418,7 +445,6 @@ def benchmark_default(self, func, gpu_args, threads, grid, result): for obs in self.benchmark_observers: result.update(obs.get_results()) - def benchmark_continuous(self, func, gpu_args, threads, grid, result, duration): """Benchmark continuously for at least 'duration' seconds.""" iterations = int(np.ceil(duration / (result["time"] / 1000))) @@ -442,7 +468,6 @@ def benchmark_continuous(self, func, gpu_args, threads, grid, result, duration): for obs in self.continuous_observers: result.update(obs.get_results()) - def set_nvml_parameters(self, instance): """Set the NVML parameters. Avoids setting time leaking into benchmark time.""" if self.use_nvml: @@ -461,7 +486,6 @@ def set_nvml_parameters(self, instance): if "tegra_gr_clock" in instance.params: self.tegra.gr_clock = instance.params["tegra_gr_clock"] - def benchmark(self, func, gpu_args, instance, verbose, objective, skip_nvml_setting=False): """Benchmark the kernel instance.""" logging.debug("benchmark " + instance.name) @@ -484,12 +508,11 @@ def benchmark(self, func, gpu_args, instance, verbose, objective, skip_nvml_sett self.benchmark_prologue(func, gpu_args, instance.threads, instance.grid, result) self.benchmark_default(func, gpu_args, instance.threads, instance.grid, result) - if self.continuous_observers: - duration = 1 - for obs in self.continuous_observers: - obs.results = result - duration = max(duration, obs.continuous_duration) - + duration = 1 + for obs in self.continuous_observers: + obs.results = result + duration = max(duration, obs.continuous_duration) + if len(self.continuous_observers) > 0: self.benchmark_continuous(func, gpu_args, instance.threads, instance.grid, result, duration) except Exception as e: @@ -501,10 +524,16 @@ def benchmark(self, func, gpu_args, instance, verbose, objective, skip_nvml_sett "too many resources requested for launch", "OUT_OF_RESOURCES", "INVALID_WORK_GROUP_SIZE", + "a bounds error was thrown during kernel execution", + "Julia kernel launch failed", ] if any([skip_str in str(e) for skip_str in skippable_exceptions]): - logging.debug("benchmark fails due to runtime failure too many resources required") - if verbose: + logging.debug("benchmark fails due to runtime failure / too many resources required") + if "julia" in str(e).lower() and verbose: + warn( + f"skipping config {util.get_instance_string(instance.params)} reason: Julia kernel launch failed because of:\n{e}" + ) + elif verbose: print( f"skipping config {util.get_instance_string(instance.params)} reason: too many resources requested for launch" ) @@ -518,23 +547,35 @@ def benchmark(self, func, gpu_args, instance, verbose, objective, skip_nvml_sett return result - def check_kernel_output( - self, func, gpu_args, instance, answer, atol, verify, verbose - ): + def check_kernel_output(self, func, gpu_args, instance, answer, atol, verify, verbose): """Runs the kernel once and checks the result against answer.""" logging.debug("check_kernel_output") + # get the answer for this parameter configuration + if isinstance(answer, Tunable): + answer = answer.select_for_configuration(instance.params) + + # convert juliacall arrays to numpy arrays where necessary + if answer is not None: + answer = [None if a is None else np.array(a) for a in util.possible_julia_vector_to_list(answer)] + for i, arg in enumerate(instance.arguments): + if isinstance(answer[i], np.ndarray) and "ArrayValue" in str(type(arg)): + instance.arguments[i] = np.array(arg, dtype=answer[i].dtype) + # if not using custom verify function, check if the length is the same if answer: if len(instance.arguments) != len(answer): raise TypeError("The length of argument list and provided results do not match.") - should_sync = [answer[i] is not None for i, arg in enumerate(instance.arguments)] + # for Julia arrays, we always want to sync + should_sync = [ + answer[i] is not None or "ArrayValue" in str(type(arg)) for i, arg in enumerate(instance.arguments) + ] else: cp = _get_cupy() cupy_ndarray = (cp.ndarray,) if cp is not None else () should_sync = [ - isinstance(arg, (np.ndarray, torch.Tensor, DeviceArray) + cupy_ndarray) + isinstance(arg, (np.ndarray, cp.ndarray, torch.Tensor, DeviceArray) + cupy_ndarray) or "ArrayValue" in str(type(arg)) for arg in instance.arguments ] @@ -549,27 +590,7 @@ def check_kernel_output( return # retrieve gpu results to host memory - result_host = [] - for i, arg in enumerate(instance.arguments): - if should_sync[i]: - cp = _get_cupy() - cupy_ndarray = (cp.ndarray,) if cp is not None else () - if isinstance(arg, (np.ndarray,) + cupy_ndarray): - result_host.append(np.zeros_like(arg)) - self.dev.memcpy_dtoh(result_host[-1], gpu_args[i]) - elif isinstance(arg, torch.Tensor) and isinstance(answer[i], torch.Tensor): - if not answer[i].is_cuda: - # if the answer is on the host, copy gpu output to host as well - result_host.append(torch.zeros_like(answer[i])) - self.dev.memcpy_dtoh(result_host[-1], gpu_args[i].tensor) - else: - result_host.append(gpu_args[i].tensor) - else: - # We should sync this argument, but we do not know how to transfer this type of argument - # What do we do? Should we throw an error? - result_host.append(None) - else: - result_host.append(None) + result_host = self.retrieve_results_to_host(instance.arguments, should_sync, gpu_args, answer) # Call the output observers for obs in self.output_observers: @@ -683,9 +704,7 @@ def compile_kernel(self, instance, verbose): ] error_message = str(e.stderr) if hasattr(e, "stderr") else str(e) if any(re.search(msg, error_message) for msg in shared_mem_error_messages): - logging.debug( - "compile_kernel failed due to kernel using too much shared memory" - ) + logging.debug("compile_kernel failed due to kernel using too much shared memory") if verbose: print( f"skipping config {util.get_instance_string(instance.params)} reason: too much shared memory used" @@ -808,6 +827,32 @@ def run_kernel(self, func, gpu_args, instance): raise e return True + def retrieve_results_to_host(self, arguments: list, should_sync: list[bool], gpu_args, answer: list): + """Retrieve results from device to host memory for all arguments that should be synchronized.""" + result_host = [] + for i, arg in enumerate(arguments): + if not should_sync[i]: + result_host.append(None) + continue + cp = _get_cupy() + cupy_ndarray = (cp.ndarray,) if cp is not None else () + if isinstance(arg, (np.ndarray,) + cupy_ndarray) or arg.__class__.__name__ == "VectorValue": + result_host.append(np.zeros_like(arg)) + self.dev.memcpy_dtoh(result_host[-1], gpu_args[i]) + elif isinstance(arg, torch.Tensor) and isinstance(answer[i], torch.Tensor): + if not answer[i].is_cuda: + # if the answer is on the host, copy gpu output to host as well + result_host.append(torch.zeros_like(answer[i])) + self.dev.memcpy_dtoh(result_host[-1], gpu_args[i].tensor) + else: + result_host.append(gpu_args[i].tensor) + else: + # We should sync this argument, but we do not know how to transfer this type of argument + # What do we do? Should we throw an error? + warn(f"Argument {i} is of type {type(arg)} and should be synchronized, but is not implemented.") + result_host.append(None) + return result_host + def _preprocess_gpu_arguments(old_arguments, params): """Get a flat list of arguments based on the configuration given by `params`.""" @@ -830,11 +875,15 @@ def _default_verify_function(instance, answer, result_host, atol, verbose): # for each element in the argument list, check if the types match for i, arg in enumerate(instance.arguments): if answer[i] is not None: # skip None elements in the answer list + # convert Julia VectorValues to numpy arrays for verification + if arg.__class__.__name__ == "VectorValue": + arg = np.array(arg) + if answer[i].__class__.__name__ == "VectorValue": + answer[i] = np.array(answer[i]) + cp = _get_cupy() cupy_ndarray = (cp.ndarray,) if cp is not None else () - if isinstance(answer[i], (np.ndarray,) + cupy_ndarray) and isinstance( - arg, (np.ndarray,) + cupy_ndarray - ): + if isinstance(answer[i], (np.ndarray,) + cupy_ndarray) and isinstance(arg, (np.ndarray,) + cupy_ndarray): if not np.can_cast(arg.dtype, answer[i].dtype): raise TypeError( f"Element {i} of the expected results list has a dtype that is not compatible with the dtype of the kernel output: " @@ -844,7 +893,7 @@ def _default_verify_function(instance, answer, result_host, atol, verbose): + "." ) if answer[i].size != arg.size: - raise TypeError( + raise ValueError( f"Element {i} of the expected results list has a size different from " + "the kernel argument: " + str(answer[i].size) @@ -862,7 +911,7 @@ def _default_verify_function(instance, answer, result_host, atol, verbose): + "." ) if answer[i].size() != arg.size(): - raise TypeError( + raise ValueError( f"Element {i} of the expected results list has a size different from " + "the kernel argument: " + str(answer[i].size) @@ -888,7 +937,7 @@ def _default_verify_function(instance, answer, result_host, atol, verbose): answer[i], np.number ): raise TypeError( - f"Element {i} of expected results list is not a numpy/cupy ndarray, torch Tensor or numpy scalar." + f"Arg or Element {i} of expected results list is {type(arg)} / {type(answer[i])}, not a numpy/cupy ndarray, torch Tensor or numpy scalar." # noqa: E501 ) else: raise TypeError(f"Element {i} of expected results list and kernel arguments have different types.") @@ -910,31 +959,34 @@ def _flatten(a): result = _ravel(result_host[i]) expected = _flatten(expected) cp = _get_cupy() - if cp is not None and any([isinstance(array, cp.ndarray) for array in [expected, result]]): - output_test = cp.allclose(expected, result, atol=atol) - elif isinstance(expected, torch.Tensor) and isinstance(result, torch.Tensor): - output_test = torch.allclose(expected, result, atol=atol) - else: - output_test = np.allclose(expected, result, atol=atol) + has_cp_array = False if not cp else any([isinstance(array, cp.ndarray) for array in [expected, result]]) + lib = cp if has_cp_array else torch if isinstance(expected, torch.Tensor) and isinstance(result, torch.Tensor) else np + expected_nan = lib.isnan(expected) + output_test = lib.allclose(expected, result, atol=atol, equal_nan=expected_nan.any()) + if expected_nan.any(): + warn( + f"Answer contains {expected_nan.sum()} NaNs. NaN values will now be considered equal in comparison." + ) if not output_test and verbose: - print( - "Error: " - + util.get_config_string(instance.params) - + " detected during correctness check" - ) - print( - "this error occurred when checking value of the %oth kernel argument" - % (i,) - ) - print( - "Printing kernel output and expected result, set verbose=False to suppress this debug print" - ) - np.set_printoptions(edgeitems=50) - print("Kernel output:") + print("Error: " + util.get_config_string(instance.params) + " detected during correctness check") + print("this error occurred when checking value of the %oth kernel argument" % (i,)) + print("Printing kernel output and expected result, set verbose=False to suppress this debug print") + np.set_printoptions(edgeitems=30) + print(f"Kernel output ({np.shape(result)}):") print(result) - print("Expected:") + print(f"Expected ({np.shape(expected)}):") print(expected) + # check if there are NaNs in the output or expected, if so, print where they are + if lib.isnan(result).any(): + print("NaNs in kernel output at indices:", lib.where(lib.isnan(result))) + if lib.isnan(expected).any(): + print("NaNs in expected result at indices:", lib.where(lib.isnan(expected))) + # print only the elements that are different + print("Difference at specific elements:") + diff = lib.abs(expected - result) + indices = lib.where(diff > atol) + print(diff[indices]) correct = correct and output_test if not correct: diff --git a/kernel_tuner/interface.py b/kernel_tuner/interface.py index a1d14bf0c..387aec344 100644 --- a/kernel_tuner/interface.py +++ b/kernel_tuner/interface.py @@ -28,14 +28,13 @@ import importlib from argparse import ArgumentParser from ast import literal_eval +from copy import deepcopy from datetime import datetime import os from pathlib import Path from time import perf_counter -from copy import deepcopy import numpy -from constraint import Constraint import kernel_tuner.core as core import kernel_tuner.util as util @@ -633,6 +632,22 @@ def tune_kernel( kernelsource = core.KernelSource(kernel_name, kernel_source, lang, defines) + # Convert Julia types + if lang is not None and lang.upper() == "JULIA": + # TODO implement & test the case where Kernel Tuner is called from Julia but the target language is not Julia + tune_params = util.julia_list_of_pairs_to_dict(tune_params) + if answer is not None: + answer = [ + numpy.array(a) if isinstance(a, (list, tuple)) else a for a in util.possible_julia_vector_to_list(answer) + ] + grid_div_x = util.possible_julia_vector_to_list(grid_div_x) + grid_div_y = util.possible_julia_vector_to_list(grid_div_y) + grid_div_z = util.possible_julia_vector_to_list(grid_div_z) + if strategy_options is not None: + strategy_options = dict([tuple([k, util.possible_julia_vector_to_list(o)]) for k, o in strategy_options]) + restrictions = util.possible_julia_vector_to_list(restrictions) + block_size_names = util.possible_julia_vector_to_list(block_size_names) + _check_user_input(kernel_name, kernelsource, arguments, block_size_names) if objectives: @@ -661,6 +676,14 @@ def tune_kernel( # ensure there is always at least three names util.append_default_block_size_names(block_size_names) + # if Julia, infer the Julia backend from the kernelsource + if kernelsource.lang == "JULIA": + if compiler_options is None: + try: + compiler_options = [kernelsource.infer_julia_backend()] + except ValueError: + pass + # sort all the options into separate dicts opts = locals() kernel_options = Options([(k, opts[k]) for k in _kernel_options.keys()]) @@ -708,7 +731,6 @@ def tune_kernel( strategy = strategy_map["brute_force"] # select the runner for this job based on input - # TODO: we could use the "match case" syntax when removing support for 3.9 tuning_options.simulated_time = 0 # Get runner from environment if possible @@ -859,7 +881,8 @@ def tune_cache( return tune_kernel(**tune_args, cache=cache_path, restrictions=_restrictions, simulation_mode=True) -_run_kernel_docstring = """Compile and run a single kernel +_run_kernel_docstring = ( + """Compile and run a single kernel Compiles and runs a single kernel once, given a specific instance of the kernels tuning parameters. However, instead of measuring execution time run_kernel returns the output of the kernel. @@ -885,10 +908,9 @@ def tune_cache( :returns: A list of numpy arrays, similar to the arguments passed to this function, containing the output after kernel execution. :rtype: list -""" % _get_docstring( - _kernel_options -) + _get_docstring( - _device_options +""" + % _get_docstring(_kernel_options) + + _get_docstring(_device_options) ) @@ -919,6 +941,12 @@ def run_kernel( kernelsource = core.KernelSource(kernel_name, kernel_source, lang, defines) + if lang is not None and lang.upper() == "JULIA": + params = util.julia_list_of_pairs_to_dict(params) + block_size_names = util.possible_julia_vector_to_list(block_size_names) + # ensure there is always at least three names + util.append_default_block_size_names(block_size_names) + _check_user_input(kernel_name, kernelsource, arguments, block_size_names) # sort options into separate dicts @@ -943,7 +971,7 @@ def run_kernel( raise RuntimeError("cannot create kernel instance, too many threads per block") # see if the kernel arguments have correct type - util.check_argument_list(instance.name, instance.kernel_string, arguments) + util.check_argument_list(instance.name, instance.kernel_string, arguments, lang=lang) # compile the kernel func = dev.compile_kernel(instance, False) @@ -1002,8 +1030,8 @@ def tune_kernel_T1( output_T4=True, iterations=7, device=None, - strategy: str=None, - strategy_options: dict={}, + strategy: str = None, + strategy_options: dict = {}, ) -> tuple: """Call the tune function with a T1 input file. @@ -1013,12 +1041,8 @@ def tune_kernel_T1( kernelspec: dict = inputs["KernelSpecification"] kernel_name: str = kernelspec["KernelName"] kernel_filepath = Path(kernelspec["KernelFile"]) - kernel_source = ( - kernel_filepath if kernel_filepath.exists() else Path(input_filepath).parent / kernel_filepath - ) - kernel_source = ( - kernel_source if kernel_source.exists() else Path(input_filepath).parent.parent / kernel_filepath - ) + kernel_source = kernel_filepath if kernel_filepath.exists() else Path(input_filepath).parent / kernel_filepath + kernel_source = kernel_source if kernel_source.exists() else Path(input_filepath).parent.parent / kernel_filepath assert kernel_source.exists(), f"KernelFile '{kernel_source}' does not exist at {kernel_source.resolve()}" language: str = kernelspec["Language"] problem_size = kernelspec["ProblemSize"] @@ -1043,10 +1067,12 @@ def tune_kernel_T1( # if it is a path, import the strategy from the file opt_path: Path = Path(strategy_options["custom_search_method_path"]) class_name: str = strategy - assert opt_path.exists(), f"Custom search method path '{opt_path}' does not exist relative to current working directory {Path.cwd()}" + assert opt_path.exists(), ( + f"Custom search method path '{opt_path}' does not exist relative to current working directory {Path.cwd()}" + ) optimizer_class = import_class_from_file(opt_path, class_name) filter_keys = ["custom_search_method_path", "max_fevals", "time_limit", "constraint_aware"] - adjusted_strategy_options = {k:v for k, v in strategy_options.items() if k not in filter_keys} + adjusted_strategy_options = {k: v for k, v in strategy_options.items() if k not in filter_keys} optimizer_instance = optimizer_class(**adjusted_strategy_options) strategy = OptAlgWrapper(optimizer_instance) if "constraint_aware" not in strategy_options and hasattr(optimizer_instance, "constraint_aware"): diff --git a/kernel_tuner/observers/hip.py b/kernel_tuner/observers/hip.py index 5e27fb1b5..311d54376 100644 --- a/kernel_tuner/observers/hip.py +++ b/kernel_tuner/observers/hip.py @@ -1,7 +1,7 @@ import numpy as np -from kernel_tuner.observers.observer import BenchmarkObserver from kernel_tuner.backends.hip.util import hip_check +from kernel_tuner.observers.observer import BenchmarkObserver try: from hip import hip, hiprtc @@ -15,7 +15,9 @@ class HipRuntimeObserver(BenchmarkObserver): def __init__(self, dev): if not hip or not hiprtc: - raise ImportError("Unable to import HIP Python, or check https://kerneltuner.github.io/kernel_tuner/stable/install.html#hip-and-hip-python.") + raise ImportError( + "Unable to import HIP Python, or check https://kerneltuner.github.io/kernel_tuner/stable/install.html#hip-and-hip-python." + ) self.dev = dev self.stream = dev.stream diff --git a/kernel_tuner/observers/julia.py b/kernel_tuner/observers/julia.py new file mode 100644 index 000000000..16fa11991 --- /dev/null +++ b/kernel_tuner/observers/julia.py @@ -0,0 +1,130 @@ +from time import perf_counter +from warnings import warn + +import numpy as np + +from kernel_tuner.observers.observer import BenchmarkObserver, PrologueObserver + + +class JuliaRuntimeObserver(BenchmarkObserver): + """Cross-backend GPU timing for KernelAbstractions. + + - CUDA: CuEvent timing + - ROCBackend: HIPEvent timing + - OneAPI: host-side timing (less accurate, no events available) + - Metal: timing by wrapping the kernel launch between two command buffers and measuring the time between them + """ + + def __init__( + self, + kernelabstractions, + kt_backend, + jl_backend, + jl_backend_mod, + jl_backend_name, + stream=None, + start_event=None, + end_event=None, + ): + """Observer that measures GPU time depending on the Julia backend used.""" + self.kernelabstractions = kernelabstractions + self.kt_backend = kt_backend + self.backend = jl_backend + self.backend_mod = jl_backend_mod + self.name = jl_backend_name.lower() + self.stream = stream + self.start = start_event + self.end = end_event + self.times = [] + self.t0 = None + + if self.name == "cuda": + # initialize events for this instance of the observer + self.stream = self.backend_mod.stream() + self.start = self.start() + self.end = self.end() + elif self.name == "amdgpu": + self.stream = self.backend_mod.stream() + self.start = self.start(self.stream, do_record=False, timing=True) + self.end = self.end(self.stream, do_record=False, timing=True) + + def before_start(self): + if self.start is not None: + if self.name == "metal": + self.t0 = self.start() + elif self.name == "cuda": + # the events are recorded in the julia_helper kernel launch code + pass + elif self.name == "amdgpu": + # the events are recorded in the julia_helper kernel launch code + pass + else: + raise ValueError(f"Unsupported backend for timing: {self.name}") + else: + # fallback: host-side timestamp + self.t0 = perf_counter() + + def after_finish(self): + ms = None + if self.end is not None: + if self.name == "metal": + ms_observer = float((self.end() - self.t0) * 1000.0) + ms_helper = self.kt_backend.host_time + # take the minimum of the two measurements to mitigate overhead of command buffer timing + ms = min( + ms_observer, ms_helper + ) + elif self.name == "cuda": + # the events are recorded in the julia_helper kernel launch code + ms = float(self.backend_mod.elapsed(self.start, self.end) * 1000.0) + elif self.name == "amdgpu": + # the events are recorded in the julia_helper kernel launch code + ms = float(self.backend_mod.HIP.elapsed(self.start, self.end) * 1000.0) + else: + raise ValueError(f"Unsupported backend for timing: {self.name}") + else: + self.kernelabstractions.synchronize(self.backend) + dt = perf_counter() - self.t0 + ms = dt * 1000.0 + warn(f"Using host-side timing for Julia {self.name} backend; results may be less accurate.") + + if ms > self.kt_backend.host_time: + if ms < 1 and (ms > 1.5 * self.kt_backend.host_time and self.end is not None): + warn( + f"Measured GPU time {ms:.3f} ms is greater than host time {self.kt_backend.host_time:.3f} ms; " + "this may happen with very short execution times." + ) + elif ms > 1.5 * self.kt_backend.host_time and self.end is not None: + warn( + f"Measured GPU time {ms:.3f} ms is substantially greater than host time {self.kt_backend.host_time:.3f} ms; " + "this may indicate an issue with the timing measurement." + ) + + self.times.append(ms) + + def get_results(self): + results = { + "time": np.average(self.times), + "times": self.times.copy(), + } + self.times = [] + return results + + +class JuliaJITWarmup(PrologueObserver): + """Prologue observer to enforce warmup before every configuration to trigger JIT.""" + + def __init__(self, backend): + """Not implemented, just to trigger JIT.""" + pass + + def before_start(self): + """Not implemented, just to trigger JIT.""" + pass + + def after_finish(self): + """Not implemented, just to trigger JIT.""" + pass + + def get_results(self): + return {} diff --git a/kernel_tuner/runners/sequential.py b/kernel_tuner/runners/sequential.py index a4f2ab8f7..c987524f7 100644 --- a/kernel_tuner/runners/sequential.py +++ b/kernel_tuner/runners/sequential.py @@ -1,4 +1,5 @@ """The default runner for sequentially tuning the parameter space.""" + import logging from datetime import datetime, timezone from time import perf_counter diff --git a/kernel_tuner/util.py b/kernel_tuner/util.py index e074f8b92..b68eb04fb 100644 --- a/kernel_tuner/util.py +++ b/kernel_tuner/util.py @@ -1,4 +1,5 @@ """Module for kernel tuner utility functions.""" + import ast from datetime import timedelta import errno @@ -12,6 +13,10 @@ import time import warnings from inspect import getsource, signature +from math import ( + ceil, # noqa: F401 + floor, # noqa: F401 +) # importing here to make available for eval in restrictions / metrics / problem size etc. from pathlib import Path from types import FunctionType from typing import Union @@ -157,56 +162,52 @@ def check_argument_type(dtype, kernel_argument): return False # unknown dtype. do not throw exception to still allow kernel to run. -def check_argument_list(kernel_name, kernel_string, args): +def check_argument_list(kernel_name, kernel_string, args, lang=None): """Raise an exception if kernel arguments do not match host arguments.""" - cp = _get_cupy() - cupy_ndarray = (cp.ndarray,) if cp is not None else () kernel_arguments = list() collected_errors = list() + # Find all kernel argument lists in the kernel string + if lang and lang.upper() == "JULIA": + # for Julia, multiple kernels may be specified in one file, and remove the normally included tunable parameters + kernel_string = kernel_string.split(kernel_name)[1] + kernel_string = kernel_name + kernel_string + kernel_string = kernel_string.split("::Val")[0].rstrip() for iterator in re.finditer(kernel_name + "[ \n\t]*" + r"\(", kernel_string): kernel_start = iterator.end() - kernel_end = kernel_string.find(")", kernel_start) + # for Julia, search until the last ',' as there may be e.g. `@Const(Min)` arguments + kernel_end = kernel_string.find(r".*(\,)" if lang and lang.upper() == "JULIA" else ")", kernel_start) if kernel_start != 0: kernel_arguments.append(kernel_string[kernel_start:kernel_end].split(",")) + # Check each set of kernel arguments for arguments_set, arguments in enumerate(kernel_arguments): + + # check arguments and signature lengths + if lang and lang.upper() == "JULIA" and len(arguments) > len(args): + # for Julia additional parameters may be passed + continue collected_errors.append(list()) if len(arguments) != len(args): - collected_errors[arguments_set].append("Kernel and host argument lists do not match in size.") + collected_errors[arguments_set].append( + f"Kernel ({len(arguments)}) and host argument ({len(args)}) lists do not match in size." + ) + continue + + # skip checking for Julia, as types are commonly not specified in the kernel arguments + if lang and lang.upper() == "JULIA": + collected_errors.pop(arguments_set) continue + # Check each argument in the kernel argument list for i, arg in enumerate(args): kernel_argument = arguments[i] - - # Handle tunable arguments - if isinstance(arg, Tunable): - continue - - # Handle numpy arrays and other array types - if not isinstance(arg, (np.ndarray, np.generic, torch.Tensor, DeviceArray) + cupy_ndarray): - raise TypeError( - f"Argument at position {i} of type: {type(arg)} should be of type " - "np.ndarray, numpy scalar, or HIP Python DeviceArray type" + correct, str_dtype = check_individual_arguments(i, arg, kernel_argument) + if not correct: + collected_errors[arguments_set].append( + f"Argument at position {str(i)} of dtype: {str_dtype} does not match {kernel_argument}." ) - correct = True - if isinstance(arg, np.ndarray): - if "*" not in kernel_argument: - correct = False - - if isinstance(arg, DeviceArray): - str_dtype = str(np.dtype(arg.typestr)) - else: - str_dtype = str(arg.dtype) - - if correct and check_argument_type(str_dtype, kernel_argument): - continue - - collected_errors[arguments_set].append( - f"Argument at position {i} of dtype: {str_dtype} does not match {kernel_argument}." - ) - if not collected_errors[arguments_set]: # We assume that if there is a possible list of arguments that matches with the provided one # it is the right one @@ -216,6 +217,40 @@ def check_argument_list(kernel_name, kernel_string, args): warnings.warn(errors[0], UserWarning) +def check_individual_arguments(i, arg, kernel_argument): + """Check whether the host argument matches the kernel argument.""" + correct = True + cp = _get_cupy() + cupy_ndarray = (cp.ndarray,) if cp is not None else () + + # Handle tunable arguments + if isinstance(arg, Tunable): + return correct, "" + + # Handle numpy arrays and other array types + if not isinstance(arg, (np.ndarray, np.generic, torch.Tensor, DeviceArray) + cupy_ndarray): + if arg.__class__.__name__ == "VectorValue": + # skip for Julia, types are commonly not specified in the kernel arguments + return correct, "" + raise TypeError( + f"Argument at position {i} of type: {type(arg)} should be of type " + "np.ndarray, numpy scalar, HIP Python DeviceArray, Julia VectorValue type" + ) + + if isinstance(arg, np.ndarray) and "*" not in kernel_argument: + correct = False + + if isinstance(arg, DeviceArray): + str_dtype = str(np.dtype(arg.typestr)) + else: + str_dtype = str(arg.dtype) + + if correct: + return check_argument_type(str_dtype, kernel_argument), str_dtype + + return False, str_dtype + + class Timer: """Measures elapsed wall-clock time.""" def __init__(self): @@ -338,11 +373,15 @@ def check_block_size_names(block_size_names): if block_size_names is not None: # do some type checks for the user input if not isinstance(block_size_names, list): - raise ValueError("block_size_names should be a list of strings!") + raise ValueError("block_size_names should be a list of strings!", block_size_names, type(block_size_names)) if len(block_size_names) > 3: - raise ValueError("block_size_names should not contain more than 3 names!") + raise ValueError("block_size_names should not contain more than 3 names!", block_size_names) if not all([isinstance(name, "".__class__) for name in block_size_names]): - raise ValueError("block_size_names should contain only strings!") + raise ValueError( + "block_size_names should contain only strings!", + block_size_names, + [type(name) for name in block_size_names], + ) def append_default_block_size_names(block_size_names): @@ -378,7 +417,6 @@ def check_block_size_params_names_list(block_size_names, tune_params): return block_size_names - def check_restriction(restrict, params: dict) -> bool: """Check whether a configuration meets a search space restriction.""" # if it's a function python-constraint it can be called directly @@ -521,10 +559,13 @@ def delete_temp_file(filename): def detect_language(kernel_string): """Attempt to detect language from the kernel_string.""" + kernel_string = kernel_string.lower().strip() if "__global__" in kernel_string: lang = "CUDA" elif "__kernel" in kernel_string: lang = "OpenCL" + elif any(token in kernel_string for token in ["@cuda", "@kernel", "@device_code"]): + lang = "Julia" else: lang = "C" return lang @@ -547,7 +588,7 @@ def get_pareto_results( objective_higher_is_better: list[bool], mark_optima=True ): - from pymoo.util.nds.find_non_dominated import find_non_dominated + from pymoo.util.nds.non_dominated_sorting import find_non_dominated assert isinstance(results, list) assert isinstance(objectives, list) @@ -644,7 +685,9 @@ def get_dimension_divisor(divisor, default, params): for div in divisor: divisor_num *= get_dimension_divisor(div, 1, params) else: - raise ValueError("Error: unrecognized type in grid divisor list, should be any of int, str, callable, or iterable") + raise ValueError( + "Error: unrecognized type in grid divisor list, should be any of int, str, callable, or iterable" + ) return divisor_num @@ -757,9 +800,7 @@ def get_smem_args(smem_args, params): def get_temp_filename(suffix=None): """Return a string in the form of temp_X, where X is a large integer.""" - tmp_file = tempfile.mkstemp( - suffix=suffix or "", prefix="temp_", dir=os.getcwd() - ) # or "" for Python 2 compatibility + tmp_file = tempfile.mkstemp(suffix=suffix, prefix="temp_", dir=os.getcwd()) os.close(tmp_file[0]) return tmp_file[1] @@ -901,8 +942,8 @@ def looks_like_a_filename(kernel_source): for s in ["__global__ ", "__kernel ", "void ", "float "]: if s in kernel_source: result = False - # string must contain substring ".c", ".opencl", or ".F" - result = result and any([s in kernel_source for s in (".c", ".opencl", ".F")]) + # string must contain substring ".c", ".opencl", ".F", ".jl" + result = result and any([s in kernel_source for s in (".c", ".opencl", ".F", ".jl")]) logging.debug("kernel_source is a filename: %s" % str(result)) return result @@ -993,12 +1034,18 @@ def prepare_kernel_string(kernel_name, kernel_string, params, grid, threads, blo kernel_string = re.sub(r"\n\s*#pragma\s+unroll\s+" + k, "\n", kernel_string) # + r"[^\S]*" else: kernel_prefix += f"constexpr int {k} = {v};\n" + elif lang.upper() == "JULIA": + # kernel_prefix += f"const {k} = {v}\n" + # in Julia, we can't redefine constants like this, so we skip it and give it as arguments on the kernel launch + pass else: kernel_prefix += f"#define {k} {v}\n" # since we insert defines above the original kernel code, the line numbers will be incorrect # the following preprocessor directive informs the compiler that lines should be counted from 1 - if kernel_prefix: + if lang.upper() == "JULIA": + kernel_prefix += "\n" + elif kernel_prefix: kernel_prefix += "#line 1\n" # Also replace parameter occurrences inside the kernel name @@ -1206,7 +1253,7 @@ def get_all_lambda_asts(func): if not res: raise ValueError(f"No lambda node found in the source {source}.") except SyntaxError: - """ Ignore syntax errors on the lambda """ + """Ignore syntax errors on the lambda""" return res except OSError: raise ValueError("Could not retrieve source. Is this defined interactively or dynamically?") @@ -1217,15 +1264,18 @@ class ConstraintLambdaTransformer(ast.NodeTransformer): """Replaces any `NAME['string']` subscript with just `'string'`, if `NAME` matches the lambda argument name. """ + def __init__(self, dict_arg_name): self.dict_arg_name = dict_arg_name def visit_Subscript(self, node): # We only replace subscript expressions of the form ['some_string'] - if (isinstance(node.value, ast.Name) - and node.value.id == self.dict_arg_name - and isinstance(node.slice, ast.Constant) - and isinstance(node.slice.value, str)): + if ( + isinstance(node.value, ast.Name) + and node.value.id == self.dict_arg_name + and isinstance(node.slice, ast.Constant) + and isinstance(node.slice.value, str) + ): # Replace `dict_arg_name['some_key']` with the string used as key return ast.Name(node.slice.value) return self.generic_visit(node) @@ -1261,7 +1311,7 @@ def convert_constraint_lambdas(restrictions): try: lambda_asts = get_all_lambda_asts(c) except ValueError: - res.append(c) # it's just a plain function, not a lambda + res.append(c) # it's just a plain function, not a lambda continue for lambda_ast in lambda_asts: @@ -1270,7 +1320,9 @@ def convert_constraint_lambdas(restrictions): result = list(set(res)) if not len(result) == len(restrictions): - raise ValueError("An error occured when parsing restrictions. If you mix lambdas and string-based restrictions, please define the lambda first.") + raise ValueError( + "An error occured when parsing restrictions. If you mix lambdas and string-based restrictions, please define the lambda first." + ) return result @@ -1319,12 +1371,16 @@ def compile_restrictions( noncompiled_restrictions.append((r, [], r)) return noncompiled_restrictions + compiled_restrictions + def check_matching_problem_size(cached_problem_size, problem_size): """Check the if requested problem size matches the problem size in the cache.""" cached_problem_size_arr = np.array(cached_problem_size) problem_size_arr = np.array(problem_size) if cached_problem_size_arr.size != problem_size_arr.size or not (cached_problem_size_arr == problem_size_arr).all(): - raise ValueError(f"Cannot load cache which contains results for different problem_size, cache: {cached_problem_size}, requested: {problem_size}") + raise ValueError( + f"Cannot load cache which contains results for different problem_size, cache: {cached_problem_size}, requested: {problem_size}" + ) + def process_cache(cachefile, kernel_options, tuning_options, runner): """Cache file for storing tuned configurations. @@ -1500,6 +1556,25 @@ def dump_cache(obj: str, tuning_options): cachefile.write(obj) +def possible_julia_vector_to_list(obj): + """Convert a Julia vector to a Python list if needed.""" + if obj.__class__.__name__ == "VectorValue": + l = list(obj) + l = [possible_julia_vector_to_list(e) for e in l] + return l + return obj + + +def julia_list_of_pairs_to_dict(params): + if isinstance(params, dict) or "DictValue" in params.__class__.__name__: + raise ValueError( + f"params {params} should not be a Julia dict, because it does not preserve order. Use a list of pairs instead." + ) + params = [tuple([k, possible_julia_vector_to_list(tp)]) for k, tp in params] + params = dict(params) + return params + + def infer_restrictions_from_cache(cache: dict): param_names = cache["tune_params_keys"] valid_param_config_set = set( diff --git a/kernel_tuner/utils/directives.py b/kernel_tuner/utils/directives.py index 36cd219b0..f3930d6b7 100644 --- a/kernel_tuner/utils/directives.py +++ b/kernel_tuner/utils/directives.py @@ -1,5 +1,8 @@ -from typing import Any, Tuple +"""Utility functions and classes for handling directives.""" + from abc import ABC, abstractmethod +from typing import Any, Tuple + import numpy as np # Function templates @@ -34,80 +37,82 @@ class Directive(ABC): - """Base class for all directives""" + """Base class for all directives.""" @abstractmethod - def get(self) -> str: + def get(self) -> str: # noqa: D102 pass class Language(ABC): - """Base class for all languages""" + """Base class for all languages.""" @abstractmethod - def get(self) -> str: + def get(self) -> str: # noqa: D102 pass class OpenACC(Directive): - """Class to represent OpenACC""" + """Class to represent OpenACC.""" - def get(self) -> str: + def get(self) -> str: # noqa: D102 return "openacc" class OpenMP(Directive): - """Class to represent OpenMP""" + """Class to represent OpenMP.""" - def get(self) -> str: + def get(self) -> str: # noqa: D102 return "openmp" class Cxx(Language): - """Class to represent C++ code""" + """Class to represent C++ code.""" - def get(self) -> str: + def get(self) -> str: # noqa: D102 return "cxx" - def end_string(self) -> str: + def end_string(self) -> str: # noqa: D102 return "#pragma tuner stop" class Fortran(Language): - """Class to represent Fortran code""" + """Class to represent Fortran code.""" - def get(self) -> str: + def get(self) -> str: # noqa: D102 return "fortran" - def end_string(self) -> str: + def end_string(self) -> str: # noqa: D102 return "!$tuner stop" class Code(object): - """Class to represent the directive and host code of the application""" + """Class to represent the directive and host code of the application.""" - def __init__(self, directive: Directive, lang: Language): + def __init__(self, directive: Directive, lang: Language): # noqa: D107 self.directive = directive self.language = lang class ArraySize(object): - """Size of an array""" + """Size of an array.""" - def __init__(self): + def __init__(self): # noqa: D102, D107 self.size = list() - def __iter__(self): + def __iter__(self): # noqa: D105 for i in self.size: yield i - def __len__(self): + def __len__(self): # noqa: D105 return len(self.size) def clear(self): + """Clear the size dimensions.""" self.size.clear() def get(self) -> int: + """Get the total size represented by this ArraySize.""" length = len(self.size) if length == 0: return 0 @@ -120,13 +125,13 @@ def get(self) -> int: return product def add(self, dim: int) -> None: - # Only allow adding valid dimensions + """Only allow adding valid dimensions.""" if dim >= 1: self.size.append(dim) def fortran_md_size(size: ArraySize) -> list: - """Format a multidimensional size into the correct Fortran string""" + """Format a multidimensional size into the correct Fortran string.""" md_size = list() for dim in size: md_size.append(f":{dim}") @@ -134,32 +139,32 @@ def fortran_md_size(size: ArraySize) -> list: def is_openacc(directive: Directive) -> bool: - """Check if a directive is OpenACC""" + """Check if a directive is OpenACC.""" return isinstance(directive, OpenACC) def is_openmp(directive: Directive) -> bool: - """Check if a directive is OpenMP""" + """Check if a directive is OpenMP.""" return isinstance(directive, OpenMP) def is_cxx(lang: Language) -> bool: - """Check if language is C++""" + """Check if language is C++.""" return isinstance(lang, Cxx) def is_fortran(lang: Language) -> bool: - """Check if language is Fortran""" + """Check if language is Fortran.""" return isinstance(lang, Fortran) def line_contains(line: str, target: str) -> bool: - """Generic helper to check if a line contains the target""" + """Generic helper to check if a line contains the target.""" return target in line def directive_contains_clause(line: str, clauses: list) -> bool: - """Check if a directive contains one clause from a list""" + """Check if a directive contains one clause from a list.""" for clause in clauses: if clause in line: return True @@ -167,7 +172,7 @@ def directive_contains_clause(line: str, clauses: list) -> bool: def line_contains_openacc_directive(line: str, lang: Language) -> bool: - """Check if line contains an OpenACC directive or not""" + """Check if line contains an OpenACC directive or not.""" if is_cxx(lang): return line_contains_openacc_directive_cxx(line) elif is_fortran(lang): @@ -176,17 +181,17 @@ def line_contains_openacc_directive(line: str, lang: Language) -> bool: def line_contains_openacc_directive_cxx(line: str) -> bool: - """Check if a line of code contains a C++ OpenACC directive or not""" + """Check if a line of code contains a C++ OpenACC directive or not.""" return line_contains(line, "#pragma acc") def line_contains_openacc_directive_fortran(line: str) -> bool: - """Check if a line of code contains a Fortran OpenACC directive or not""" + """Check if a line of code contains a Fortran OpenACC directive or not.""" return line_contains(line, "!$acc") def line_contains_openmp_directive(line: str, lang: Language) -> bool: - """Check if line contains an OpenMP directive or not""" + """Check if line contains an OpenMP directive or not.""" if is_cxx(lang): return line_contains_openmp_directive_cxx(line) elif is_fortran(lang): @@ -195,17 +200,17 @@ def line_contains_openmp_directive(line: str, lang: Language) -> bool: def line_contains_openmp_directive_cxx(line: str) -> bool: - """Check if a line of code contains a C++ OpenMP directive or not""" + """Check if a line of code contains a C++ OpenMP directive or not.""" return line_contains(line, "#pragma omp") def line_contains_openmp_directive_fortran(line: str) -> bool: - """Check if a line of code contains a Fortran OpenMP directive or not""" + """Check if a line of code contains a Fortran OpenMP directive or not.""" return line_contains(line, "!$omp") def line_contains_openacc_parallel_directive(line: str, lang: Language) -> bool: - """Check if line contains an OpenACC parallel directive or not""" + """Check if line contains an OpenACC parallel directive or not.""" if is_cxx(lang): return line_contains_openacc_parallel_directive_cxx(line) elif is_fortran(lang): @@ -214,17 +219,17 @@ def line_contains_openacc_parallel_directive(line: str, lang: Language) -> bool: def line_contains_openacc_parallel_directive_cxx(line: str) -> bool: - """Check if a line of code contains a C++ OpenACC parallel directive or not""" + """Check if a line of code contains a C++ OpenACC parallel directive or not.""" return line_contains(line, "#pragma acc parallel") def line_contains_openacc_parallel_directive_fortran(line: str) -> bool: - """Check if a line of code contains a Fortran OpenACC parallel directive or not""" + """Check if a line of code contains a Fortran OpenACC parallel directive or not.""" return line_contains(line, "!$acc parallel") def line_contains_openmp_target_directive(line: str, lang: Language) -> bool: - """Check if line contains an OpenMP target directive or not""" + """Check if line contains an OpenMP target directive or not.""" if is_cxx(lang): return line_contains_openmp_target_directive_cxx(line) elif is_fortran(lang): @@ -233,29 +238,29 @@ def line_contains_openmp_target_directive(line: str, lang: Language) -> bool: def line_contains_openmp_target_directive_cxx(line: str) -> bool: - """Check if a line of code contains a C++ OpenMP target directive or not""" + """Check if a line of code contains a C++ OpenMP target directive or not.""" return line_contains(line, "#pragma omp target") def line_contains_openmp_target_directive_fortran(line: str) -> bool: - """Check if a line of code contains a Fortran OpenMP target directive or not""" + """Check if a line of code contains a Fortran OpenMP target directive or not.""" return line_contains(line, "!$omp target") def openacc_directive_contains_data_clause(line: str) -> bool: - """Check if an OpenACC directive contains one data clause""" + """Check if an OpenACC directive contains one data clause.""" data_clauses = ["copy", "copyin", "copyout", "create", "no_create", "present", "device_ptr", "attach"] return directive_contains_clause(line, data_clauses) def openmp_directive_contains_data_clause(line: str) -> bool: - """Check if an OpenMP directive contains one data clause""" + """Check if an OpenMP directive contains one data clause.""" data_clauses = ["map"] return directive_contains_clause(line, data_clauses) def create_data_directive_openacc(name: str, size: ArraySize, lang: Language) -> str: - """Create a data directive for a given language""" + """Create a data directive for a given language.""" if is_cxx(lang): return create_data_directive_openacc_cxx(name, size) elif is_fortran(lang): @@ -264,12 +269,12 @@ def create_data_directive_openacc(name: str, size: ArraySize, lang: Language) -> def create_data_directive_openacc_cxx(name: str, size: ArraySize) -> str: - """Create C++ OpenACC code to allocate and copy data""" + """Create C++ OpenACC code to allocate and copy data.""" return f"#pragma acc enter data create({name}[:{size.get()}])\n#pragma acc update device({name}[:{size.get()}])\n" def create_data_directive_openacc_fortran(name: str, size: ArraySize) -> str: - """Create Fortran OpenACC code to allocate and copy data""" + """Create Fortran OpenACC code to allocate and copy data.""" if len(size) == 1: return f"!$acc enter data create({name}(:{size.get()}))\n!$acc update device({name}(:{size.get()}))\n" else: @@ -280,7 +285,7 @@ def create_data_directive_openacc_fortran(name: str, size: ArraySize) -> str: def create_data_directive_openmp(name: str, size: ArraySize, lang: Language) -> str: - """Create a data directive for a given language""" + """Create a data directive for a given language.""" if is_cxx(lang): return create_data_directive_openmp_cxx(name, size) elif is_fortran(lang): @@ -289,12 +294,12 @@ def create_data_directive_openmp(name: str, size: ArraySize, lang: Language) -> def create_data_directive_openmp_cxx(name: str, size: ArraySize) -> str: - """Create C++ OpenMP code to allocate and copy data""" + """Create C++ OpenMP code to allocate and copy data.""" return f"#pragma omp target enter data map(to: {name}[:{size.get()}])\n" def create_data_directive_openmp_fortran(name: str, size: ArraySize) -> str: - """Create Fortran OpenMP code to allocate and copy data""" + """Create Fortran OpenMP code to allocate and copy data.""" if len(size) == 1: return f"!$omp target enter data map(to: {name}(:{size.get()}))\n" else: @@ -303,7 +308,7 @@ def create_data_directive_openmp_fortran(name: str, size: ArraySize) -> str: def exit_data_directive_openacc(name: str, size: ArraySize, lang: Language) -> str: - """Create code to copy data back for a given language""" + """Create code to copy data back for a given language.""" if is_cxx(lang): return exit_data_directive_openacc_cxx(name, size) elif is_fortran(lang): @@ -312,12 +317,12 @@ def exit_data_directive_openacc(name: str, size: ArraySize, lang: Language) -> s def exit_data_directive_openacc_cxx(name: str, size: ArraySize) -> str: - """Create C++ OpenACC code to copy back data""" + """Create C++ OpenACC code to copy back data.""" return f"#pragma acc exit data copyout({name}[:{size.get()}])\n" def exit_data_directive_openacc_fortran(name: str, size: ArraySize) -> str: - """Create Fortran OpenACC code to copy back data""" + """Create Fortran OpenACC code to copy back data.""" if len(size) == 1: return f"!$acc exit data copyout({name}(:{size.get()}))\n" else: @@ -326,7 +331,7 @@ def exit_data_directive_openacc_fortran(name: str, size: ArraySize) -> str: def exit_data_directive_openmp(name: str, size: ArraySize, lang: Language) -> str: - """Create code to copy data back for a given language""" + """Create code to copy data back for a given language.""" if is_cxx(lang): return exit_data_directive_openmp_cxx(name, size) elif is_fortran(lang): @@ -335,12 +340,12 @@ def exit_data_directive_openmp(name: str, size: ArraySize, lang: Language) -> st def exit_data_directive_openmp_cxx(name: str, size: ArraySize) -> str: - """Create C++ OpenMP code to copy back data""" + """Create C++ OpenMP code to copy back data.""" return f"#pragma omp target exit data map(from: {name}[:{size.get()}])\n" def exit_data_directive_openmp_fortran(name: str, size: ArraySize) -> str: - """Create Fortran OpenMP code to copy back data""" + """Create Fortran OpenMP code to copy back data.""" if len(size) == 1: return f"!$omp target exit data map(from: {name}(:{size.get()}))\n" else: @@ -349,12 +354,12 @@ def exit_data_directive_openmp_fortran(name: str, size: ArraySize) -> str: def correct_kernel(kernel_name: str, line: str) -> bool: - """Checks if the line contains the correct kernel name""" + """Checks if the line contains the correct kernel name.""" return f" {kernel_name} " in line or (kernel_name in line and len(line.partition(kernel_name)[2]) == 0) def find_size_in_preprocessor(dimension: str, preprocessor: list) -> int: - """Find the dimension of a directive defined value in the preprocessor""" + """Find the dimension of a directive defined value in the preprocessor.""" ret_size = 0 for line in preprocessor: if f"#define {dimension}" in line: @@ -367,7 +372,7 @@ def find_size_in_preprocessor(dimension: str, preprocessor: list) -> int: def extract_code(start: str, stop: str, code: str, langs: Code, kernel_name: str = None) -> dict: - """Extract an arbitrary section of code""" + """Extract an arbitrary section of code.""" found_section = False sections = dict() tmp_string = list() @@ -400,7 +405,7 @@ def extract_code(start: str, stop: str, code: str, langs: Code, kernel_name: str def parse_size(size: Any, preprocessor: list = None, dimensions: dict = None) -> ArraySize: - """Converts an arbitrary object into an integer representing memory size""" + """Converts an arbitrary object into an integer representing memory size.""" ret_size = ArraySize() if type(size) is not int: try: @@ -441,7 +446,7 @@ def parse_size(size: Any, preprocessor: list = None, dimensions: dict = None) -> def wrap_timing(code: str, lang: Language) -> str: - """Helper to wrap timing code around the provided code""" + """Helper to wrap timing code around the provided code.""" if is_cxx(lang): return end_timing_cxx(start_timing_cxx(code)) elif is_fortran(lang): @@ -450,8 +455,7 @@ def wrap_timing(code: str, lang: Language) -> str: def start_timing_cxx(code: str) -> str: - """Wrap C++ timing code around the provided code""" - + """Wrap C++ timing code around the provided code.""" start = "auto kt_timing_start = std::chrono::steady_clock::now();" end = "auto kt_timing_end = std::chrono::steady_clock::now();" timing = "std::chrono::duration elapsed_time = kt_timing_end - kt_timing_start;" @@ -460,8 +464,7 @@ def start_timing_cxx(code: str) -> str: def wrap_timing_fortran(code: str) -> str: - """Wrap Fortran timing code around the provided code""" - + """Wrap Fortran timing code around the provided code.""" start = "call system_clock(kt_timing_start, kt_rate)" end = "call system_clock(kt_timing_end)" timing = "timing = (real(kt_timing_end - kt_timing_start) / real(kt_rate)) * 1e3" @@ -470,12 +473,12 @@ def wrap_timing_fortran(code: str) -> str: def end_timing_cxx(code: str) -> str: - """In C++ we need to return the measured time""" + """In C++ we need to return the measured time.""" return "\n".join([code, "return elapsed_time.count();\n"]) def wrap_data(code: str, langs: Code, data: dict, preprocessor: list = None, user_dimensions: dict = None) -> str: - """Insert data directives before and after the timed code""" + """Insert data directives before and after the timed code.""" intro = str() outro = str() for name in data.keys(): @@ -492,7 +495,7 @@ def wrap_data(code: str, langs: Code, data: dict, preprocessor: list = None, use def wrap_data_openacc(name: str, size: int, langs: Code) -> Tuple[str, str]: - """Create language specific data directives""" + """Create language specific data directives.""" if is_cxx(langs.language): intro = create_data_directive_openacc_cxx(name, size) outro = exit_data_directive_openacc_cxx(name, size) @@ -503,7 +506,7 @@ def wrap_data_openacc(name: str, size: int, langs: Code) -> Tuple[str, str]: def wrap_data_openmp(name: str, size: int, langs: Code) -> Tuple[str, str]: - """Create language specific data directives""" + """Create language specific data directives.""" if is_cxx(langs.language): intro = create_data_directive_openmp_cxx(name, size) outro = exit_data_directive_openmp_cxx(name, size) @@ -514,7 +517,7 @@ def wrap_data_openmp(name: str, size: int, langs: Code) -> Tuple[str, str]: def extract_directive_code(code: str, langs: Code, kernel_name: str = None) -> dict: - """Extract explicitly marked directive sections from code""" + """Extract explicitly marked directive sections from code.""" if is_cxx(langs.language): start_string = "#pragma tuner start" elif is_fortran(langs.language): @@ -524,7 +527,7 @@ def extract_directive_code(code: str, langs: Code, kernel_name: str = None) -> d def extract_initialization_code(code: str, langs: Code) -> str: - """Extract the initialization section from code""" + """Extract the initialization section from code.""" if is_cxx(langs.language): start_string = "#pragma tuner initialize" elif is_fortran(langs.language): @@ -538,7 +541,7 @@ def extract_initialization_code(code: str, langs: Code) -> str: def extract_deinitialization_code(code: str, langs: Code) -> str: - """Extract the deinitialization section from code""" + """Extract the deinitialization section from code.""" if is_cxx(langs.language): start_string = "#pragma tuner deinitialize" elif is_fortran(langs.language): @@ -552,7 +555,7 @@ def extract_deinitialization_code(code: str, langs: Code) -> str: def format_argument_fortran(p_type: str, p_size: int, p_name: str) -> str: - """Format the argument for Fortran code""" + """Format the argument for Fortran code.""" argument = "" if "float*" in p_type: argument = f"real (c_float), dimension({p_size}) :: {p_name}" @@ -570,8 +573,7 @@ def format_argument_fortran(p_type: str, p_size: int, p_name: str) -> str: def extract_directive_signature(code: str, langs: Code, kernel_name: str = None) -> dict: - """Extract the user defined signature for directive sections""" - + """Extract the user defined signature for directive sections.""" if is_cxx(langs.language): start_string = "#pragma tuner start" elif is_fortran(langs.language): @@ -605,9 +607,9 @@ def extract_directive_signature(code: str, langs: Code, kernel_name: str = None) if is_cxx(langs.language): signatures[name] = f"float {name}({', '.join(params)})" elif is_fortran(langs.language): - signatures[ - name - ] = f"function {name}({', '.join(params)}) result(timing)\nuse iso_c_binding\nimplicit none\n" + signatures[name] = ( + f"function {name}({', '.join(params)}) result(timing)\nuse iso_c_binding\nimplicit none\n" + ) params = list() for param in tmp_string: if len(param) == 0: @@ -619,16 +621,15 @@ def extract_directive_signature(code: str, langs: Code, kernel_name: str = None) p_type = p_type.split(":")[0] params.append(format_argument_fortran(p_type, p_size, p_name)) signatures[name] += "\n".join(params) + "\n" - signatures[ - name - ] += "integer(c_int):: kt_timing_start\nreal(c_float):: kt_rate\ninteger(c_int):: kt_timing_end\nreal(c_float):: timing\n" + signatures[name] += ( + "integer(c_int):: kt_timing_start\nreal(c_float):: kt_rate\ninteger(c_int):: kt_timing_end\nreal(c_float):: timing\n" # noqa: E501 + ) return signatures def extract_directive_data(code: str, langs: Code, kernel_name: str = None) -> dict: - """Extract the data used in the directive section""" - + """Extract the data used in the directive section.""" if is_cxx(langs.language): start_string = "#pragma tuner start" elif is_fortran(langs.language): @@ -662,7 +663,7 @@ def extract_directive_data(code: str, langs: Code, kernel_name: str = None) -> d def extract_preprocessor(code: str) -> list: - """Extract include and define statements from code""" + """Extract include and define statements from code.""" preprocessor = list() for line in code.replace("\\\n", "").split("\n"): @@ -682,8 +683,7 @@ def generate_directive_function( deinitialization: str = "", user_dimensions: dict = None, ) -> str: - """Generate tunable function for one directive""" - + """Generate tunable function for one directive.""" if is_cxx(langs.language): code = cpp_template body = start_timing_cxx(body) @@ -717,7 +717,7 @@ def generate_directive_function( def allocate_array(p_type: str, size: int) -> np.ndarray: - """Allocate a Numpy array""" + """Allocate a Numpy array.""" max_int = 1024 array = None if p_type == "float*": @@ -733,7 +733,7 @@ def allocate_array(p_type: str, size: int) -> np.ndarray: def allocate_scalar(p_type: str, size: int) -> np.number: - """Allocate a Numpy scalar""" + """Allocate a Numpy scalar.""" scalar = None if p_type == "float": scalar = np.float32(size) @@ -748,7 +748,7 @@ def allocate_scalar(p_type: str, size: int) -> np.number: def allocate_signature_memory(data: dict, preprocessor: list = None, user_dimensions: dict = None) -> list: - """Allocates the data needed by a kernel and returns the arguments array""" + """Allocates the data needed by a kernel and returns the arguments array.""" args = [] for parameter in data.keys(): @@ -763,7 +763,7 @@ def allocate_signature_memory(data: dict, preprocessor: list = None, user_dimens def add_new_line(line: str) -> str: - """Adds the new line character to the end of the line if not present""" + """Adds the new line character to the end of the line if not present.""" if line.rfind("\n") != len(line) - 1: return line + "\n" return line @@ -772,7 +772,7 @@ def add_new_line(line: str) -> str: def add_present_openacc( code: str, langs: Code, data: dict, preprocessor: list = None, user_dimensions: dict = None ) -> str: - """Add the present clause to OpenACC directive""" + """Add the present clause to OpenACC directive.""" new_body = "" for line in code.replace("\\\n", "").split("\n"): if not line_contains_openacc_parallel_directive(line, langs.language): @@ -798,12 +798,12 @@ def add_present_openacc( def add_present_openacc_cxx(name: str, size: ArraySize) -> str: - """Create present clause for C++ OpenACC directive""" + """Create present clause for C++ OpenACC directive.""" return f" present({name}[:{size.get()}]) " def add_present_openacc_fortran(name: str, size: ArraySize) -> str: - """Create present clause for Fortran OpenACC directive""" + """Create present clause for Fortran OpenACC directive.""" if len(size) == 1: return f" present({name}(:{size.get()})) " else: @@ -812,7 +812,7 @@ def add_present_openacc_fortran(name: str, size: ArraySize) -> str: def process_directives(langs: Code, source: str, user_dimensions: dict = None) -> Tuple[dict, dict]: - """Helper functions to process all the directives in the code and create tunable functions""" + """Helper functions to process all the directives in the code and create tunable functions.""" kernel_strings = dict() kernel_args = dict() preprocessor = extract_preprocessor(source) diff --git a/noxfile.py b/noxfile.py index 97fdd2f7d..f93da3e5a 100644 --- a/noxfile.py +++ b/noxfile.py @@ -5,27 +5,41 @@ Be careful that the general setup of tests is left to pyproject.toml. """ - import platform import re +import sys from pathlib import Path import nox from nox_poetry import Session, session +# Create a repository root path to access the Julia helper functions +# TODO this should be changed to something more robust +REPO_ROOT = Path(__file__).parent.joinpath(Path("kernel_tuner/backends")).resolve() +sys.path.append(str(REPO_ROOT)) +from julia_helper import backend_map, detect_julia_gpu_backends # noqa: E402, F401 + # set the test parameters verbose = False -python_versions_to_test = ["3.11", "3.12", "3.13", "3.14"] +python_versions_to_test = ["3.14", "3.13", "3.12", "3.11"] nox.options.stop_on_first_error = True nox.options.error_on_missing_interpreters = True -nox.options.default_venv_backend = 'virtualenv' +nox.options.default_venv_backend = "virtualenv" +nox.options.reuse_existing_virtualenvs = False # workspace level settings settings_file_path = Path("./noxsettings.toml") -venvbackend_values = ('none', 'virtualenv', 'conda', 'mamba', 'venv') # from https://nox.thea.codes/en/stable/usage.html#changing-the-sessions-default-backend +venvbackend_values = ( + "none", + "virtualenv", + "conda", + "mamba", + "venv", +) # from https://nox.thea.codes/en/stable/usage.html#changing-the-sessions-default-backend + # TODO remove this from a session function, session is only needed to receive trigger argument -@session # to only run on the current python interpreter +@session # to only run on the current python interpreter def create_settings(session: Session) -> None: """One-time creation of noxsettings.toml.""" arg_trigger = False @@ -44,47 +58,63 @@ def create_settings(session: Session) -> None: noxenv_file_path.unlink() # write the settings assert venvbackend in venvbackend_values, f"{venvbackend=}, must be one of {','.join(venvbackend_values)}" - settings = (f'venvbackend = "{venvbackend}"\n' - f'envdir = "{envdir}"\n') + settings = f'venvbackend = "{venvbackend}"\nenvdir = "{envdir}"\n' settings_file_path.write_text(settings) # exit to make sure the user checks the settings are correct if arg_trigger: - session.warn(f"Settings file '{settings_file_path}' created, exiting. Please check settings are correct before running Nox again.") + session.warn( + f"Settings file '{settings_file_path}' created, exiting. " + "Please check settings are correct before running Nox again." + ) exit(1) + # obtain workspace level settings from the 'noxsettings.toml' file if settings_file_path.exists(): with settings_file_path.open(mode="rb") as fp: import tomli + nox_settings = tomli.load(fp) - venvbackend = nox_settings['venvbackend'] - envdir = nox_settings['envdir'] - assert venvbackend in venvbackend_values, f"File '{settings_file_path}' has {venvbackend=}, must be one of {','.join(venvbackend_values)}" + venvbackend = nox_settings["venvbackend"] + envdir = nox_settings["envdir"] + assert venvbackend in venvbackend_values, ( + f"File '{settings_file_path}' has {venvbackend=}, must be one of {','.join(venvbackend_values)}" + ) nox.options.default_venv_backend = venvbackend nox.options.venvbackend = venvbackend if envdir is not None and len(envdir) > 0: nox.options.envdir = envdir -# @session # to only run on the current python interpreter + +# @session # to only run on the current python interpreter # def lint(session: Session) -> None: # """Ensure the code is formatted as expected.""" # session.install("ruff") -# session.run("ruff", "--output-format=github", "--config=pyproject.toml", ".") +# session.warn("Linting errors detected:") +# session.run(*"ruff check . --config=pyproject.toml --statistics --exit-zero".split()) +# session.run(*"ruff check --config=pyproject.toml --output-format=github .".split()) + -@session # to only run on the current python interpreter +@session # to only run on the current python interpreter def check_poetry(session: Session) -> None: """Check whether Poetry is correctly configured.""" session.run("poetry", "check", "--no-interaction", external=True) -@session # to only run on the current python interpreter + +@session # to only run on the current python interpreter def check_development_environment(session: Session) -> None: """Check whether the development environment is up to date with the dependencies, and try to update if necessary.""" if session.posargs: - if 'github-action' in session.posargs: - session.log("Skipping development environment check on the GitHub Actions runner, as this is always up to date.") + if "github-action" in session.posargs: + session.log( + "Skipping development environment check on the GitHub Actions runner, as this is always up to date." + ) return None output: str = session.run("poetry", "install", "--sync", "--dry-run", "--with", "test", silent=True, external=True) - match = re.search(r"Package operations: (\d+) (?:install|installs), (\d+) (?:update|updates), (\d+) (?:removal|removals), \d+ skipped", output) + match = re.search( + r"Package operations: (\d+) (?:install|installs), (\d+) (?:update|updates), (\d+) (?:removal|removals), \d+ skipped", # noqa: E501 + output, + ) assert match is not None, f"Could not check development environment, reason: {output}" groups = match.groups() installs, updates, removals = int(groups[0]), int(groups[1]), int(groups[2]) @@ -95,54 +125,85 @@ def check_development_environment(session: Session) -> None: Your development environment is out of date ({installs} installs, {updates} updates). Update with 'poetry install --sync', using '--with' and '-E' for optional dependencies, extras respectively. Note: {removals} packages are not in the specification (i.e. installed manually) and may be removed. - To preview changes, run 'poetry install --sync --dry-run' (with optional dependencies and extras).""") + To preview changes, run 'poetry install --sync --dry-run' (with optional dependencies and extras).""" + ) + @session(python=python_versions_to_test) # missing versions can be installed with `pyenv install ...` # do not forget check / set the versions with `pyenv global`, or `pyenv local` in case of virtual environment def tests(session: Session) -> None: """Run the tests for the specified Python versions.""" session.log(f"Testing on Python {session.python}") + env_vars = {} # check if optional dependencies have been disabled by user arguments (e.g. `nox -- skip-gpu`, `nox -- skip-cuda`) install_cuda = True install_hip = True install_opencl = True + install_julia = True + julia_use_gpu = True install_additional_tests = False small_disk = False + skip_gpu = False + github_action = False if session.posargs: for arg in session.posargs: if arg.lower() == "skip-gpu": install_cuda = False install_hip = False install_opencl = False - break + julia_use_gpu = False + skip_gpu = True elif arg.lower() == "skip-cuda": install_cuda = False elif arg.lower() == "skip-hip": install_hip = False elif arg.lower() == "skip-opencl": install_opencl = False + elif arg.lower() == "skip-julia": + install_julia = False + julia_use_gpu = False + elif arg.lower() == "julia-use-gpu": + julia_use_gpu = True elif arg.lower() == "additional-tests": install_additional_tests = True elif arg.lower() == "small-disk": small_disk = True + elif arg.lower() == "github-action": + github_action = True else: raise ValueError(f"Unrecognized argument {arg}") + if install_julia == False and julia_use_gpu == True: + raise ValueError("Cannot use Julia GPU backend if Julia is disabled") + if skip_gpu == True and julia_use_gpu == True: + print("With both `skip-gpu` and `julia-use-gpu` given, the latter takes precedence. Julia GPU backends will be used if available.") # check if there are optional dependencies that can not be installed if install_hip: - if platform.system().lower() != 'linux': + if platform.system().lower() != "linux": session.warn("HIP is only available on Linux, disabling dependency and tests") install_hip = False - full_install = install_cuda and install_hip and install_opencl and install_additional_tests + full_install = install_cuda and install_hip and install_opencl and install_julia and install_additional_tests # if the user has a small disk, remove the other environment caches before each session is ran if small_disk: try: - session_folder = session.name.replace('.', '*').strip() + session_folder = session.name.replace(".", "*").strip() folders_to_delete: str = session.run( - "find", "./.nox", "-mindepth", "1", "-maxdepth", "1", "-type", "d", "-not", "-name", session_folder, - silent=True, external=True) - folders_to_delete: list[str] = folders_to_delete.split('\n') + "find", + "./.nox", + "-mindepth", + "1", + "-maxdepth", + "1", + "-type", + "d", + "-not", + "-name", + session_folder, + silent=True, + external=True, + ) + folders_to_delete: list[str] = folders_to_delete.split("\n") for folder_to_delete in folders_to_delete: if len(folder_to_delete) > 0: session.warn(f"Removing environment cache {folder_to_delete} because of 'small-disk' argument") @@ -153,6 +214,9 @@ def tests(session: Session) -> None: # remove temporary files leftover from the previous session session.run("rm", "-f", "temp_*.c", external=True) + session.run("rm", "-f", "temp_*.cu", external=True) + session.run("rm", "-f", "temp_*.cl", external=True) + session.run("rm", "-f", "temp_*.jl", external=True) # set extra arguments based on optional dependencies extras_args = [] @@ -162,60 +226,107 @@ def tests(session: Session) -> None: extras_args.extend(["-E", "hip"]) if install_opencl: extras_args.extend(["-E", "opencl"]) + if install_julia: + extras_args.extend(["-E", "julia"]) + # set the paths to Julia install, environment and project + session_envdir = nox.options.envdir if nox.options.envdir is not None else session.env["VIRTUAL_ENV"] + julia_envdir = Path(session_envdir) / ".julia" + if julia_envdir is not None: + session.env["JULIA_DEPOT_PATH"] = str(Path(julia_envdir).resolve()) + env_vars["JULIA_DEPOT_PATH"] = str(Path(julia_envdir).resolve()) + # session.env["JULIAUP_DEPOT_PATH"] = session.env["JULIA_DEPOT_PATH"] + # session.env["PYTHON_JULIACALL_PROJECT"] = str(Path(julia_envdir).resolve()) # separately install optional dependencies with weird dependencies / build process install_warning = """Installation failed, this likely means that the required hardware or drivers are missing. - Run with `-- skip-gpu` or one of the more specific options (e.g. `-- skip-cuda`) to avoid this.""" + Run with `-- skip-gpu` / `-- skip-julia` or one of the more specific options (e.g. `-- skip-cuda`) to avoid this.""" # noqa: E501 if install_cuda: # use NVCC to get the CUDA version import re + nvcc_output: str = session.run("nvcc", "--version", silent=True, external=True) nvcc_output = "".join(nvcc_output.splitlines()) # convert to single string for easier REGEX cuda_version = re.match(r"^.*release ([0-9]+.[0-9]+).*$", nvcc_output, flags=re.IGNORECASE).group(1).strip() session.warn(f"Detected CUDA version: {cuda_version}") - # if we need to install the CUDA extras, first install pycuda seperately, reason: - # since version 2022.2 it has `oldest-supported-numpy` as a build dependency which doesn't work with Poetry - if " not found: " in session.run("pip", "show", "pycuda", external=True, silent=True, success_codes=[0,1]): - # if PyCUDA is not installed, install it - session.warn("PyCUDA not installed") - try: - session.install("pycuda", "--no-cache-dir", "--force-reinstall") # Attention: if changed, check `pycuda` in pyproject.toml as well - except Exception as error: - session.log(error) - session.warn(install_warning) + if " not found: " in session.run("pip", "show", "pycuda", external=True, silent=True, success_codes=[0, 1]): + # if we need to install the CUDA extras, first install pycuda seperately, reason: + # between version 2022.2 and 2025.1 PyCUDA had `oldest-supported-numpy` as a build dependency which doesn't work with Poetry + # No longer needed as of 2025.1 + # commented as no longer needed + # # if PyCUDA is not installed, install it + # session.warn("PyCUDA not installed") + # try: + # session.install( + # "pycuda", "--no-cache-dir", "--force-reinstall" + # ) # Attention: if changed, check `pycuda` in pyproject.toml as well + # except Exception as error: + # session.log(error) + # session.warn(install_warning) + pass else: session.warn("PyCUDA installed") - # if PyCUDA is already installed, check whether the CUDA version PyCUDA was installed with matches the current CUDA version - session.install("numpy") # required by pycuda.driver - pycuda_version = session.run("python", "-c", "import pycuda.driver as drv; drv.init(); print('.'.join(list(str(d) for d in drv.get_version())))", silent=True) - shortest_string, longest_string = (pycuda_version, cuda_version) if len(pycuda_version) < len(cuda_version) else (cuda_version, pycuda_version) - if longest_string[:len(shortest_string)] != shortest_string: - session.warn(f"PyCUDA was compiled with a version of CUDA ({pycuda_version}) that does not match the current version ({cuda_version}). Re-installing.") + # if PyCUDA is already installed, check whether the CUDA version PyCUDA was installed with matches the current CUDA version # noqa: E501 + session.install("numpy") # required by pycuda.driver + pycuda_version = session.run( + "python", + "-c", + "import pycuda.driver as drv; drv.init(); print('.'.join(list(str(d) for d in drv.get_version())))", + silent=True, + ) + shortest_string, longest_string = ( + (pycuda_version, cuda_version) + if len(pycuda_version) < len(cuda_version) + else (cuda_version, pycuda_version) + ) + if longest_string[: len(shortest_string)] != shortest_string: + session.warn( + f"PyCUDA was compiled with a version of CUDA ({pycuda_version}) that does not match the current version ({cuda_version}). Re-installing." # noqa: E501 + ) try: - session.install("pycuda", "--no-cache-dir", "--force-reinstall") # Attention: if changed, check `pycuda` in pyproject.toml as well + session.install( + "pycuda>=2025.1", "--no-cache-dir", "--force-reinstall" + ) # Attention: if changed, check `pycuda` in pyproject.toml as well except Exception as error: session.log(error) session.warn(install_warning) # finally, install the dependencies, optional dependencies and the package itself - poetry_env = Path(session.run_always("poetry", "env", "info", "--executable", silent=not verbose, external=True).splitlines()[-1].strip()).resolve() + poetry_env = Path( + session.run_always("poetry", "env", "info", "--executable", silent=not verbose, external=True) + .splitlines()[-1] + .strip() + ).resolve() session_env = Path(session.bin, "python/").resolve() assert poetry_env.exists(), f"{poetry_env=} does not exist" assert session_env.exists(), f"{session_env=} does not exist" # if the poetry virtualenv is not set to the session env, use requirements file export instead of Poetry install if poetry_env != session_env: - session.warn(f"Poetry env ({str(poetry_env)}) is not session env ({str(session_env)}), falling back to install via requirements export") + session.warn( + f"Poetry env ({str(poetry_env)}) is not session env ({str(session_env)}), falling back to install via requirements export" # noqa: E501 + ) requirements_file = Path(f"tmp_test_requirements_{session.name}.txt") if requirements_file.exists(): requirements_file.unlink() if verbose: - print(session.run_always('conda', 'list')) - session.run_always('poetry', 'export', '-f', 'requirements.txt', '-o', requirements_file.name, '--with=test', '--without-hashes', *extras_args, external=True, silent=not verbose) - session.install('-r', requirements_file.name) - session.install('.') + print(session.run_always("conda", "list")) + session.run_always( + "poetry", + "export", + "-f", + "requirements.txt", + "-o", + requirements_file.name, + "--with=test", + "--without-hashes", + *extras_args, + external=True, + silent=not verbose, + ) + session.install("-r", requirements_file.name) + session.install(".") requirements_file.unlink() if verbose: - print(session.run_always('conda', 'list')) + print(session.run_always("conda", "list")) else: try: session.run_always("poetry", "install", "--with", "test", *extras_args, external=True, silent=False) @@ -223,6 +334,48 @@ def tests(session: Session) -> None: session.warn(install_warning) raise error + # if applicable, install julia dependencies in the session environment + if install_julia: + # sanitize loader env to avoid issues with Julia loading libraries from the wrong environment + for v in ["DYLD_LIBRARY_PATH", "DYLD_FALLBACK_LIBRARY_PATH"]: + session.env.pop(v, None) + # set the Julia version + # when changed, also see `require_julia` in Project.toml and the Julia version in the GitHub Actions workflow + if github_action: + preamble = "import juliapkg; juliapkg.require_julia('1.11')" # must match the setup-julia action + else: + preamble = "import juliapkg; juliapkg.require_julia('1.11, 2')" # meaning >= 1.11, < 2.0 + # call JuliaPKG to precompile packages in the session environment + session.run( + "python", "-c", + f"{preamble}; juliapkg.resolve(update=True)", + ) + # retrieve the project path for this isolated session environment and pass it as an environment variable + julia_project_path = session.run( + "python", "-c", + f"{preamble}; print(juliapkg.project())", + silent=True + ).strip() + # create the .julia/registries directory if it doesn't exist to avoid juliapkg.add() crash + session.run( + "bash", "-c", + "[ -d ~ ] && mkdir -p ~/.julia/registries", + external=True + ) + session.env["PYTHON_JULIAPKG_PROJECT"] = julia_project_path + env_vars["PYTHON_JULIAPKG_PROJECT"] = julia_project_path + # install any additional dependencies used by the tests, as `check_package_and_install` won't work from Nox + if julia_use_gpu: + gpu_backends_string = "".join( + f'juliapkg.add("{backend_map[backend]["pkg"]}"); ' if backend_map[backend]["pkg"] else "" for backend in detect_julia_gpu_backends() + ) + else: + gpu_backends_string = "" + session.run( + "python", "-c", + f"{preamble}; juliapkg.resolve(); juliapkg.add('KernelAbstractions'); {gpu_backends_string} juliapkg.resolve();", + ) + # if applicable, install the dependencies for additional tests if install_additional_tests and install_cuda: install_additional_warning = """ @@ -253,16 +406,18 @@ def tests(session: Session) -> None: # for the last Python version session if all optional dependencies are enabled: if session.python == python_versions_to_test[-1] and full_install: # run pytest on the package to generate the correct coverage report - session.run("pytest", external=False) + session.run("pytest", external=False, env=env_vars) else: # for the other Python version sessions: # run pytest without coverage reporting - session.run("pytest", "--no-cov", external=False) + session.run("pytest", "--no-cov", external=False, env=env_vars) # warn if no coverage report if not full_install: - session.warn(""" + session.warn( + """ Tests ran successfully, but only a subset. Coverage file not generated. Run with 'additional-tests' and without 'skip-gpu', 'skip-cuda' etc. to avoid this. - """) + """ + ) diff --git a/pyproject.toml b/pyproject.toml index 479ec8a05..ec9542ba3 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -5,7 +5,7 @@ build-backend = "poetry.core.masonry.api" [project] name = "kernel_tuner" description = "An easy to use CUDA/OpenCL kernel tuner in Python" -version = "1.4.0" # adhere to PEP440 versioning: https://packaging.python.org/en/latest/guides/distributing-packages-using-setuptools/#id55 +version = "1.5.0" # adhere to PEP440 versioning: https://packaging.python.org/en/latest/guides/distributing-packages-using-setuptools/#id55 readme = "README.md" license = "Apache-2.0" authors = [ @@ -54,7 +54,7 @@ dependencies = [ "scipy>=1.14.1", # Python >=3.13 needs scipy >=1.14 "packaging", # required by file_utils "jsonschema", - "python-constraint2>=2.4.0", + "python-constraint2>=2.5.0", "xmltodict", "pandas>=2.0.0", "scikit-learn>=1.0.2", @@ -99,6 +99,7 @@ cuda = ["cuda-python>=12.6.0", "nvidia-ml-py>=12.535.108", "pynvml>=11.4.1"] # A opencl = ["pyopencl"] # Attention: if pyopencl is changed here, also change `session.install("pyopencl")` in the Noxfile cuda_opencl = ["cuda-python>=12.6.0", "pyopencl"] # Attention: if pycuda is changed here, also change `session.install("pycuda")` in the Noxfile hip = ["hip-python"] +julia = ["juliapkg>=0.1.23", "juliacall>=0.9.31"] tutorial = ["jupyter>=1.0.0", "matplotlib>=3.5.0", "nvidia-ml-py>=12.535.108"] # ATTENTION: if anything is changed here, run `poetry update` and `poetry export --with docs --without-hashes --format=requirements.txt --output doc/requirements.txt` @@ -126,7 +127,7 @@ pytest-cov = "^5.0.0" mock = "^5.1.0" nox = "^2024.4.15" nox-poetry = "^1.0.3" -ruff = "^0.4.8" +ruff = "^0.15.0" pep440 = "^0.1.2" tomli = "^2.0.1" # held back by Python <= 3.10, can be replaced by built-in [tomllib](https://docs.python.org/3.11/library/tomllib.html) from Python 3.11 onwards scikit-optimize = "0.10.2" @@ -148,13 +149,13 @@ line-length = 120 [tool.ruff] line-length = 120 respect-gitignore = true -exclude = ["doc", "examples"] +exclude = ["doc", "examples", "test"] +[tool.ruff.lint] select = [ + "NPY201", # numpy docstring style "E", # pycodestyle "F", # pyflakes, "D", # pydocstyle, ] -[tool.ruff.pydocstyle] -convention = "google" -[tool.ruff.lint] -select = ["NPY201"] +[tool.ruff.lint.pydocstyle] +convention = "google" \ No newline at end of file diff --git a/test/conftest.py b/test/conftest.py index 1539a6cdf..a1867210e 100644 --- a/test/conftest.py +++ b/test/conftest.py @@ -4,4 +4,4 @@ def pytest_collection_modifyitems(items): for item in items: if item.get_closest_marker('timeout') is None: - item.add_marker(pytest.mark.timeout(60)) \ No newline at end of file + item.add_marker(pytest.mark.timeout(180)) \ No newline at end of file diff --git a/test/context.py b/test/context.py index 9ed2efa25..c49a762f3 100644 --- a/test/context.py +++ b/test/context.py @@ -2,6 +2,7 @@ import subprocess import sys import ctypes.util +from os import environ import pytest @@ -35,19 +36,19 @@ gfortran_present = shutil.which("gfortran") is not None openmp_present = ctypes.util.find_library('gomp') is not None openacc_present = shutil.which("nvc++") is not None +running_on_ci = any([environ.get(CI, "false").lower() == "true" for CI in ["GITHUB_ACTIONS", "TRAVIS", "CIRCLECI", "GITLAB_CI"]]) try: import cupy - cupy.cuda.Device( - 0 - ).attributes # triggers exception if there are no CUDA-capable devices + cupy.cuda.Device(0).attributes # triggers exception if there are no CUDA-capable devices cupy_present = True except Exception: cupy_present = False try: import cuda + print(cuda) cuda_present = True except Exception: @@ -55,6 +56,7 @@ try: from hip import hip + hip.hipDriverGetVersion() hip_present = True except (ImportError, RuntimeError): @@ -63,6 +65,7 @@ try: import botorch import torch + bayes_opt_botorch_present = True except ImportError: bayes_opt_botorch_present = False @@ -70,12 +73,14 @@ try: import gpytorch import torch + bayes_opt_gpytorch_present = True except ImportError: bayes_opt_gpytorch_present = False try: import pyatf + pyatf_present = True except ImportError: pyatf_present = False @@ -86,33 +91,35 @@ except ImportError: pymoo_present = False +try: + import juliacall + julia_present = True +except ImportError: + julia_present = False + try: from autotuning_methodology.report_experiments import get_strategy_scores + methodology_present = True except ImportError: methodology_present = False -skip_if_no_pycuda = pytest.mark.skipif( - not pycuda_present, reason="PyCuda not installed or no CUDA device detected" -) +skip_if_no_pycuda = pytest.mark.skipif(not pycuda_present, reason="PyCuda not installed or no CUDA device detected") skip_if_no_pynvml = pytest.mark.skipif(not pynvml_present, reason="NVML not installed") -skip_if_no_cupy = pytest.mark.skipif( - not cupy_present, reason="CuPy not installed or no CUDA device detected" -) -skip_if_no_cuda = pytest.mark.skipif( - not cuda_present, reason="NVIDIA CUDA not installed" -) -skip_if_no_opencl = pytest.mark.skipif( - not opencl_present, reason="PyOpenCL not installed or no OpenCL device detected" -) +skip_if_no_cupy = pytest.mark.skipif(not cupy_present, reason="CuPy not installed or no CUDA device detected") +skip_if_no_cuda = pytest.mark.skipif(not cuda_present, reason="NVIDIA CUDA not installed") +skip_if_no_opencl = pytest.mark.skipif(not opencl_present, reason="PyOpenCL not installed or no OpenCL device detected") skip_if_no_gcc = pytest.mark.skipif(not gcc_present, reason="No gcc on PATH") -skip_if_no_gfortran = pytest.mark.skipif( - not gfortran_present, reason="No gfortran on PATH" -) +skip_if_no_gfortran = pytest.mark.skipif(not gfortran_present, reason="No gfortran on PATH") +skip_if_no_julia = pytest.mark.skipif(not shutil.which("julia") or not julia_present, reason="No Julia on PATH or juliacall not installed") skip_if_no_openmp = pytest.mark.skipif(not openmp_present, reason="No OpenMP found") skip_if_no_openacc = pytest.mark.skipif(not openacc_present, reason="No nvc++ on PATH") -skip_if_no_bayesopt_gpytorch = pytest.mark.skipif(not bayes_opt_gpytorch_present, reason="Torch and GPyTorch not installed") -skip_if_no_bayesopt_botorch = pytest.mark.skipif(not bayes_opt_botorch_present, reason="Torch and BOTorch not installed") +skip_if_no_bayesopt_gpytorch = pytest.mark.skipif( + not bayes_opt_gpytorch_present, reason="Torch and GPyTorch not installed" +) +skip_if_no_bayesopt_botorch = pytest.mark.skipif( + not bayes_opt_botorch_present, reason="Torch and BOTorch not installed" +) skip_if_no_hip = pytest.mark.skipif(not hip_present, reason="No HIP Python found") skip_if_no_pyatf = pytest.mark.skipif(not pyatf_present, reason="PyATF not installed") skip_if_no_methodology = pytest.mark.skipif(not methodology_present, reason="Autotuning Methodology not found") @@ -136,3 +143,5 @@ def skip_backend(backend: str): pytest.skip("No nvc++ on PATH") elif backend.upper() == "HIP" and not hip_present: pytest.skip("HIP Python not installed") + elif backend.upper() == "JULIA" and not shutil.which("julia"): + pytest.skip("No Julia on PATH") diff --git a/test/test_backend.py b/test/test_backend.py index e694649c1..d56bdcce0 100644 --- a/test/test_backend.py +++ b/test/test_backend.py @@ -4,8 +4,9 @@ skip_if_no_cuda, skip_if_no_opencl, skip_if_no_pycuda, + skip_if_no_julia, ) -from kernel_tuner.backends import backend, compiler, cupy, nvcuda, opencl, pycuda +from kernel_tuner.backends import backend, compiler, cupy, nvcuda, opencl, pycuda, julia class WrongBackend(backend.Backend): @@ -45,3 +46,8 @@ def test_opencl_backend(): @skip_if_no_pycuda def test_pycuda_backend(): dev = pycuda.PyCudaFunctions() + + +@skip_if_no_julia +def test_julia_backend(): + dev = julia.JuliaFunctions() diff --git a/test/test_core.py b/test/test_core.py index 67c8b7c6a..3a807b31e 100644 --- a/test/test_core.py +++ b/test/test_core.py @@ -96,9 +96,6 @@ def __getattr__(self, name): dev.check_kernel_output(func, gpu_args, instance, answer, 1e-6, None, True) - - - def test_default_verify_function_arrays(): answer = [np.zeros(4).astype(np.float32), None, np.ones(5).astype(np.int32)] @@ -117,7 +114,7 @@ def test_default_verify_function_arrays(): core._default_verify_function(instance, ans, result_host, 0, False) print("_default_verify_function failed to throw an exception") assert False - except TypeError: + except (TypeError, ValueError): assert True for result_host in [result_host, result_host2]: @@ -140,7 +137,7 @@ def test_default_verify_function_scalar(): core._default_verify_function(instance, ans, result_host, 0, False) print("_default_verify_function failed to throw an exception") assert False - except TypeError: + except (TypeError, ValueError): assert True assert core._default_verify_function(instance, answer, result_host, 0.1, False) diff --git a/test/test_julia_functions.py b/test/test_julia_functions.py new file mode 100644 index 000000000..d8666cb57 --- /dev/null +++ b/test/test_julia_functions.py @@ -0,0 +1,73 @@ +from warnings import warn +import numpy as np +import pytest + +from kernel_tuner import tune_kernel +from kernel_tuner.backends.julia import JuliaFunctions +from kernel_tuner.core import KernelInstance, KernelSource + +from .test_runners import env # noqa: F401 +from .context import skip_if_no_julia + + +kernel_name = "vector_add!" +kernel_string = r""" + using KernelAbstractions + + @kernel function vector_add!( + c, a, b, n, ::Val{block_size_x} = Val(128) + ) where {block_size_x} + i = @index(Global) + if i <= n + c[i] = a[i] + b[i] + end + end + """ + + +@skip_if_no_julia +def test_ready_argument_list(): + """Ensure Julia backend correctly converts arguments into Julia objects.""" + from juliacall import ValueBase + + size = 1000 + a = np.int32(75) + b = np.random.randn(size).astype(np.float32) + c = np.zeros_like(b) + + arguments = [c, a, b] + + dev = JuliaFunctions(0) + gpu_args = dev.ready_argument_list(arguments) + + # Julia Array maps back through PythonCall as pyjl_pointer-like proxies + # Scalars remain scalars + assert isinstance(gpu_args[0], ValueBase) # Julia GPU Array proxy + assert isinstance(gpu_args[1], (int, np.int32)) # scalar unchanged + assert isinstance(gpu_args[2], ValueBase) # Julia GPU Array proxy + + +@skip_if_no_julia +def test_compile(): + """Check that Julia kernel code successfully compiles.""" + kernel_sources = KernelSource(kernel_name, kernel_string, "julia") + kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), []) + + dev = JuliaFunctions(0) + + try: + dev.compile(kernel_instance) + except Exception as e: + pytest.fail("Did not expect any exception: " + str(e)) + + +@skip_if_no_julia +def test_tune_kernel(env): + """Run a minimal Julia kernel tuner example.""" + env[0] = kernel_name + env[1] = kernel_string + env[4] = list(env[4].items()) # convert from a dict to a list of tuples to preserve order + + result, _ = tune_kernel(*env, lang="julia", verbose=True) + + assert len(result) > 0 diff --git a/test/test_time_budgets.py b/test/test_time_budgets.py index 8773801c8..16b10fd0f 100644 --- a/test/test_time_budgets.py +++ b/test/test_time_budgets.py @@ -1,3 +1,4 @@ +# ruff: noqa from itertools import product from time import perf_counter @@ -7,7 +8,7 @@ from kernel_tuner import tune_kernel -from .context import skip_if_no_gcc +from .context import skip_if_no_gcc, running_on_ci @pytest.fixture @@ -44,9 +45,10 @@ def env(): @skip_if_no_gcc def test_no_time_budget(env): """Ensure that a RuntimeError is raised if the startup takes longer than the time budget.""" - with raises(RuntimeError, match='startup time of the tuning process'): + with raises(RuntimeError, match="startup time of the tuning process"): tune_kernel(*env, strategy="random_sample", strategy_options={"strategy": "random_sample", "time_limit": 0.0}) + @skip_if_no_gcc def test_some_time_budget(env): """Ensure that the time limit is respected.""" @@ -64,7 +66,8 @@ def test_some_time_budget(env): assert 0 < len(res) < size_all # Ensure that the time limit was respected by some margin. - assert perf_counter() - start_time < time_limit * 2 + assert perf_counter() - start_time < (time_limit * 2) * (1.5 if running_on_ci else 1.0) # allow for more overhead on CI + @skip_if_no_gcc def test_full_time_budget(env): diff --git a/test/test_util_functions.py b/test/test_util_functions.py index a4bd16f6a..2ab80253f 100644 --- a/test/test_util_functions.py +++ b/test/test_util_functions.py @@ -1,3 +1,4 @@ +# ruff: noqa from __future__ import print_function import json @@ -36,17 +37,13 @@ def test_get_grid_dimensions1(): assert grid[1] == 28 assert grid[2] == 1 - grid = get_grid_dimensions( - problem_size, params, (grid_div[0], None, None), block_size_names - ) + grid = get_grid_dimensions(problem_size, params, (grid_div[0], None, None), block_size_names) assert grid[0] == 25 assert grid[1] == 1024 assert grid[2] == 1 - grid = get_grid_dimensions( - problem_size, params, (None, grid_div[1], None), block_size_names - ) + grid = get_grid_dimensions(problem_size, params, (None, grid_div[1], None), block_size_names) assert grid[0] == 1024 assert grid[1] == 28 @@ -60,17 +57,13 @@ def test_get_grid_dimensions1(): assert grid[1] == 25 assert grid[2] == 1 - grid = get_grid_dimensions( - problem_size, params, ("41", 37, None), block_size_names - ) + grid = get_grid_dimensions(problem_size, params, ("41", 37, None), block_size_names) assert grid[0] == 25 assert grid[1] == 28 assert grid[2] == 1 - grid = get_grid_dimensions( - problem_size, params, (None, [2, "block_y"], None), block_size_names - ) + grid = get_grid_dimensions(problem_size, params, (None, [2, "block_y"], None), block_size_names) assert grid[0] == 1024 assert grid[1] == 14 @@ -84,9 +77,7 @@ def test_get_grid_dimensions2(): grid_div_x = ["block_x*8"] grid_div_y = ["(block_y+2)/8"] - grid = get_grid_dimensions( - problem_size, params, (grid_div_x, grid_div_y, None), block_size_names - ) + grid = get_grid_dimensions(problem_size, params, (grid_div_x, grid_div_y, None), block_size_names) assert grid[0] == 4 assert grid[1] == 256 @@ -100,9 +91,7 @@ def test_get_grid_dimensions3(): grid_div_y = ["(block_y+2)/8"] def assert_grid_dimensions(problem_size): - grid = get_grid_dimensions( - problem_size, params, (grid_div_x, grid_div_y, None), block_size_names - ) + grid = get_grid_dimensions(problem_size, params, (grid_div_x, grid_div_y, None), block_size_names) assert grid[0] == 1 assert grid[1] == 256 assert grid[2] == 1 @@ -188,15 +177,13 @@ def test_prepare_kernel_string(): defines = dict(foo=1, bar="custom", baz=lambda config: config["is"] * 5) _, output = prepare_kernel_string("this", kernel, params, grid, threads, block_size_names, "", defines) - expected = "#define foo 1\n" "#define bar custom\n" "#define baz 40\n" "#line 1\n" "this is a weird kernel" + expected = "#define foo 1\n#define bar custom\n#define baz 40\n#line 1\nthis is a weird kernel" assert output == expected # Throw exception on invalid name (for instance, a space in the name) invalid_defines = {"invalid name": "1"} with pytest.raises(ValueError): - prepare_kernel_string( - "this", kernel, params, grid, threads, block_size_names, "", invalid_defines - ) + prepare_kernel_string("this", kernel, params, grid, threads, block_size_names, "", invalid_defines) def test_prepare_kernel_string_partial_loop_unrolling(): @@ -211,9 +198,7 @@ def test_prepare_kernel_string_partial_loop_unrolling(): params = dict() params["loop_unroll_factor_monkey"] = 8 - _, output = prepare_kernel_string( - "this", kernel, params, grid, threads, block_size_names, "CUDA", None - ) + _, output = prepare_kernel_string("this", kernel, params, grid, threads, block_size_names, "CUDA", None) assert "constexpr int loop_unroll_factor_monkey = 8;" in output params["loop_unroll_factor_monkey"] = 0 @@ -221,6 +206,7 @@ def test_prepare_kernel_string_partial_loop_unrolling(): assert "constexpr int loop_unroll_factor_monkey" not in output assert "#pragma unroll loop_unroll_factor_monkey" not in output + def test_replace_param_occurrences(): kernel = "this is a weird kernel" params = dict() @@ -228,9 +214,7 @@ def test_replace_param_occurrences(): params["weird"] = 14 new_kernel = replace_param_occurrences(kernel, params) - assert ( - new_kernel == "this 8 a 14 kernel" - ) # Note: The "is" in "this" should not be replaced + assert new_kernel == "this 8 a 14 kernel" # Note: The "is" in "this" should not be replaced new_kernel = replace_param_occurrences(kernel, dict()) assert kernel == new_kernel @@ -358,9 +342,7 @@ def test_check_argument_list3(): } """ args = [np.uint16(42), np.float16([3, 4, 6]), np.int32([300])] - assert_user_warning( - check_argument_list, [kernel_name, kernel_string, args], "at position 2" - ) + assert_user_warning(check_argument_list, [kernel_name, kernel_string, args], "at position 2") def test_check_argument_list4(): @@ -370,9 +352,7 @@ def test_check_argument_list4(): } """ args = [np.uint16(42), np.float16([3, 4, 6]), np.int64([300]), np.ubyte(32)] - assert_user_warning( - check_argument_list, [kernel_name, kernel_string, args], "do not match in size" - ) + assert_user_warning(check_argument_list, [kernel_name, kernel_string, args], "do not match in size") def test_check_argument_list5(): @@ -576,18 +556,12 @@ def test_warnings(function, args, number, warning_type): # check warning does not triger when nondefault block size names are used correctly block_size_names = ["block_size_a", "block_size_b"] - tune_params = dict( - zip(["block_size_a", "block_size_b", "many_other_things"], [1, 2, 3]) - ) - test_warnings( - check_block_size_params_names_list, [block_size_names, tune_params], 0, None - ) + tune_params = dict(zip(["block_size_a", "block_size_b", "many_other_things"], [1, 2, 3])) + test_warnings(check_block_size_params_names_list, [block_size_names, tune_params], 0, None) # check that a warning is issued when none of the default names are used and no alternative names are specified block_size_names = None - tune_params = dict( - zip(["block_size_a", "block_size_b", "many_other_things"], [1, 2, 3]) - ) + tune_params = dict(zip(["block_size_a", "block_size_b", "many_other_things"], [1, 2, 3])) test_warnings( check_block_size_params_names_list, [block_size_names, tune_params], @@ -597,12 +571,8 @@ def test_warnings(function, args, number, warning_type): # check that no error is raised when any of the default block size names is being used block_size_names = None - tune_params = dict( - zip(["block_size_x", "several_other_things"], [[1, 2, 3, 4], [2, 4]]) - ) - test_warnings( - check_block_size_params_names_list, [block_size_names, tune_params], 0, None - ) + tune_params = dict(zip(["block_size_x", "several_other_things"], [[1, 2, 3, 4], [2, 4]])) + test_warnings(check_block_size_params_names_list, [block_size_names, tune_params], 0, None) def test_get_kernel_string_func(): @@ -802,10 +772,7 @@ def test_process_metrics(): # assert params["b"] == 15 # test if a metric overrides any existing metrics - params = { - "x": 15, - "b": 12 - } + params = {"x": 15, "b": 12} metrics = dict() metrics["b"] = "x" params = process_metrics(params, metrics) @@ -815,7 +782,11 @@ def test_process_metrics(): def test_parse_restrictions(): tune_params = {"block_size_x": [50, 100], "use_padding": [0, 1]} restrict = ["block_size_x != 320"] - restrictions = ["block_size_x != 320", "use_padding == 0 or block_size_x % 32 != 0", "50 <= block_size_x * use_padding < 100"] + restrictions = [ + "block_size_x != 320", + "use_padding == 0 or block_size_x % 32 != 0", + "50 <= block_size_x * use_padding < 100", + ] # test the monolithic parsed function parsed = parse_restrictions(restrict, tune_params, monolithic=True)[0] @@ -842,9 +813,9 @@ def test_check_matching_problem_size(): with pytest.raises(ValueError): check_matching_problem_size(42, 1000) with pytest.raises(ValueError): - check_matching_problem_size([42,1], 42) + check_matching_problem_size([42, 1], 42) with pytest.raises(ValueError): - check_matching_problem_size([42,0], 42) + check_matching_problem_size([42, 0], 42) with pytest.raises(ValueError): check_matching_problem_size(None, 42) # these should not error @@ -852,28 +823,39 @@ def test_check_matching_problem_size(): check_matching_problem_size([1000], 1000) check_matching_problem_size(1000, 1000) check_matching_problem_size(1000, [1000]) - check_matching_problem_size([1000,], 1000) + check_matching_problem_size( + [ + 1000, + ], + 1000, + ) def test_convert_constraint_lambdas(): - restrictions = [lambda p: 32 <= p["block_size_x"]*p["block_size_y"] <= 1024, - "32 <= block_size_x*block_size_y <= 512", - lambda p: p["block_size_z"] < 8] + restrictions = [ + lambda p: 32 <= p["block_size_x"] * p["block_size_y"] <= 1024, + "32 <= block_size_x*block_size_y <= 512", + lambda p: p["block_size_z"] < 8, + ] result = convert_constraint_lambdas(restrictions) print(result) - expected = ['32 <= block_size_x * block_size_y <= 1024', 'block_size_z < 8', '32 <= block_size_x*block_size_y <= 512'] + expected = [ + "32 <= block_size_x * block_size_y <= 1024", + "block_size_z < 8", + "32 <= block_size_x*block_size_y <= 512", + ] assert sorted(result) == sorted(expected) restrictions2 = [] - restrictions2 += [lambda p: 32 <= p["block_size_x"]*p["block_size_y"] <= 1024] + restrictions2 += [lambda p: 32 <= p["block_size_x"] * p["block_size_y"] <= 1024] restrictions2 += [lambda p: p["block_size_z"] < 8] result2 = convert_constraint_lambdas(restrictions2) print(result2) - expected2 = ['block_size_z < 8', '32 <= block_size_x * block_size_y <= 1024'] + expected2 = ["block_size_z < 8", "32 <= block_size_x * block_size_y <= 1024"] assert sorted(result2) == sorted(expected2) @@ -885,16 +867,20 @@ def test_convert_constraint_lambdas_illformatted(): That is why this test expects an exception """ - restrictions = ["32 <= block_size_x*block_size_y <= 512", - lambda p: 32 <= p["block_size_x"]*p["block_size_y"] <= 1024, - lambda p: p["block_size_z"] < 8] + restrictions = [ + "32 <= block_size_x*block_size_y <= 512", + lambda p: 32 <= p["block_size_x"] * p["block_size_y"] <= 1024, + lambda p: p["block_size_z"] < 8, + ] - expected = ['32 <= block_size_x * block_size_y <= 1024', 'block_size_z < 8', '32 <= block_size_x*block_size_y <= 512'] + expected = [ + "32 <= block_size_x * block_size_y <= 1024", + "block_size_z < 8", + "32 <= block_size_x*block_size_y <= 512", + ] try: result = convert_constraint_lambdas(restrictions) print(result) except ValueError: pass - - diff --git a/test/utils/nvcuda.py b/test/utils/nvcuda.py index 4d05440cf..73ec4c949 100644 --- a/test/utils/nvcuda.py +++ b/test/utils/nvcuda.py @@ -1,3 +1,4 @@ +# ruff: noqa: D100, D103 from kernel_tuner.utils.nvcuda import to_valid_nvrtc_gpu_arch_cc diff --git a/test/utils/test_directives.py b/test/utils/test_directives.py index 759a20a54..b96867bd6 100644 --- a/test/utils/test_directives.py +++ b/test/utils/test_directives.py @@ -1,3 +1,4 @@ +# ruff: noqa from kernel_tuner.utils.directives import * @@ -283,16 +284,10 @@ def test_extract_directive_signature(): code = "#pragma tuner start vector_add a(float*:VECTOR_SIZE) b(float*:VECTOR_SIZE) c(float*:VECTOR_SIZE) size(int:VECTOR_SIZE) \n#pragma acc" signatures = extract_directive_signature(code, acc_cxx) assert len(signatures) == 1 - assert ( - "float vector_add(float * a, float * b, float * c, int size)" - in signatures["vector_add"] - ) + assert "float vector_add(float * a, float * b, float * c, int size)" in signatures["vector_add"] signatures = extract_directive_signature(code, acc_cxx, "vector_add") assert len(signatures) == 1 - assert ( - "float vector_add(float * a, float * b, float * c, int size)" - in signatures["vector_add"] - ) + assert "float vector_add(float * a, float * b, float * c, int size)" in signatures["vector_add"] signatures = extract_directive_signature(code, acc_cxx, "vector_add_ext") assert len(signatures) == 0 code = "!$tuner start vector_add A(float*:VECTOR_SIZE) B(float*:VECTOR_SIZE) C(float*:VECTOR_SIZE) n(int:VECTOR_SIZE)\n!$acc"