From 8e92a65fcd3263ece2e7cc9e8d9e46f9bf9ec6a9 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 18 Mar 2026 03:15:28 +0000 Subject: [PATCH 1/7] Initial plan From eeda712fab4df7c8346a96001b09aff9016f1721 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 18 Mar 2026 03:25:27 +0000 Subject: [PATCH 2/7] Add initial pytest test suite and CI integration - tests/conftest.py: session-scoped HIP binary fixtures, GPU/hipcc detection helpers - tests/test_imports.py: verify all public API symbols are importable - tests/test_kernels.py: kernel discovery, line mapping, instruction extraction, filtering, basic blocks, Kernel wrapper - tests/test_arguments.py: argument metadata, nested structs, typedef resolution, template instantiations - pyproject.toml: add [tool.pytest.ini_options] with testpaths and -v flag - .github/workflows/ci.yml: add 'Run pytest' step after editable install Co-authored-by: mawad-amd <112003944+mawad-amd@users.noreply.github.com> --- .github/workflows/ci.yml | 6 + pyproject.toml | 5 + tests/conftest.py | 260 +++++++++++++++++++++++++ tests/test_arguments.py | 362 +++++++++++++++++++++++++++++++++++ tests/test_imports.py | 125 ++++++++++++ tests/test_kernels.py | 401 +++++++++++++++++++++++++++++++++++++++ 6 files changed, 1159 insertions(+) create mode 100644 tests/conftest.py create mode 100644 tests/test_arguments.py create mode 100644 tests/test_imports.py create mode 100644 tests/test_kernels.py diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 30dc1b7..8899f12 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -38,3 +38,9 @@ jobs: pip install -e '.[dev]' python3 -c 'import kerneldb' " + - name: Run pytest + run: | + bash .github/scripts/container_exec.sh " + pip install -e '.[dev]' + pytest tests/ -v + " diff --git a/pyproject.toml b/pyproject.toml index ad167e0..be2068f 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -36,6 +36,11 @@ kerneldb = ["*.so", "*.so.*"] [tool.setuptools.dynamic] version = {attr = "setuptools_scm.get_version"} +# ---- pytest configuration ---- +[tool.pytest.ini_options] +testpaths = ["tests"] +addopts = "-v" + # ---- setuptools-scm versioning ---- [tool.setuptools_scm] version_scheme = "post-release" # .postN after last tag diff --git a/tests/conftest.py b/tests/conftest.py new file mode 100644 index 0000000..4c5e76a --- /dev/null +++ b/tests/conftest.py @@ -0,0 +1,260 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +""" +Shared pytest fixtures and helpers for kernelDB tests. + +Tests in this suite require a ROCm environment with: + - hipcc compiler (to compile HIP test kernels) + - An HSA-compatible GPU (AMD CDNA architecture) + +Tests are automatically skipped if these requirements are not met. +""" + +import shutil +import subprocess +import sys +import pytest +from pathlib import Path + + +# --------------------------------------------------------------------------- +# Environment detection helpers +# --------------------------------------------------------------------------- + + +def _hipcc_available() -> bool: + """Return True if the hipcc compiler is on PATH.""" + return shutil.which("hipcc") is not None + + +def _gpu_available() -> bool: + """Return True if kernelDB can initialise HSA and find a GPU agent.""" + try: + from kerneldb import _kerneldb # type: ignore[import] + + status = _kerneldb.hsa_init() + if status != 0: + return False + agent = _kerneldb.get_first_gpu_agent() + return agent.handle != 0 + except (ImportError, AttributeError): + return False + except RuntimeError: + return False + + +def _compile_hip(src_path: Path, out_path: Path, extra_flags: list[str] | None = None) -> None: + """Compile *src_path* with hipcc and write the executable to *out_path*. + + Raises ``pytest.skip`` on compilation failure so individual tests can + choose how to handle the case. + """ + cmd = ["hipcc", "-g", str(src_path), "-o", str(out_path)] + if extra_flags: + cmd.extend(extra_flags) + result = subprocess.run(cmd, capture_output=True, text=True) + if result.returncode != 0: + pytest.skip(f"HIP compilation failed:\n{result.stderr}") + + +# --------------------------------------------------------------------------- +# pytest marks / skip decorators +# --------------------------------------------------------------------------- + +requires_hipcc = pytest.mark.skipif( + not _hipcc_available(), + reason="hipcc not available – ROCm toolchain required", +) + +requires_gpu = pytest.mark.skipif( + not _gpu_available(), + reason="No HSA GPU agent found – AMD GPU with ROCm required", +) + +requires_rocm = pytest.mark.skipif( + not (_hipcc_available() and _gpu_available()), + reason="ROCm environment (hipcc + GPU) required", +) + + +# --------------------------------------------------------------------------- +# HIP source snippets used across multiple test modules +# --------------------------------------------------------------------------- + +SIMPLE_KERNELS_SOURCE = r""" +#include + +__global__ void vector_add(float* a, float* b, float* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] + b[idx]; + } +} + +__global__ void vector_multiply(float* a, float* b, float* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] * b[idx]; + } +} + +int main() { return 0; } +""" + +ARGUMENTS_SOURCE = r""" +#include + +__global__ void kernel_with_args(double* a, double* b, double* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] + b[idx]; + } +} + +int main() { return 0; } +""" + +NESTED_STRUCTS_SOURCE = r""" +#include + +struct Point3D { + float x; + float y; + float z; +}; + +struct BoundingBox { + Point3D min_pt; + Point3D max_pt; +}; + +__global__ void update_bounds(BoundingBox* boxes, int count) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < count) { + boxes[idx].min_pt.x = 0.0f; + } +} + +int main() { return 0; } +""" + +TYPEDEF_SOURCE = r""" +#include + +using MyInt = int; +using MyFloat = float; +typedef double MyDouble; + +__global__ void typedef_kernel(MyInt a, MyFloat b, MyDouble c, int* d) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx == 0) { + d[0] = a + static_cast(b) + static_cast(c); + } +} + +int main() { return 0; } +""" + +TEMPLATE_SOURCE = r""" +#include + +template +__global__ void scale_values(T* input, T* output, T factor, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + output[idx] = input[idx] * factor; + } +} + +// Explicit instantiations so the kernels appear in the binary +template __global__ void scale_values(float*, float*, float, int); +template __global__ void scale_values(double*, double*, double, int); + +int main() { return 0; } +""" + + +# --------------------------------------------------------------------------- +# Session-scoped compiled binary fixtures +# --------------------------------------------------------------------------- + + +@pytest.fixture(scope="session") +def simple_binary(tmp_path_factory): + """Path to a compiled HIP binary containing vector_add and vector_multiply.""" + if not _hipcc_available(): + pytest.skip("hipcc not available") + if not _gpu_available(): + pytest.skip("No GPU agent found") + + src_dir = tmp_path_factory.mktemp("hip_simple") + cpp_file = src_dir / "simple_kernels.cpp" + exe_file = src_dir / "simple_kernels" + cpp_file.write_text(SIMPLE_KERNELS_SOURCE) + _compile_hip(cpp_file, exe_file) + return str(exe_file) + + +@pytest.fixture(scope="session") +def arguments_binary(tmp_path_factory): + """Path to a compiled HIP binary with typed kernel arguments.""" + if not _hipcc_available(): + pytest.skip("hipcc not available") + if not _gpu_available(): + pytest.skip("No GPU agent found") + + src_dir = tmp_path_factory.mktemp("hip_args") + cpp_file = src_dir / "args_kernel.cpp" + exe_file = src_dir / "args_kernel" + cpp_file.write_text(ARGUMENTS_SOURCE) + _compile_hip(cpp_file, exe_file) + return str(exe_file) + + +@pytest.fixture(scope="session") +def nested_structs_binary(tmp_path_factory): + """Path to a compiled HIP binary with nested struct kernel arguments.""" + if not _hipcc_available(): + pytest.skip("hipcc not available") + if not _gpu_available(): + pytest.skip("No GPU agent found") + + src_dir = tmp_path_factory.mktemp("hip_nested") + cpp_file = src_dir / "nested_structs.cpp" + exe_file = src_dir / "nested_structs" + cpp_file.write_text(NESTED_STRUCTS_SOURCE) + _compile_hip(cpp_file, exe_file) + return str(exe_file) + + +@pytest.fixture(scope="session") +def typedef_binary(tmp_path_factory): + """Path to a compiled HIP binary with typedef/using-alias kernel arguments.""" + if not _hipcc_available(): + pytest.skip("hipcc not available") + if not _gpu_available(): + pytest.skip("No GPU agent found") + + src_dir = tmp_path_factory.mktemp("hip_typedef") + cpp_file = src_dir / "typedef_kernel.cpp" + exe_file = src_dir / "typedef_kernel" + cpp_file.write_text(TYPEDEF_SOURCE) + _compile_hip(cpp_file, exe_file) + return str(exe_file) + + +@pytest.fixture(scope="session") +def template_binary(tmp_path_factory): + """Path to a compiled HIP binary with template kernel instantiations.""" + if not _hipcc_available(): + pytest.skip("hipcc not available") + if not _gpu_available(): + pytest.skip("No GPU agent found") + + src_dir = tmp_path_factory.mktemp("hip_template") + cpp_file = src_dir / "template_kernel.cpp" + exe_file = src_dir / "template_kernel" + cpp_file.write_text(TEMPLATE_SOURCE) + _compile_hip(cpp_file, exe_file) + return str(exe_file) diff --git a/tests/test_arguments.py b/tests/test_arguments.py new file mode 100644 index 0000000..cdeae76 --- /dev/null +++ b/tests/test_arguments.py @@ -0,0 +1,362 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +""" +Tests for KernelDB kernel argument extraction: + - Basic argument metadata (name, type, size, alignment) + - Pointer argument sizes + - Nested struct member recursion + - Typedef / using-alias resolution + +All tests in this module require a ROCm environment with hipcc and a GPU. +""" + +import pytest +from conftest import requires_rocm + +kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available", exc_type=ImportError) +KernelDB = kerneldb.KernelDB + + +# --------------------------------------------------------------------------- +# Helpers +# --------------------------------------------------------------------------- + + +def _find_kernel(kernels: list[str], fragment: str) -> str: + matches = [k for k in kernels if fragment in k] + if not matches: + pytest.skip(f"No kernel containing {fragment!r} found in binary") + return matches[0] + + +def _arg_by_name(arguments, name: str): + for arg in arguments: + if arg.name == name: + return arg + return None + + +# --------------------------------------------------------------------------- +# Basic argument metadata +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_get_kernel_arguments_returns_list(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + assert isinstance(args, list) + + +@requires_rocm +def test_get_kernel_arguments_non_empty(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + assert len(args) > 0, "Expected kernel_with_args to have argument information" + + +@requires_rocm +def test_get_kernel_arguments_correct_count(arguments_binary): + """kernel_with_args(double*, double*, double*, int) has 4 parameters.""" + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + assert len(args) == 4, f"Expected 4 arguments, got {len(args)}" + + +@requires_rocm +def test_arguments_have_name_attribute(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + for arg in kdb.get_kernel_arguments(kernel_name): + assert hasattr(arg, "name") + assert isinstance(arg.name, str) + assert len(arg.name) > 0 + + +@requires_rocm +def test_arguments_have_type_name_attribute(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + for arg in kdb.get_kernel_arguments(kernel_name): + assert hasattr(arg, "type_name") + assert isinstance(arg.type_name, str) + assert len(arg.type_name) > 0 + + +@requires_rocm +def test_arguments_have_size_attribute(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + for arg in kdb.get_kernel_arguments(kernel_name): + assert hasattr(arg, "size") + assert isinstance(arg.size, int) + assert arg.size > 0 + + +@requires_rocm +def test_arguments_have_alignment_attribute(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + for arg in kdb.get_kernel_arguments(kernel_name): + assert hasattr(arg, "alignment") + assert isinstance(arg.alignment, int) + assert arg.alignment > 0 + + +@requires_rocm +def test_argument_names_match_source(arguments_binary): + """Argument names should match the source parameter names a, b, c, n.""" + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + names = [arg.name for arg in args] + for expected in ("a", "b", "c", "n"): + assert expected in names, f"Expected argument {expected!r} in {names}" + + +@requires_rocm +def test_pointer_argument_size(arguments_binary): + """Pointer arguments (double*) must be 8 bytes on a 64-bit platform.""" + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + ptr_arg = _arg_by_name(args, "a") + assert ptr_arg is not None, "Argument 'a' not found" + assert ptr_arg.size == 8, f"Expected pointer size 8, got {ptr_arg.size}" + + +@requires_rocm +def test_int_argument_size(arguments_binary): + """The 'n' argument is an int, which must be 4 bytes.""" + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + n_arg = _arg_by_name(args, "n") + assert n_arg is not None, "Argument 'n' not found" + assert n_arg.size == 4, f"Expected int size 4, got {n_arg.size}" + + +@requires_rocm +def test_int_argument_type_name(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(kernel_name) + n_arg = _arg_by_name(args, "n") + assert n_arg is not None, "Argument 'n' not found" + assert "int" in n_arg.type_name.lower() + + +# --------------------------------------------------------------------------- +# Kernel wrapper: has_arguments / arguments property +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_kernel_wrapper_has_arguments_true(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + kernel = kdb.get_kernel(kernel_name) + assert kernel.has_arguments() is True + + +@requires_rocm +def test_kernel_wrapper_arguments_property(arguments_binary): + kdb = KernelDB(arguments_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + kernel = kdb.get_kernel(kernel_name) + args = kernel.arguments + assert isinstance(args, list) + assert len(args) > 0 + + +# --------------------------------------------------------------------------- +# Nested struct member recursion +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_nested_structs_kernel_found(nested_structs_binary): + kdb = KernelDB(nested_structs_binary) + kernels = kdb.get_kernels() + assert any("update_bounds" in k for k in kernels), ( + f"'update_bounds' not found in {kernels}" + ) + + +@requires_rocm +def test_nested_structs_argument_has_members(nested_structs_binary): + """The BoundingBox argument must expose its Point3D sub-members.""" + kdb = KernelDB(nested_structs_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") + args = kdb.get_kernel_arguments(kernel_name) + + # Find the BoundingBox argument (passed by pointer as first arg) + bbox_arg = None + for arg in args: + if "BoundingBox" in arg.type_name or "boxes" == arg.name: + bbox_arg = arg + break + + assert bbox_arg is not None, ( + f"Expected a BoundingBox/boxes argument, got {[a.type_name for a in args]}" + ) + + +@requires_rocm +def test_nested_structs_members_attribute_is_list(nested_structs_binary): + """Every KernelArgument must expose a members attribute.""" + kdb = KernelDB(nested_structs_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") + for arg in kdb.get_kernel_arguments(kernel_name): + assert hasattr(arg, "members") + + +@requires_rocm +def test_nested_structs_point3d_has_three_members(nested_structs_binary): + """Point3D has x, y, z – any Point3D sub-argument must have 3 members.""" + kdb = KernelDB(nested_structs_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") + args = kdb.get_kernel_arguments(kernel_name) + + def find_point3d(arg_list): + for arg in arg_list: + if "Point3D" in arg.type_name: + return arg + if arg.members: + result = find_point3d(arg.members) + if result: + return result + return None + + point = find_point3d(args) + if point is None: + pytest.skip("Point3D type not found in argument tree (may be optimized away)") + + assert len(point.members) == 3, ( + f"Expected 3 members in Point3D, got {len(point.members)}" + ) + + +# --------------------------------------------------------------------------- +# Typedef / using-alias resolution +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_typedef_argument_count(typedef_binary): + """typedef_kernel(MyInt, MyFloat, MyDouble, int*) has 4 parameters.""" + kdb = KernelDB(typedef_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") + args = kdb.get_kernel_arguments(kernel_name) + assert len(args) == 4, f"Expected 4 arguments, got {len(args)}" + + +@requires_rocm +def test_typedef_not_resolved_by_default(typedef_binary): + """With resolve_typedefs=False the alias names (MyInt etc.) are preserved.""" + kdb = KernelDB(typedef_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") + args = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=False) + type_names = [arg.type_name for arg in args] + + # At least one argument should retain its typedef name + has_typedef_name = any( + name in ("MyInt", "MyFloat", "MyDouble") for name in type_names + ) + assert has_typedef_name, ( + f"Expected typedef names to be preserved, got types: {type_names}" + ) + + +@requires_rocm +def test_typedef_resolved_when_requested(typedef_binary): + """With resolve_typedefs=True, typedef names map to their underlying types.""" + kdb = KernelDB(typedef_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") + args_resolved = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=True) + type_names = [arg.type_name for arg in args_resolved] + + # After resolution, original typedef names like MyInt/MyFloat should not appear + for alias in ("MyInt", "MyFloat", "MyDouble"): + assert alias not in type_names, ( + f"Typedef alias {alias!r} still present after resolution: {type_names}" + ) + + +@requires_rocm +def test_typedef_resolved_types_are_primitives(typedef_binary): + """After typedef resolution the types should be C primitive type names.""" + kdb = KernelDB(typedef_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") + args_resolved = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=True) + type_names = [arg.type_name for arg in args_resolved] + + primitive_types = {"int", "float", "double", "unsigned int"} + for type_name in type_names: + # Check that each resolved type is recognisably a C/C++ primitive or pointer to one. + # Strip pointer qualifiers (* and spaces) before matching. + base_type = type_name.replace("*", "").strip() + assert any(prim in base_type for prim in primitive_types), ( + f"Resolved type {type_name!r} does not look like a primitive type" + ) + + +# --------------------------------------------------------------------------- +# Template kernel instantiations +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_template_kernels_found(template_binary): + """Both scale_values and scale_values must appear.""" + kdb = KernelDB(template_binary) + kernels = kdb.get_kernels() + scale_kernels = [k for k in kernels if "scale_values" in k] + assert len(scale_kernels) >= 2, ( + f"Expected at least 2 template instantiations, got {scale_kernels}" + ) + + +@requires_rocm +def test_template_float_instantiation_arguments(template_binary): + """scale_values must have 4 arguments.""" + kdb = KernelDB(template_binary) + kernels = kdb.get_kernels() + float_kernels = [k for k in kernels if "scale_values" in k and "float" in k] + if not float_kernels: + pytest.skip("scale_values not found in binary") + args = kdb.get_kernel_arguments(float_kernels[0]) + assert len(args) == 4, f"Expected 4 arguments for scale_values, got {len(args)}" + + +@requires_rocm +def test_template_float_argument_types(template_binary): + """scale_values argument types should reference float.""" + kdb = KernelDB(template_binary) + kernels = kdb.get_kernels() + float_kernels = [k for k in kernels if "scale_values" in k and "float" in k] + if not float_kernels: + pytest.skip("scale_values not found in binary") + args = kdb.get_kernel_arguments(float_kernels[0]) + # Exclude 'n' (the int count parameter) + non_int_args = [a for a in args if a.name != "n"] + for arg in non_int_args: + assert "float" in arg.type_name.lower(), ( + f"Expected float type for {arg.name!r}, got {arg.type_name!r}" + ) + + +@requires_rocm +def test_template_double_instantiation_arguments(template_binary): + """scale_values must have 4 arguments.""" + kdb = KernelDB(template_binary) + kernels = kdb.get_kernels() + double_kernels = [k for k in kernels if "scale_values" in k and "double" in k] + if not double_kernels: + pytest.skip("scale_values not found in binary") + args = kdb.get_kernel_arguments(double_kernels[0]) + assert len(args) == 4, f"Expected 4 arguments for scale_values, got {len(args)}" diff --git a/tests/test_imports.py b/tests/test_imports.py new file mode 100644 index 0000000..452f192 --- /dev/null +++ b/tests/test_imports.py @@ -0,0 +1,125 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +""" +Tests that verify the public Python API of kernelDB can be imported correctly. + +These tests require the kernelDB C++ extension to be compiled and installed. +They are automatically skipped when the extension is not available. +""" + +import types +import pytest + +# Skip the entire module if the kernelDB C++ extension is not available +kerneldb = pytest.importorskip( + "kerneldb", + reason="kernelDB C++ extension not available – build with ROCm toolchain", + exc_type=ImportError, +) + + +# --------------------------------------------------------------------------- +# Module-level import tests +# --------------------------------------------------------------------------- + + +def test_import_kerneldb_top_level(): + """The top-level kerneldb package must be importable.""" + import kerneldb # noqa: F401 + + +def test_version_attribute_exists(): + """kerneldb.__version__ must be a non-empty string.""" + import kerneldb + + assert hasattr(kerneldb, "__version__") + assert isinstance(kerneldb.__version__, str) + assert len(kerneldb.__version__) > 0 + + +def test_all_attribute_exists(): + """kerneldb.__all__ must be a non-empty list of strings.""" + import kerneldb + + assert hasattr(kerneldb, "__all__") + assert isinstance(kerneldb.__all__, list) + assert len(kerneldb.__all__) > 0 + for name in kerneldb.__all__: + assert isinstance(name, str), f"__all__ entry {name!r} is not a string" + + +def test_all_entries_are_accessible(): + """Every name listed in __all__ must be accessible as an attribute.""" + import kerneldb + + for name in kerneldb.__all__: + assert hasattr(kerneldb, name), f"kerneldb.{name} listed in __all__ but not accessible" + + +# --------------------------------------------------------------------------- +# Public class import tests +# --------------------------------------------------------------------------- + + +def test_import_KernelDB(): + from kerneldb import KernelDB # noqa: F401 + + +def test_import_Kernel(): + from kerneldb import Kernel # noqa: F401 + + +def test_import_Instruction(): + from kerneldb import Instruction # noqa: F401 + + +def test_import_BasicBlock(): + from kerneldb import BasicBlock # noqa: F401 + + +def test_import_CDNAKernel(): + from kerneldb import CDNAKernel # noqa: F401 + + +def test_import_ArchDescriptor(): + from kerneldb import ArchDescriptor # noqa: F401 + + +def test_import_HsaAgent(): + from kerneldb import HsaAgent # noqa: F401 + + +def test_import_KernelArgument(): + from kerneldb import KernelArgument # noqa: F401 + + +def test_import_all_public_classes_at_once(): + """All public classes can be imported in a single statement.""" + from kerneldb import ( # noqa: F401 + KernelDB, + Kernel, + Instruction, + BasicBlock, + CDNAKernel, + ArchDescriptor, + HsaAgent, + KernelArgument, + ) + + +# --------------------------------------------------------------------------- +# Type / callable tests (no GPU required) +# --------------------------------------------------------------------------- + + +def test_KernelDB_is_class(): + from kerneldb import KernelDB + + assert isinstance(KernelDB, type) + + +def test_Kernel_is_class(): + from kerneldb import Kernel + + assert isinstance(Kernel, type) diff --git a/tests/test_kernels.py b/tests/test_kernels.py new file mode 100644 index 0000000..c3367e2 --- /dev/null +++ b/tests/test_kernels.py @@ -0,0 +1,401 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +""" +Tests for KernelDB kernel-level operations: + - Kernel discovery (get_kernels) + - Source-line mapping (get_kernel_lines) + - Instruction extraction (get_instructions_for_line) + - Instruction regex filtering + - Basic block extraction + - High-level Kernel wrapper properties + +All tests in this module require a ROCm environment with hipcc and a GPU. +""" + +import re +import pytest +from conftest import requires_rocm + +kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available", exc_type=ImportError) +KernelDB = kerneldb.KernelDB + + +# --------------------------------------------------------------------------- +# Helpers +# --------------------------------------------------------------------------- + + +def _find_kernel(kernels: list[str], fragment: str) -> str: + """Return the first kernel name that contains *fragment*, or skip.""" + matches = [k for k in kernels if fragment in k] + if not matches: + pytest.skip(f"No kernel containing {fragment!r} found in binary") + return matches[0] + + +# --------------------------------------------------------------------------- +# KernelDB initialisation +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_kerneldb_init_with_binary(simple_binary): + """KernelDB must initialise without raising given a valid binary path.""" + kdb = KernelDB(simple_binary) + assert kdb is not None + + +@requires_rocm +def test_kerneldb_binary_path_stored(simple_binary): + """The binary_path attribute must reflect the path that was given.""" + kdb = KernelDB(simple_binary) + assert kdb.binary_path == simple_binary + + +# --------------------------------------------------------------------------- +# Kernel discovery +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_get_kernels_returns_list(simple_binary): + kdb = KernelDB(simple_binary) + kernels = kdb.get_kernels() + assert isinstance(kernels, list) + + +@requires_rocm +def test_get_kernels_non_empty(simple_binary): + kdb = KernelDB(simple_binary) + kernels = kdb.get_kernels() + assert len(kernels) > 0, "Expected at least one kernel in the binary" + + +@requires_rocm +def test_get_kernels_names_are_strings(simple_binary): + kdb = KernelDB(simple_binary) + for name in kdb.get_kernels(): + assert isinstance(name, str) + assert len(name) > 0 + + +@requires_rocm +def test_get_kernels_contains_vector_add(simple_binary): + kdb = KernelDB(simple_binary) + kernels = kdb.get_kernels() + assert any("vector_add" in k for k in kernels), ( + f"'vector_add' not found in kernel list: {kernels}" + ) + + +@requires_rocm +def test_get_kernels_contains_vector_multiply(simple_binary): + kdb = KernelDB(simple_binary) + kernels = kdb.get_kernels() + assert any("vector_multiply" in k for k in kernels), ( + f"'vector_multiply' not found in kernel list: {kernels}" + ) + + +@requires_rocm +def test_get_kernels_two_kernels(simple_binary): + """The simple binary has exactly two kernels.""" + kdb = KernelDB(simple_binary) + kernels = kdb.get_kernels() + assert len(kernels) == 2, f"Expected 2 kernels, got {len(kernels)}: {kernels}" + + +# --------------------------------------------------------------------------- +# Source-line mapping +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_get_kernel_lines_returns_list(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + lines = kdb.get_kernel_lines(kernel_name) + assert isinstance(lines, list) + + +@requires_rocm +def test_get_kernel_lines_non_empty(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + lines = kdb.get_kernel_lines(kernel_name) + assert len(lines) > 0, "Expected at least one source line in kernel" + + +@requires_rocm +def test_get_kernel_lines_are_positive_integers(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + for line in kdb.get_kernel_lines(kernel_name): + assert isinstance(line, int) + assert line > 0, f"Line number must be positive, got {line}" + + +@requires_rocm +def test_different_kernels_have_different_lines(simple_binary): + """vector_add and vector_multiply live on different source lines.""" + kdb = KernelDB(simple_binary) + add_name = _find_kernel(kdb.get_kernels(), "vector_add") + mul_name = _find_kernel(kdb.get_kernels(), "vector_multiply") + + add_lines = set(kdb.get_kernel_lines(add_name)) + mul_lines = set(kdb.get_kernel_lines(mul_name)) + + # They may share some lines (e.g., the common expression line) but not all + assert add_lines != mul_lines, ( + "vector_add and vector_multiply reported identical line sets" + ) + + +# --------------------------------------------------------------------------- +# Instruction extraction +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_get_instructions_for_line_returns_list(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + first_line = kdb.get_kernel_lines(kernel_name)[0] + instructions = kdb.get_instructions_for_line(kernel_name, first_line) + assert isinstance(instructions, list) + + +@requires_rocm +def test_get_instructions_for_line_non_empty(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + first_line = kdb.get_kernel_lines(kernel_name)[0] + instructions = kdb.get_instructions_for_line(kernel_name, first_line) + assert len(instructions) > 0 + + +@requires_rocm +def test_instructions_have_disassembly_attribute(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + first_line = kdb.get_kernel_lines(kernel_name)[0] + instructions = kdb.get_instructions_for_line(kernel_name, first_line) + for inst in instructions: + assert hasattr(inst, "disassembly") + assert isinstance(inst.disassembly, str) + assert len(inst.disassembly) > 0 + + +@requires_rocm +def test_instructions_have_line_attribute(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + first_line = kdb.get_kernel_lines(kernel_name)[0] + instructions = kdb.get_instructions_for_line(kernel_name, first_line) + for inst in instructions: + assert hasattr(inst, "line") + assert isinstance(inst.line, int) + + +@requires_rocm +def test_instructions_have_column_attribute(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + first_line = kdb.get_kernel_lines(kernel_name)[0] + instructions = kdb.get_instructions_for_line(kernel_name, first_line) + for inst in instructions: + assert hasattr(inst, "column") + assert isinstance(inst.column, int) + + +@requires_rocm +def test_instructions_have_file_name_attribute(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + first_line = kdb.get_kernel_lines(kernel_name)[0] + instructions = kdb.get_instructions_for_line(kernel_name, first_line) + for inst in instructions: + assert hasattr(inst, "file_name") + + +@requires_rocm +def test_instruction_line_numbers_match_queried_line(simple_binary): + """Instructions returned for a line should report that same line number.""" + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + for line in kdb.get_kernel_lines(kernel_name): + for inst in kdb.get_instructions_for_line(kernel_name, line): + assert inst.line == line, ( + f"Expected inst.line == {line}, got {inst.line}" + ) + + +# --------------------------------------------------------------------------- +# Instruction filtering +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_filter_pattern_returns_subset(simple_binary): + """Filtering instructions reduces or keeps the same count.""" + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + all_instructions = [] + filtered_instructions = [] + for line in kdb.get_kernel_lines(kernel_name): + all_instructions.extend(kdb.get_instructions_for_line(kernel_name, line)) + filtered_instructions.extend( + kdb.get_instructions_for_line(kernel_name, line, ".*load.*") + ) + assert len(filtered_instructions) <= len(all_instructions) + + +@requires_rocm +def test_filter_pattern_matches_disassembly(simple_binary): + """Each instruction returned by a filter must match the pattern.""" + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + pattern = ".*load.*" + for line in kdb.get_kernel_lines(kernel_name): + for inst in kdb.get_instructions_for_line(kernel_name, line, pattern): + assert re.search(pattern, inst.disassembly, re.IGNORECASE), ( + f"Instruction {inst.disassembly!r} does not match pattern {pattern!r}" + ) + + +@requires_rocm +def test_load_store_filter_finds_memory_ops(simple_binary): + """The vector_add kernel must have at least one load/store instruction.""" + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + mem_ops = [] + for line in kdb.get_kernel_lines(kernel_name): + mem_ops.extend( + kdb.get_instructions_for_line(kernel_name, line, ".*(load|store).*") + ) + assert len(mem_ops) > 0, "Expected at least one memory operation in vector_add" + + +@requires_rocm +def test_no_match_filter_returns_empty(simple_binary): + """A filter that matches nothing must return an empty list per line.""" + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + for line in kdb.get_kernel_lines(kernel_name): + result = kdb.get_instructions_for_line( + kernel_name, line, "THISDOESNOTEXISTINANYASM__xyz" + ) + assert result == [], ( + f"Expected empty list for impossible filter, got {result}" + ) + + +# --------------------------------------------------------------------------- +# Basic blocks +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_get_basic_blocks_returns_list(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + blocks = kernel.get_basic_blocks() + assert isinstance(blocks, list) + + +@requires_rocm +def test_get_basic_blocks_non_empty(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + blocks = kernel.get_basic_blocks() + assert len(blocks) > 0, "Expected at least one basic block in vector_add" + + +# --------------------------------------------------------------------------- +# Kernel wrapper (high-level API) +# --------------------------------------------------------------------------- + + +@requires_rocm +def test_kernel_wrapper_name(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + assert kernel.name == kernel_name + + +@requires_rocm +def test_kernel_wrapper_lines_property(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + assert isinstance(kernel.lines, list) + assert len(kernel.lines) > 0 + for line in kernel.lines: + assert isinstance(line, int) and line > 0 + + +@requires_rocm +def test_kernel_wrapper_assembly_property(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + asm = kernel.assembly + assert isinstance(asm, list) + assert len(asm) > 0 + for line in asm: + assert isinstance(line, str) and len(line) > 0 + + +@requires_rocm +def test_kernel_wrapper_files_property(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + files = kernel.files + assert isinstance(files, list) + # Source files should be strings + for f in files: + assert isinstance(f, str) + + +@requires_rocm +def test_kernel_wrapper_signature_property(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + # signature is an alias for name + assert isinstance(kernel.signature, str) + assert len(kernel.signature) > 0 + + +@requires_rocm +def test_kernel_wrapper_get_instructions_for_line(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + first_line = kernel.lines[0] + instructions = kernel.get_instructions_for_line(first_line) + assert isinstance(instructions, list) + assert len(instructions) > 0 + + +@requires_rocm +def test_kernel_wrapper_get_instructions_with_filter(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(kernel_name) + all_insts = kernel.get_instructions_for_line(kernel.lines[0]) + filtered = kernel.get_instructions_for_line(kernel.lines[0], ".*load.*") + assert len(filtered) <= len(all_insts) + + +@requires_rocm +def test_kernel_wrapper_get_file_name(simple_binary): + kdb = KernelDB(simple_binary) + kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") + result = kdb.get_file_name(kernel_name, 0) + assert isinstance(result, str) From 98cec6253e1e9f807dc0097d85264c0142d3a566 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 18 Mar 2026 03:35:45 +0000 Subject: [PATCH 3/7] Simplify tests and fix pytest 7.x compatibility (remove exc_type arg) - Remove exc_type=ImportError from pytest.importorskip() calls (added in pytest 8.x; CI runs 7.3.2) - Consolidate granular single-assertion tests into broader per-feature tests: test_imports.py: 15 -> 3 tests test_kernels.py: 35 -> 5 tests test_arguments.py: 25 -> 5 tests Co-authored-by: mawad-amd <112003944+mawad-amd@users.noreply.github.com> --- tests/test_arguments.py | 318 +++++++--------------------------- tests/test_imports.py | 124 +++----------- tests/test_kernels.py | 367 +++++++--------------------------------- 3 files changed, 140 insertions(+), 669 deletions(-) diff --git a/tests/test_arguments.py b/tests/test_arguments.py index cdeae76..2c67c15 100644 --- a/tests/test_arguments.py +++ b/tests/test_arguments.py @@ -4,9 +4,9 @@ """ Tests for KernelDB kernel argument extraction: - Basic argument metadata (name, type, size, alignment) - - Pointer argument sizes - Nested struct member recursion - Typedef / using-alias resolution + - Template kernel instantiations All tests in this module require a ROCm environment with hipcc and a GPU. """ @@ -14,27 +14,19 @@ import pytest from conftest import requires_rocm -kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available", exc_type=ImportError) +kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB -# --------------------------------------------------------------------------- -# Helpers -# --------------------------------------------------------------------------- - - -def _find_kernel(kernels: list[str], fragment: str) -> str: +def _find_kernel(kernels, fragment): matches = [k for k in kernels if fragment in k] if not matches: pytest.skip(f"No kernel containing {fragment!r} found in binary") return matches[0] -def _arg_by_name(arguments, name: str): - for arg in arguments: - if arg.name == name: - return arg - return None +def _arg_by_name(arguments, name): + return next((a for a in arguments if a.name == name), None) # --------------------------------------------------------------------------- @@ -43,134 +35,40 @@ def _arg_by_name(arguments, name: str): @requires_rocm -def test_get_kernel_arguments_returns_list(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) - assert isinstance(args, list) - - -@requires_rocm -def test_get_kernel_arguments_non_empty(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) - assert len(args) > 0, "Expected kernel_with_args to have argument information" - - -@requires_rocm -def test_get_kernel_arguments_correct_count(arguments_binary): - """kernel_with_args(double*, double*, double*, int) has 4 parameters.""" +def test_kernel_arguments(arguments_binary): + """kernel_with_args must expose 4 named arguments with expected types/sizes.""" kdb = KernelDB(arguments_binary) kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") args = kdb.get_kernel_arguments(kernel_name) - assert len(args) == 4, f"Expected 4 arguments, got {len(args)}" - - -@requires_rocm -def test_arguments_have_name_attribute(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - for arg in kdb.get_kernel_arguments(kernel_name): - assert hasattr(arg, "name") - assert isinstance(arg.name, str) - assert len(arg.name) > 0 - - -@requires_rocm -def test_arguments_have_type_name_attribute(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - for arg in kdb.get_kernel_arguments(kernel_name): - assert hasattr(arg, "type_name") - assert isinstance(arg.type_name, str) - assert len(arg.type_name) > 0 - - -@requires_rocm -def test_arguments_have_size_attribute(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - for arg in kdb.get_kernel_arguments(kernel_name): - assert hasattr(arg, "size") - assert isinstance(arg.size, int) - assert arg.size > 0 - - -@requires_rocm -def test_arguments_have_alignment_attribute(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - for arg in kdb.get_kernel_arguments(kernel_name): - assert hasattr(arg, "alignment") - assert isinstance(arg.alignment, int) - assert arg.alignment > 0 + assert isinstance(args, list) and len(args) == 4 + for arg in args: + assert isinstance(arg.name, str) and arg.name + assert isinstance(arg.type_name, str) and arg.type_name + assert isinstance(arg.size, int) and arg.size > 0 + assert isinstance(arg.alignment, int) and arg.alignment > 0 -@requires_rocm -def test_argument_names_match_source(arguments_binary): - """Argument names should match the source parameter names a, b, c, n.""" - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) - names = [arg.name for arg in args] + names = [a.name for a in args] for expected in ("a", "b", "c", "n"): assert expected in names, f"Expected argument {expected!r} in {names}" - -@requires_rocm -def test_pointer_argument_size(arguments_binary): - """Pointer arguments (double*) must be 8 bytes on a 64-bit platform.""" - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) ptr_arg = _arg_by_name(args, "a") - assert ptr_arg is not None, "Argument 'a' not found" - assert ptr_arg.size == 8, f"Expected pointer size 8, got {ptr_arg.size}" - + assert ptr_arg.size == 8, f"Pointer size must be 8 bytes, got {ptr_arg.size}" -@requires_rocm -def test_int_argument_size(arguments_binary): - """The 'n' argument is an int, which must be 4 bytes.""" - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) n_arg = _arg_by_name(args, "n") - assert n_arg is not None, "Argument 'n' not found" - assert n_arg.size == 4, f"Expected int size 4, got {n_arg.size}" - - -@requires_rocm -def test_int_argument_type_name(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) - n_arg = _arg_by_name(args, "n") - assert n_arg is not None, "Argument 'n' not found" + assert n_arg.size == 4, f"int size must be 4 bytes, got {n_arg.size}" assert "int" in n_arg.type_name.lower() -# --------------------------------------------------------------------------- -# Kernel wrapper: has_arguments / arguments property -# --------------------------------------------------------------------------- - - @requires_rocm -def test_kernel_wrapper_has_arguments_true(arguments_binary): +def test_kernel_wrapper_arguments(arguments_binary): + """Kernel.has_arguments() and .arguments must reflect the extracted args.""" kdb = KernelDB(arguments_binary) kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") kernel = kdb.get_kernel(kernel_name) - assert kernel.has_arguments() is True - -@requires_rocm -def test_kernel_wrapper_arguments_property(arguments_binary): - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - kernel = kdb.get_kernel(kernel_name) - args = kernel.arguments - assert isinstance(args, list) - assert len(args) > 0 + assert kernel.has_arguments() is True + assert isinstance(kernel.arguments, list) and kernel.arguments # --------------------------------------------------------------------------- @@ -179,66 +77,31 @@ def test_kernel_wrapper_arguments_property(arguments_binary): @requires_rocm -def test_nested_structs_kernel_found(nested_structs_binary): - kdb = KernelDB(nested_structs_binary) - kernels = kdb.get_kernels() - assert any("update_bounds" in k for k in kernels), ( - f"'update_bounds' not found in {kernels}" - ) - - -@requires_rocm -def test_nested_structs_argument_has_members(nested_structs_binary): - """The BoundingBox argument must expose its Point3D sub-members.""" +def test_nested_struct_arguments(nested_structs_binary): + """update_bounds args must include a BoundingBox arg exposing Point3D members.""" kdb = KernelDB(nested_structs_binary) kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") args = kdb.get_kernel_arguments(kernel_name) - # Find the BoundingBox argument (passed by pointer as first arg) - bbox_arg = None + # Every argument must have a members attribute for arg in args: - if "BoundingBox" in arg.type_name or "boxes" == arg.name: - bbox_arg = arg - break - - assert bbox_arg is not None, ( - f"Expected a BoundingBox/boxes argument, got {[a.type_name for a in args]}" - ) - - -@requires_rocm -def test_nested_structs_members_attribute_is_list(nested_structs_binary): - """Every KernelArgument must expose a members attribute.""" - kdb = KernelDB(nested_structs_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") - for arg in kdb.get_kernel_arguments(kernel_name): assert hasattr(arg, "members") - -@requires_rocm -def test_nested_structs_point3d_has_three_members(nested_structs_binary): - """Point3D has x, y, z – any Point3D sub-argument must have 3 members.""" - kdb = KernelDB(nested_structs_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") - args = kdb.get_kernel_arguments(kernel_name) - - def find_point3d(arg_list): + # Helper: recursively search for a type by name + def find_type(arg_list, type_fragment): for arg in arg_list: - if "Point3D" in arg.type_name: + if type_fragment in arg.type_name: return arg if arg.members: - result = find_point3d(arg.members) - if result: - return result + found = find_type(arg.members, type_fragment) + if found: + return found return None - point = find_point3d(args) + point = find_type(args, "Point3D") if point is None: - pytest.skip("Point3D type not found in argument tree (may be optimized away)") - - assert len(point.members) == 3, ( - f"Expected 3 members in Point3D, got {len(point.members)}" - ) + pytest.skip("Point3D not found in argument tree (may be optimised away)") + assert len(point.members) == 3, f"Point3D should have 3 members, got {len(point.members)}" # --------------------------------------------------------------------------- @@ -247,61 +110,31 @@ def find_point3d(arg_list): @requires_rocm -def test_typedef_argument_count(typedef_binary): - """typedef_kernel(MyInt, MyFloat, MyDouble, int*) has 4 parameters.""" +def test_typedef_resolution(typedef_binary): + """typedef_kernel must have 4 args; typedefs preserved or resolved correctly.""" kdb = KernelDB(typedef_binary) kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") - args = kdb.get_kernel_arguments(kernel_name) - assert len(args) == 4, f"Expected 4 arguments, got {len(args)}" - -@requires_rocm -def test_typedef_not_resolved_by_default(typedef_binary): - """With resolve_typedefs=False the alias names (MyInt etc.) are preserved.""" - kdb = KernelDB(typedef_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") - args = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=False) - type_names = [arg.type_name for arg in args] + args_raw = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=False) + assert len(args_raw) == 4 - # At least one argument should retain its typedef name - has_typedef_name = any( - name in ("MyInt", "MyFloat", "MyDouble") for name in type_names + type_names_raw = [a.type_name for a in args_raw] + assert any(n in ("MyInt", "MyFloat", "MyDouble") for n in type_names_raw), ( + f"Expected typedef names to be preserved, got: {type_names_raw}" ) - assert has_typedef_name, ( - f"Expected typedef names to be preserved, got types: {type_names}" - ) - -@requires_rocm -def test_typedef_resolved_when_requested(typedef_binary): - """With resolve_typedefs=True, typedef names map to their underlying types.""" - kdb = KernelDB(typedef_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") args_resolved = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=True) - type_names = [arg.type_name for arg in args_resolved] - - # After resolution, original typedef names like MyInt/MyFloat should not appear + type_names_resolved = [a.type_name for a in args_resolved] for alias in ("MyInt", "MyFloat", "MyDouble"): - assert alias not in type_names, ( - f"Typedef alias {alias!r} still present after resolution: {type_names}" + assert alias not in type_names_resolved, ( + f"Alias {alias!r} should have been resolved, got: {type_names_resolved}" ) - -@requires_rocm -def test_typedef_resolved_types_are_primitives(typedef_binary): - """After typedef resolution the types should be C primitive type names.""" - kdb = KernelDB(typedef_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") - args_resolved = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=True) - type_names = [arg.type_name for arg in args_resolved] - primitive_types = {"int", "float", "double", "unsigned int"} - for type_name in type_names: - # Check that each resolved type is recognisably a C/C++ primitive or pointer to one. - # Strip pointer qualifiers (* and spaces) before matching. - base_type = type_name.replace("*", "").strip() - assert any(prim in base_type for prim in primitive_types), ( - f"Resolved type {type_name!r} does not look like a primitive type" + for type_name in type_names_resolved: + base = type_name.replace("*", "").strip() + assert any(p in base for p in primitive_types), ( + f"Resolved type {type_name!r} does not look like a C primitive" ) @@ -311,52 +144,21 @@ def test_typedef_resolved_types_are_primitives(typedef_binary): @requires_rocm -def test_template_kernels_found(template_binary): - """Both scale_values and scale_values must appear.""" +def test_template_kernel_arguments(template_binary): + """scale_values must be instantiated for at least float and double.""" kdb = KernelDB(template_binary) kernels = kdb.get_kernels() - scale_kernels = [k for k in kernels if "scale_values" in k] - assert len(scale_kernels) >= 2, ( - f"Expected at least 2 template instantiations, got {scale_kernels}" - ) - -@requires_rocm -def test_template_float_instantiation_arguments(template_binary): - """scale_values must have 4 arguments.""" - kdb = KernelDB(template_binary) - kernels = kdb.get_kernels() - float_kernels = [k for k in kernels if "scale_values" in k and "float" in k] - if not float_kernels: - pytest.skip("scale_values not found in binary") - args = kdb.get_kernel_arguments(float_kernels[0]) - assert len(args) == 4, f"Expected 4 arguments for scale_values, got {len(args)}" - - -@requires_rocm -def test_template_float_argument_types(template_binary): - """scale_values argument types should reference float.""" - kdb = KernelDB(template_binary) - kernels = kdb.get_kernels() - float_kernels = [k for k in kernels if "scale_values" in k and "float" in k] - if not float_kernels: - pytest.skip("scale_values not found in binary") - args = kdb.get_kernel_arguments(float_kernels[0]) - # Exclude 'n' (the int count parameter) - non_int_args = [a for a in args if a.name != "n"] - for arg in non_int_args: - assert "float" in arg.type_name.lower(), ( - f"Expected float type for {arg.name!r}, got {arg.type_name!r}" + scale_kernels = [k for k in kernels if "scale_values" in k] + assert len(scale_kernels) >= 2, f"Expected ≥2 template instantiations, got {scale_kernels}" + + for type_fragment in ("float", "double"): + matching = [k for k in scale_kernels if type_fragment in k] + if not matching: + pytest.skip(f"scale_values<{type_fragment}> not found in binary") + args = kdb.get_kernel_arguments(matching[0]) + assert len(args) == 4, f"Expected 4 args for scale_values<{type_fragment}>" + non_int = [a for a in args if a.name != "n"] + assert all(type_fragment in a.type_name.lower() for a in non_int), ( + f"Non-count args should be {type_fragment}: {[a.type_name for a in non_int]}" ) - - -@requires_rocm -def test_template_double_instantiation_arguments(template_binary): - """scale_values must have 4 arguments.""" - kdb = KernelDB(template_binary) - kernels = kdb.get_kernels() - double_kernels = [k for k in kernels if "scale_values" in k and "double" in k] - if not double_kernels: - pytest.skip("scale_values not found in binary") - args = kdb.get_kernel_arguments(double_kernels[0]) - assert len(args) == 4, f"Expected 4 arguments for scale_values, got {len(args)}" diff --git a/tests/test_imports.py b/tests/test_imports.py index 452f192..4e9a284 100644 --- a/tests/test_imports.py +++ b/tests/test_imports.py @@ -8,118 +8,40 @@ They are automatically skipped when the extension is not available. """ -import types import pytest # Skip the entire module if the kernelDB C++ extension is not available kerneldb = pytest.importorskip( "kerneldb", reason="kernelDB C++ extension not available – build with ROCm toolchain", - exc_type=ImportError, ) +PUBLIC_CLASSES = [ + "KernelDB", + "Kernel", + "Instruction", + "BasicBlock", + "CDNAKernel", + "ArchDescriptor", + "HsaAgent", + "KernelArgument", +] -# --------------------------------------------------------------------------- -# Module-level import tests -# --------------------------------------------------------------------------- +def test_package_metadata(): + """kerneldb must expose a non-empty __version__ and a valid __all__.""" + assert isinstance(kerneldb.__version__, str) and kerneldb.__version__ + assert isinstance(kerneldb.__all__, list) and kerneldb.__all__ + assert all(isinstance(n, str) for n in kerneldb.__all__) -def test_import_kerneldb_top_level(): - """The top-level kerneldb package must be importable.""" - import kerneldb # noqa: F401 +def test_public_api_accessible(): + """All expected public classes must be importable from the package.""" + for name in PUBLIC_CLASSES: + obj = getattr(kerneldb, name, None) + assert obj is not None, f"kerneldb.{name} not accessible" -def test_version_attribute_exists(): - """kerneldb.__version__ must be a non-empty string.""" - import kerneldb - assert hasattr(kerneldb, "__version__") - assert isinstance(kerneldb.__version__, str) - assert len(kerneldb.__version__) > 0 - - -def test_all_attribute_exists(): - """kerneldb.__all__ must be a non-empty list of strings.""" - import kerneldb - - assert hasattr(kerneldb, "__all__") - assert isinstance(kerneldb.__all__, list) - assert len(kerneldb.__all__) > 0 - for name in kerneldb.__all__: - assert isinstance(name, str), f"__all__ entry {name!r} is not a string" - - -def test_all_entries_are_accessible(): - """Every name listed in __all__ must be accessible as an attribute.""" - import kerneldb - - for name in kerneldb.__all__: - assert hasattr(kerneldb, name), f"kerneldb.{name} listed in __all__ but not accessible" - - -# --------------------------------------------------------------------------- -# Public class import tests -# --------------------------------------------------------------------------- - - -def test_import_KernelDB(): - from kerneldb import KernelDB # noqa: F401 - - -def test_import_Kernel(): - from kerneldb import Kernel # noqa: F401 - - -def test_import_Instruction(): - from kerneldb import Instruction # noqa: F401 - - -def test_import_BasicBlock(): - from kerneldb import BasicBlock # noqa: F401 - - -def test_import_CDNAKernel(): - from kerneldb import CDNAKernel # noqa: F401 - - -def test_import_ArchDescriptor(): - from kerneldb import ArchDescriptor # noqa: F401 - - -def test_import_HsaAgent(): - from kerneldb import HsaAgent # noqa: F401 - - -def test_import_KernelArgument(): - from kerneldb import KernelArgument # noqa: F401 - - -def test_import_all_public_classes_at_once(): - """All public classes can be imported in a single statement.""" - from kerneldb import ( # noqa: F401 - KernelDB, - Kernel, - Instruction, - BasicBlock, - CDNAKernel, - ArchDescriptor, - HsaAgent, - KernelArgument, - ) - - -# --------------------------------------------------------------------------- -# Type / callable tests (no GPU required) -# --------------------------------------------------------------------------- - - -def test_KernelDB_is_class(): - from kerneldb import KernelDB - - assert isinstance(KernelDB, type) - - -def test_Kernel_is_class(): - from kerneldb import Kernel - - assert isinstance(Kernel, type) +def test_KernelDB_and_Kernel_are_classes(): + assert isinstance(kerneldb.KernelDB, type) + assert isinstance(kerneldb.Kernel, type) diff --git a/tests/test_kernels.py b/tests/test_kernels.py index c3367e2..aef08e6 100644 --- a/tests/test_kernels.py +++ b/tests/test_kernels.py @@ -5,8 +5,7 @@ Tests for KernelDB kernel-level operations: - Kernel discovery (get_kernels) - Source-line mapping (get_kernel_lines) - - Instruction extraction (get_instructions_for_line) - - Instruction regex filtering + - Instruction extraction and regex filtering (get_instructions_for_line) - Basic block extraction - High-level Kernel wrapper properties @@ -17,16 +16,11 @@ import pytest from conftest import requires_rocm -kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available", exc_type=ImportError) +kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB -# --------------------------------------------------------------------------- -# Helpers -# --------------------------------------------------------------------------- - - -def _find_kernel(kernels: list[str], fragment: str) -> str: +def _find_kernel(kernels, fragment): """Return the first kernel name that contains *fragment*, or skip.""" matches = [k for k in kernels if fragment in k] if not matches: @@ -34,76 +28,22 @@ def _find_kernel(kernels: list[str], fragment: str) -> str: return matches[0] -# --------------------------------------------------------------------------- -# KernelDB initialisation -# --------------------------------------------------------------------------- - - -@requires_rocm -def test_kerneldb_init_with_binary(simple_binary): - """KernelDB must initialise without raising given a valid binary path.""" - kdb = KernelDB(simple_binary) - assert kdb is not None - - -@requires_rocm -def test_kerneldb_binary_path_stored(simple_binary): - """The binary_path attribute must reflect the path that was given.""" - kdb = KernelDB(simple_binary) - assert kdb.binary_path == simple_binary - - # --------------------------------------------------------------------------- # Kernel discovery # --------------------------------------------------------------------------- @requires_rocm -def test_get_kernels_returns_list(simple_binary): +def test_kernel_discovery(simple_binary): + """Binary must expose two named kernels: vector_add and vector_multiply.""" kdb = KernelDB(simple_binary) kernels = kdb.get_kernels() - assert isinstance(kernels, list) - -@requires_rocm -def test_get_kernels_non_empty(simple_binary): - kdb = KernelDB(simple_binary) - kernels = kdb.get_kernels() - assert len(kernels) > 0, "Expected at least one kernel in the binary" - - -@requires_rocm -def test_get_kernels_names_are_strings(simple_binary): - kdb = KernelDB(simple_binary) - for name in kdb.get_kernels(): - assert isinstance(name, str) - assert len(name) > 0 - - -@requires_rocm -def test_get_kernels_contains_vector_add(simple_binary): - kdb = KernelDB(simple_binary) - kernels = kdb.get_kernels() - assert any("vector_add" in k for k in kernels), ( - f"'vector_add' not found in kernel list: {kernels}" - ) - - -@requires_rocm -def test_get_kernels_contains_vector_multiply(simple_binary): - kdb = KernelDB(simple_binary) - kernels = kdb.get_kernels() - assert any("vector_multiply" in k for k in kernels), ( - f"'vector_multiply' not found in kernel list: {kernels}" - ) - - -@requires_rocm -def test_get_kernels_two_kernels(simple_binary): - """The simple binary has exactly two kernels.""" - kdb = KernelDB(simple_binary) - kernels = kdb.get_kernels() - assert len(kernels) == 2, f"Expected 2 kernels, got {len(kernels)}: {kernels}" + assert isinstance(kernels, list) + assert len(kernels) == 2 + assert all(isinstance(k, str) and k for k in kernels) + assert any("vector_add" in k for k in kernels) + assert any("vector_multiply" in k for k in kernels) # --------------------------------------------------------------------------- @@ -112,44 +52,19 @@ def test_get_kernels_two_kernels(simple_binary): @requires_rocm -def test_get_kernel_lines_returns_list(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - lines = kdb.get_kernel_lines(kernel_name) - assert isinstance(lines, list) - - -@requires_rocm -def test_get_kernel_lines_non_empty(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - lines = kdb.get_kernel_lines(kernel_name) - assert len(lines) > 0, "Expected at least one source line in kernel" - - -@requires_rocm -def test_get_kernel_lines_are_positive_integers(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - for line in kdb.get_kernel_lines(kernel_name): - assert isinstance(line, int) - assert line > 0, f"Line number must be positive, got {line}" - - -@requires_rocm -def test_different_kernels_have_different_lines(simple_binary): - """vector_add and vector_multiply live on different source lines.""" +def test_kernel_lines(simple_binary): + """Each kernel must map to a non-empty list of positive source lines.""" kdb = KernelDB(simple_binary) - add_name = _find_kernel(kdb.get_kernels(), "vector_add") - mul_name = _find_kernel(kdb.get_kernels(), "vector_multiply") + kernels = kdb.get_kernels() - add_lines = set(kdb.get_kernel_lines(add_name)) - mul_lines = set(kdb.get_kernel_lines(mul_name)) + for kernel_name in kernels: + lines = kdb.get_kernel_lines(kernel_name) + assert isinstance(lines, list) and lines, f"No lines for {kernel_name}" + assert all(isinstance(ln, int) and ln > 0 for ln in lines) - # They may share some lines (e.g., the common expression line) but not all - assert add_lines != mul_lines, ( - "vector_add and vector_multiply reported identical line sets" - ) + add_name = _find_kernel(kernels, "vector_add") + mul_name = _find_kernel(kernels, "vector_multiply") + assert set(kdb.get_kernel_lines(add_name)) != set(kdb.get_kernel_lines(mul_name)) # --------------------------------------------------------------------------- @@ -158,77 +73,19 @@ def test_different_kernels_have_different_lines(simple_binary): @requires_rocm -def test_get_instructions_for_line_returns_list(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - first_line = kdb.get_kernel_lines(kernel_name)[0] - instructions = kdb.get_instructions_for_line(kernel_name, first_line) - assert isinstance(instructions, list) - - -@requires_rocm -def test_get_instructions_for_line_non_empty(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - first_line = kdb.get_kernel_lines(kernel_name)[0] - instructions = kdb.get_instructions_for_line(kernel_name, first_line) - assert len(instructions) > 0 - - -@requires_rocm -def test_instructions_have_disassembly_attribute(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - first_line = kdb.get_kernel_lines(kernel_name)[0] - instructions = kdb.get_instructions_for_line(kernel_name, first_line) - for inst in instructions: - assert hasattr(inst, "disassembly") - assert isinstance(inst.disassembly, str) - assert len(inst.disassembly) > 0 - - -@requires_rocm -def test_instructions_have_line_attribute(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - first_line = kdb.get_kernel_lines(kernel_name)[0] - instructions = kdb.get_instructions_for_line(kernel_name, first_line) - for inst in instructions: - assert hasattr(inst, "line") - assert isinstance(inst.line, int) - - -@requires_rocm -def test_instructions_have_column_attribute(simple_binary): +def test_instructions(simple_binary): + """Instructions for each source line must have the expected attributes.""" kdb = KernelDB(simple_binary) kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - first_line = kdb.get_kernel_lines(kernel_name)[0] - instructions = kdb.get_instructions_for_line(kernel_name, first_line) - for inst in instructions: - assert hasattr(inst, "column") - assert isinstance(inst.column, int) - -@requires_rocm -def test_instructions_have_file_name_attribute(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - first_line = kdb.get_kernel_lines(kernel_name)[0] - instructions = kdb.get_instructions_for_line(kernel_name, first_line) - for inst in instructions: - assert hasattr(inst, "file_name") - - -@requires_rocm -def test_instruction_line_numbers_match_queried_line(simple_binary): - """Instructions returned for a line should report that same line number.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") for line in kdb.get_kernel_lines(kernel_name): - for inst in kdb.get_instructions_for_line(kernel_name, line): - assert inst.line == line, ( - f"Expected inst.line == {line}, got {inst.line}" - ) + instructions = kdb.get_instructions_for_line(kernel_name, line) + assert isinstance(instructions, list) + for inst in instructions: + assert isinstance(inst.disassembly, str) and inst.disassembly + assert isinstance(inst.line, int) and inst.line == line + assert isinstance(inst.column, int) + assert hasattr(inst, "file_name") # --------------------------------------------------------------------------- @@ -237,165 +94,55 @@ def test_instruction_line_numbers_match_queried_line(simple_binary): @requires_rocm -def test_filter_pattern_returns_subset(simple_binary): - """Filtering instructions reduces or keeps the same count.""" +def test_instruction_filtering(simple_binary): + """Regex filtering must return a subset that each matches the pattern.""" kdb = KernelDB(simple_binary) kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - all_instructions = [] - filtered_instructions = [] - for line in kdb.get_kernel_lines(kernel_name): - all_instructions.extend(kdb.get_instructions_for_line(kernel_name, line)) - filtered_instructions.extend( - kdb.get_instructions_for_line(kernel_name, line, ".*load.*") - ) - assert len(filtered_instructions) <= len(all_instructions) + pattern = ".*(load|store).*" - -@requires_rocm -def test_filter_pattern_matches_disassembly(simple_binary): - """Each instruction returned by a filter must match the pattern.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - pattern = ".*load.*" + all_insts, mem_ops = [], [] for line in kdb.get_kernel_lines(kernel_name): - for inst in kdb.get_instructions_for_line(kernel_name, line, pattern): - assert re.search(pattern, inst.disassembly, re.IGNORECASE), ( - f"Instruction {inst.disassembly!r} does not match pattern {pattern!r}" - ) + all_insts.extend(kdb.get_instructions_for_line(kernel_name, line)) + mem_ops.extend(kdb.get_instructions_for_line(kernel_name, line, pattern)) + assert len(mem_ops) > 0, "vector_add must have at least one load/store" + assert len(mem_ops) <= len(all_insts) + assert all(re.search(pattern, inst.disassembly, re.IGNORECASE) for inst in mem_ops) -@requires_rocm -def test_load_store_filter_finds_memory_ops(simple_binary): - """The vector_add kernel must have at least one load/store instruction.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - mem_ops = [] + # A pattern that matches nothing returns empty lists for line in kdb.get_kernel_lines(kernel_name): - mem_ops.extend( - kdb.get_instructions_for_line(kernel_name, line, ".*(load|store).*") - ) - assert len(mem_ops) > 0, "Expected at least one memory operation in vector_add" - - -@requires_rocm -def test_no_match_filter_returns_empty(simple_binary): - """A filter that matches nothing must return an empty list per line.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - for line in kdb.get_kernel_lines(kernel_name): - result = kdb.get_instructions_for_line( - kernel_name, line, "THISDOESNOTEXISTINANYASM__xyz" - ) - assert result == [], ( - f"Expected empty list for impossible filter, got {result}" - ) + assert kdb.get_instructions_for_line(kernel_name, line, "NOMATCH__xyz") == [] # --------------------------------------------------------------------------- -# Basic blocks +# Basic blocks and Kernel wrapper # --------------------------------------------------------------------------- @requires_rocm -def test_get_basic_blocks_returns_list(simple_binary): +def test_basic_blocks_and_kernel_wrapper(simple_binary): + """Kernel wrapper properties and basic-block extraction must work correctly.""" kdb = KernelDB(simple_binary) kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") kernel = kdb.get_kernel(kernel_name) - blocks = kernel.get_basic_blocks() - assert isinstance(blocks, list) - -@requires_rocm -def test_get_basic_blocks_non_empty(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) + # basic blocks blocks = kernel.get_basic_blocks() - assert len(blocks) > 0, "Expected at least one basic block in vector_add" - + assert isinstance(blocks, list) and blocks -# --------------------------------------------------------------------------- -# Kernel wrapper (high-level API) -# --------------------------------------------------------------------------- - - -@requires_rocm -def test_kernel_wrapper_name(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) + # wrapper properties assert kernel.name == kernel_name + assert isinstance(kernel.signature, str) and kernel.signature + assert isinstance(kernel.lines, list) and all(ln > 0 for ln in kernel.lines) + assert isinstance(kernel.assembly, list) and all(isinstance(s, str) for s in kernel.assembly) + assert isinstance(kernel.files, list) and all(isinstance(f, str) for f in kernel.files) - -@requires_rocm -def test_kernel_wrapper_lines_property(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) - assert isinstance(kernel.lines, list) - assert len(kernel.lines) > 0 - for line in kernel.lines: - assert isinstance(line, int) and line > 0 - - -@requires_rocm -def test_kernel_wrapper_assembly_property(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) - asm = kernel.assembly - assert isinstance(asm, list) - assert len(asm) > 0 - for line in asm: - assert isinstance(line, str) and len(line) > 0 - - -@requires_rocm -def test_kernel_wrapper_files_property(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) - files = kernel.files - assert isinstance(files, list) - # Source files should be strings - for f in files: - assert isinstance(f, str) - - -@requires_rocm -def test_kernel_wrapper_signature_property(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) - # signature is an alias for name - assert isinstance(kernel.signature, str) - assert len(kernel.signature) > 0 - - -@requires_rocm -def test_kernel_wrapper_get_instructions_for_line(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) + # instructions via wrapper first_line = kernel.lines[0] - instructions = kernel.get_instructions_for_line(first_line) - assert isinstance(instructions, list) - assert len(instructions) > 0 - + insts = kernel.get_instructions_for_line(first_line) + assert isinstance(insts, list) and insts + filtered = kernel.get_instructions_for_line(first_line, ".*load.*") + assert len(filtered) <= len(insts) -@requires_rocm -def test_kernel_wrapper_get_instructions_with_filter(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) - all_insts = kernel.get_instructions_for_line(kernel.lines[0]) - filtered = kernel.get_instructions_for_line(kernel.lines[0], ".*load.*") - assert len(filtered) <= len(all_insts) - - -@requires_rocm -def test_kernel_wrapper_get_file_name(simple_binary): - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - result = kdb.get_file_name(kernel_name, 0) - assert isinstance(result, str) + # file name helper + assert isinstance(kdb.get_file_name(kernel_name, 0), str) From 0f69c52d38e12562c3078c7fe6d99907c5009506 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 18 Mar 2026 03:47:47 +0000 Subject: [PATCH 4/7] Fix failing tests and simplify fixtures/decorators - Fix test_kernel_lines and test_basic_blocks_and_kernel_wrapper: get_kernel_lines() can return line 0, relax assertion from ln > 0 to ln >= 0 - Remove unused requires_hipcc / requires_gpu / requires_rocm marks from conftest.py (fixtures already handle skipping) - Remove @requires_rocm decorator from all test functions in test_kernels.py and test_arguments.py (fixtures handle it) - Remove 'from conftest import requires_rocm' imports from test files Co-authored-by: mawad-amd <112003944+mawad-amd@users.noreply.github.com> --- tests/conftest.py | 20 -------------------- tests/test_arguments.py | 6 ------ tests/test_kernels.py | 10 ++-------- 3 files changed, 2 insertions(+), 34 deletions(-) diff --git a/tests/conftest.py b/tests/conftest.py index 4c5e76a..f6ac119 100644 --- a/tests/conftest.py +++ b/tests/conftest.py @@ -58,26 +58,6 @@ def _compile_hip(src_path: Path, out_path: Path, extra_flags: list[str] | None = pytest.skip(f"HIP compilation failed:\n{result.stderr}") -# --------------------------------------------------------------------------- -# pytest marks / skip decorators -# --------------------------------------------------------------------------- - -requires_hipcc = pytest.mark.skipif( - not _hipcc_available(), - reason="hipcc not available – ROCm toolchain required", -) - -requires_gpu = pytest.mark.skipif( - not _gpu_available(), - reason="No HSA GPU agent found – AMD GPU with ROCm required", -) - -requires_rocm = pytest.mark.skipif( - not (_hipcc_available() and _gpu_available()), - reason="ROCm environment (hipcc + GPU) required", -) - - # --------------------------------------------------------------------------- # HIP source snippets used across multiple test modules # --------------------------------------------------------------------------- diff --git a/tests/test_arguments.py b/tests/test_arguments.py index 2c67c15..ae97c42 100644 --- a/tests/test_arguments.py +++ b/tests/test_arguments.py @@ -12,7 +12,6 @@ """ import pytest -from conftest import requires_rocm kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB @@ -34,7 +33,6 @@ def _arg_by_name(arguments, name): # --------------------------------------------------------------------------- -@requires_rocm def test_kernel_arguments(arguments_binary): """kernel_with_args must expose 4 named arguments with expected types/sizes.""" kdb = KernelDB(arguments_binary) @@ -60,7 +58,6 @@ def test_kernel_arguments(arguments_binary): assert "int" in n_arg.type_name.lower() -@requires_rocm def test_kernel_wrapper_arguments(arguments_binary): """Kernel.has_arguments() and .arguments must reflect the extracted args.""" kdb = KernelDB(arguments_binary) @@ -76,7 +73,6 @@ def test_kernel_wrapper_arguments(arguments_binary): # --------------------------------------------------------------------------- -@requires_rocm def test_nested_struct_arguments(nested_structs_binary): """update_bounds args must include a BoundingBox arg exposing Point3D members.""" kdb = KernelDB(nested_structs_binary) @@ -109,7 +105,6 @@ def find_type(arg_list, type_fragment): # --------------------------------------------------------------------------- -@requires_rocm def test_typedef_resolution(typedef_binary): """typedef_kernel must have 4 args; typedefs preserved or resolved correctly.""" kdb = KernelDB(typedef_binary) @@ -143,7 +138,6 @@ def test_typedef_resolution(typedef_binary): # --------------------------------------------------------------------------- -@requires_rocm def test_template_kernel_arguments(template_binary): """scale_values must be instantiated for at least float and double.""" kdb = KernelDB(template_binary) diff --git a/tests/test_kernels.py b/tests/test_kernels.py index aef08e6..31b5757 100644 --- a/tests/test_kernels.py +++ b/tests/test_kernels.py @@ -14,7 +14,6 @@ import re import pytest -from conftest import requires_rocm kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB @@ -33,7 +32,6 @@ def _find_kernel(kernels, fragment): # --------------------------------------------------------------------------- -@requires_rocm def test_kernel_discovery(simple_binary): """Binary must expose two named kernels: vector_add and vector_multiply.""" kdb = KernelDB(simple_binary) @@ -51,7 +49,6 @@ def test_kernel_discovery(simple_binary): # --------------------------------------------------------------------------- -@requires_rocm def test_kernel_lines(simple_binary): """Each kernel must map to a non-empty list of positive source lines.""" kdb = KernelDB(simple_binary) @@ -60,7 +57,7 @@ def test_kernel_lines(simple_binary): for kernel_name in kernels: lines = kdb.get_kernel_lines(kernel_name) assert isinstance(lines, list) and lines, f"No lines for {kernel_name}" - assert all(isinstance(ln, int) and ln > 0 for ln in lines) + assert all(isinstance(ln, int) and ln >= 0 for ln in lines) add_name = _find_kernel(kernels, "vector_add") mul_name = _find_kernel(kernels, "vector_multiply") @@ -72,7 +69,6 @@ def test_kernel_lines(simple_binary): # --------------------------------------------------------------------------- -@requires_rocm def test_instructions(simple_binary): """Instructions for each source line must have the expected attributes.""" kdb = KernelDB(simple_binary) @@ -93,7 +89,6 @@ def test_instructions(simple_binary): # --------------------------------------------------------------------------- -@requires_rocm def test_instruction_filtering(simple_binary): """Regex filtering must return a subset that each matches the pattern.""" kdb = KernelDB(simple_binary) @@ -119,7 +114,6 @@ def test_instruction_filtering(simple_binary): # --------------------------------------------------------------------------- -@requires_rocm def test_basic_blocks_and_kernel_wrapper(simple_binary): """Kernel wrapper properties and basic-block extraction must work correctly.""" kdb = KernelDB(simple_binary) @@ -133,7 +127,7 @@ def test_basic_blocks_and_kernel_wrapper(simple_binary): # wrapper properties assert kernel.name == kernel_name assert isinstance(kernel.signature, str) and kernel.signature - assert isinstance(kernel.lines, list) and all(ln > 0 for ln in kernel.lines) + assert isinstance(kernel.lines, list) and all(ln >= 0 for ln in kernel.lines) assert isinstance(kernel.assembly, list) and all(isinstance(s, str) for s in kernel.assembly) assert isinstance(kernel.files, list) and all(isinstance(f, str) for f in kernel.files) From 340ffac0f924638f5db67361dd11066cac7b1c55 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 23 Mar 2026 00:04:31 -0700 Subject: [PATCH 5/7] Simplify tests: self-contained with no shared fixtures Each test file embeds its own HIP source and compiles inline. conftest.py reduced to just the requires_rocm skip marker. Co-Authored-By: Claude Opus 4.6 --- tests/conftest.py | 235 ++------------------------------------- tests/test_arguments.py | 236 +++++++++++++++++++++++++++------------- tests/test_imports.py | 33 +----- tests/test_kernels.py | 166 +++++++++++++++------------- 4 files changed, 262 insertions(+), 408 deletions(-) diff --git a/tests/conftest.py b/tests/conftest.py index f6ac119..48e409a 100644 --- a/tests/conftest.py +++ b/tests/conftest.py @@ -1,240 +1,27 @@ # SPDX-License-Identifier: MIT # Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. -""" -Shared pytest fixtures and helpers for kernelDB tests. - -Tests in this suite require a ROCm environment with: - - hipcc compiler (to compile HIP test kernels) - - An HSA-compatible GPU (AMD CDNA architecture) - -Tests are automatically skipped if these requirements are not met. -""" - import shutil -import subprocess -import sys -import pytest -from pathlib import Path - -# --------------------------------------------------------------------------- -# Environment detection helpers -# --------------------------------------------------------------------------- +import pytest -def _hipcc_available() -> bool: - """Return True if the hipcc compiler is on PATH.""" +def _hipcc_available(): return shutil.which("hipcc") is not None -def _gpu_available() -> bool: - """Return True if kernelDB can initialise HSA and find a GPU agent.""" +def _gpu_available(): try: - from kerneldb import _kerneldb # type: ignore[import] + from kerneldb import _kerneldb - status = _kerneldb.hsa_init() - if status != 0: + if _kerneldb.hsa_init() != 0: return False - agent = _kerneldb.get_first_gpu_agent() - return agent.handle != 0 - except (ImportError, AttributeError): + return _kerneldb.get_first_gpu_agent().handle != 0 + except Exception: return False - except RuntimeError: - return False - - -def _compile_hip(src_path: Path, out_path: Path, extra_flags: list[str] | None = None) -> None: - """Compile *src_path* with hipcc and write the executable to *out_path*. - - Raises ``pytest.skip`` on compilation failure so individual tests can - choose how to handle the case. - """ - cmd = ["hipcc", "-g", str(src_path), "-o", str(out_path)] - if extra_flags: - cmd.extend(extra_flags) - result = subprocess.run(cmd, capture_output=True, text=True) - if result.returncode != 0: - pytest.skip(f"HIP compilation failed:\n{result.stderr}") - - -# --------------------------------------------------------------------------- -# HIP source snippets used across multiple test modules -# --------------------------------------------------------------------------- - -SIMPLE_KERNELS_SOURCE = r""" -#include - -__global__ void vector_add(float* a, float* b, float* c, int n) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < n) { - c[idx] = a[idx] + b[idx]; - } -} - -__global__ void vector_multiply(float* a, float* b, float* c, int n) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < n) { - c[idx] = a[idx] * b[idx]; - } -} - -int main() { return 0; } -""" - -ARGUMENTS_SOURCE = r""" -#include - -__global__ void kernel_with_args(double* a, double* b, double* c, int n) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < n) { - c[idx] = a[idx] + b[idx]; - } -} - -int main() { return 0; } -""" - -NESTED_STRUCTS_SOURCE = r""" -#include - -struct Point3D { - float x; - float y; - float z; -}; - -struct BoundingBox { - Point3D min_pt; - Point3D max_pt; -}; - -__global__ void update_bounds(BoundingBox* boxes, int count) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < count) { - boxes[idx].min_pt.x = 0.0f; - } -} - -int main() { return 0; } -""" - -TYPEDEF_SOURCE = r""" -#include - -using MyInt = int; -using MyFloat = float; -typedef double MyDouble; - -__global__ void typedef_kernel(MyInt a, MyFloat b, MyDouble c, int* d) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx == 0) { - d[0] = a + static_cast(b) + static_cast(c); - } -} - -int main() { return 0; } -""" - -TEMPLATE_SOURCE = r""" -#include - -template -__global__ void scale_values(T* input, T* output, T factor, int n) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < n) { - output[idx] = input[idx] * factor; - } -} - -// Explicit instantiations so the kernels appear in the binary -template __global__ void scale_values(float*, float*, float, int); -template __global__ void scale_values(double*, double*, double, int); - -int main() { return 0; } -""" - - -# --------------------------------------------------------------------------- -# Session-scoped compiled binary fixtures -# --------------------------------------------------------------------------- - - -@pytest.fixture(scope="session") -def simple_binary(tmp_path_factory): - """Path to a compiled HIP binary containing vector_add and vector_multiply.""" - if not _hipcc_available(): - pytest.skip("hipcc not available") - if not _gpu_available(): - pytest.skip("No GPU agent found") - - src_dir = tmp_path_factory.mktemp("hip_simple") - cpp_file = src_dir / "simple_kernels.cpp" - exe_file = src_dir / "simple_kernels" - cpp_file.write_text(SIMPLE_KERNELS_SOURCE) - _compile_hip(cpp_file, exe_file) - return str(exe_file) - - -@pytest.fixture(scope="session") -def arguments_binary(tmp_path_factory): - """Path to a compiled HIP binary with typed kernel arguments.""" - if not _hipcc_available(): - pytest.skip("hipcc not available") - if not _gpu_available(): - pytest.skip("No GPU agent found") - - src_dir = tmp_path_factory.mktemp("hip_args") - cpp_file = src_dir / "args_kernel.cpp" - exe_file = src_dir / "args_kernel" - cpp_file.write_text(ARGUMENTS_SOURCE) - _compile_hip(cpp_file, exe_file) - return str(exe_file) - - -@pytest.fixture(scope="session") -def nested_structs_binary(tmp_path_factory): - """Path to a compiled HIP binary with nested struct kernel arguments.""" - if not _hipcc_available(): - pytest.skip("hipcc not available") - if not _gpu_available(): - pytest.skip("No GPU agent found") - - src_dir = tmp_path_factory.mktemp("hip_nested") - cpp_file = src_dir / "nested_structs.cpp" - exe_file = src_dir / "nested_structs" - cpp_file.write_text(NESTED_STRUCTS_SOURCE) - _compile_hip(cpp_file, exe_file) - return str(exe_file) - - -@pytest.fixture(scope="session") -def typedef_binary(tmp_path_factory): - """Path to a compiled HIP binary with typedef/using-alias kernel arguments.""" - if not _hipcc_available(): - pytest.skip("hipcc not available") - if not _gpu_available(): - pytest.skip("No GPU agent found") - - src_dir = tmp_path_factory.mktemp("hip_typedef") - cpp_file = src_dir / "typedef_kernel.cpp" - exe_file = src_dir / "typedef_kernel" - cpp_file.write_text(TYPEDEF_SOURCE) - _compile_hip(cpp_file, exe_file) - return str(exe_file) - -@pytest.fixture(scope="session") -def template_binary(tmp_path_factory): - """Path to a compiled HIP binary with template kernel instantiations.""" - if not _hipcc_available(): - pytest.skip("hipcc not available") - if not _gpu_available(): - pytest.skip("No GPU agent found") - src_dir = tmp_path_factory.mktemp("hip_template") - cpp_file = src_dir / "template_kernel.cpp" - exe_file = src_dir / "template_kernel" - cpp_file.write_text(TEMPLATE_SOURCE) - _compile_hip(cpp_file, exe_file) - return str(exe_file) +requires_rocm = pytest.mark.skipif( + not (_hipcc_available() and _gpu_available()), + reason="Requires hipcc and a GPU", +) diff --git a/tests/test_arguments.py b/tests/test_arguments.py index ae97c42..e8c0a6d 100644 --- a/tests/test_arguments.py +++ b/tests/test_arguments.py @@ -2,25 +2,43 @@ # Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. """ -Tests for KernelDB kernel argument extraction: - - Basic argument metadata (name, type, size, alignment) - - Nested struct member recursion - - Typedef / using-alias resolution - - Template kernel instantiations +Tests for KernelDB kernel argument extraction: basic metadata, nested structs, +typedef resolution, and template instantiations. -All tests in this module require a ROCm environment with hipcc and a GPU. +All tests require a ROCm environment with hipcc and a GPU. """ +import subprocess +import tempfile +from pathlib import Path + import pytest +from conftest import requires_rocm + kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB +# -- Helpers ---------------------------------------------------------------- + + +def _compile(source, name): + tmp = Path(tempfile.mkdtemp(prefix=f"kerneldb_{name}_")) + src = tmp / f"{name}.cpp" + exe = tmp / name + src.write_text(source) + r = subprocess.run(["hipcc", "-g", str(src), "-o", str(exe)], + capture_output=True, text=True) + if r.returncode != 0: + pytest.skip(f"hipcc compilation failed:\n{r.stderr}") + return str(exe) + + def _find_kernel(kernels, fragment): matches = [k for k in kernels if fragment in k] if not matches: - pytest.skip(f"No kernel containing {fragment!r} found in binary") + pytest.skip(f"No kernel containing {fragment!r} found") return matches[0] @@ -28,16 +46,98 @@ def _arg_by_name(arguments, name): return next((a for a in arguments if a.name == name), None) -# --------------------------------------------------------------------------- -# Basic argument metadata -# --------------------------------------------------------------------------- +# -- HIP sources ----------------------------------------------------------- + +_ARGS_SOURCE = r""" +#include + +__global__ void kernel_with_args(double* a, double* b, double* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] + b[idx]; + } +} + +int main() { return 0; } +""" + +_NESTED_SOURCE = r""" +#include + +struct Point3D { + float x; + float y; + float z; +}; + +struct BoundingBox { + Point3D min_pt; + Point3D max_pt; +}; + +__global__ void update_bounds(BoundingBox* boxes, int count) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < count) { + boxes[idx].min_pt.x = 0.0f; + } +} + +int main() { return 0; } +""" + +_TYPEDEF_SOURCE = r""" +#include + +using MyInt = int; +using MyFloat = float; +typedef double MyDouble; + +__global__ void typedef_kernel(MyInt a, MyFloat b, MyDouble c, int* d) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx == 0) { + d[0] = a + static_cast(b) + static_cast(c); + } +} + +int main() { return 0; } +""" + +_TEMPLATE_SOURCE = r""" +#include + +template +__global__ void scale_values(T* input, T* output, T factor, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + output[idx] = input[idx] * factor; + } +} + +template __global__ void scale_values(float*, float*, float, int); +template __global__ void scale_values(double*, double*, double, int); + +int main() { return 0; } +""" +# -- Cached binaries (compiled once per process) ---------------------------- -def test_kernel_arguments(arguments_binary): - """kernel_with_args must expose 4 named arguments with expected types/sizes.""" - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - args = kdb.get_kernel_arguments(kernel_name) +_cache = {} + + +def _get_binary(source, name): + if name not in _cache: + _cache[name] = _compile(source, name) + return _cache[name] + + +# -- Tests: basic arguments ------------------------------------------------ + + +@requires_rocm +def test_kernel_arguments(): + kdb = KernelDB(_get_binary(_ARGS_SOURCE, "args")) + name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + args = kdb.get_kernel_arguments(name) assert isinstance(args, list) and len(args) == 4 for arg in args: @@ -46,50 +146,43 @@ def test_kernel_arguments(arguments_binary): assert isinstance(arg.size, int) and arg.size > 0 assert isinstance(arg.alignment, int) and arg.alignment > 0 - names = [a.name for a in args] for expected in ("a", "b", "c", "n"): - assert expected in names, f"Expected argument {expected!r} in {names}" - - ptr_arg = _arg_by_name(args, "a") - assert ptr_arg.size == 8, f"Pointer size must be 8 bytes, got {ptr_arg.size}" + assert expected in [a.name for a in args] + assert _arg_by_name(args, "a").size == 8 # pointer = 8 bytes n_arg = _arg_by_name(args, "n") - assert n_arg.size == 4, f"int size must be 4 bytes, got {n_arg.size}" + assert n_arg.size == 4 assert "int" in n_arg.type_name.lower() -def test_kernel_wrapper_arguments(arguments_binary): - """Kernel.has_arguments() and .arguments must reflect the extracted args.""" - kdb = KernelDB(arguments_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "kernel_with_args") - kernel = kdb.get_kernel(kernel_name) +@requires_rocm +def test_kernel_wrapper_arguments(): + kdb = KernelDB(_get_binary(_ARGS_SOURCE, "args")) + name = _find_kernel(kdb.get_kernels(), "kernel_with_args") + kernel = kdb.get_kernel(name) assert kernel.has_arguments() is True assert isinstance(kernel.arguments, list) and kernel.arguments -# --------------------------------------------------------------------------- -# Nested struct member recursion -# --------------------------------------------------------------------------- +# -- Tests: nested structs ------------------------------------------------- -def test_nested_struct_arguments(nested_structs_binary): - """update_bounds args must include a BoundingBox arg exposing Point3D members.""" - kdb = KernelDB(nested_structs_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "update_bounds") - args = kdb.get_kernel_arguments(kernel_name) +@requires_rocm +def test_nested_struct_arguments(): + kdb = KernelDB(_get_binary(_NESTED_SOURCE, "nested")) + name = _find_kernel(kdb.get_kernels(), "update_bounds") + args = kdb.get_kernel_arguments(name) - # Every argument must have a members attribute for arg in args: assert hasattr(arg, "members") - # Helper: recursively search for a type by name - def find_type(arg_list, type_fragment): + def find_type(arg_list, fragment): for arg in arg_list: - if type_fragment in arg.type_name: + if fragment in arg.type_name: return arg if arg.members: - found = find_type(arg.members, type_fragment) + found = find_type(arg.members, fragment) if found: return found return None @@ -97,62 +190,49 @@ def find_type(arg_list, type_fragment): point = find_type(args, "Point3D") if point is None: pytest.skip("Point3D not found in argument tree (may be optimised away)") - assert len(point.members) == 3, f"Point3D should have 3 members, got {len(point.members)}" + assert len(point.members) == 3 -# --------------------------------------------------------------------------- -# Typedef / using-alias resolution -# --------------------------------------------------------------------------- +# -- Tests: typedef resolution --------------------------------------------- -def test_typedef_resolution(typedef_binary): - """typedef_kernel must have 4 args; typedefs preserved or resolved correctly.""" - kdb = KernelDB(typedef_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "typedef_kernel") +@requires_rocm +def test_typedef_resolution(): + kdb = KernelDB(_get_binary(_TYPEDEF_SOURCE, "typedef")) + name = _find_kernel(kdb.get_kernels(), "typedef_kernel") - args_raw = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=False) + args_raw = kdb.get_kernel_arguments(name, resolve_typedefs=False) assert len(args_raw) == 4 + raw_types = [a.type_name for a in args_raw] + assert any(n in ("MyInt", "MyFloat", "MyDouble") for n in raw_types) - type_names_raw = [a.type_name for a in args_raw] - assert any(n in ("MyInt", "MyFloat", "MyDouble") for n in type_names_raw), ( - f"Expected typedef names to be preserved, got: {type_names_raw}" - ) - - args_resolved = kdb.get_kernel_arguments(kernel_name, resolve_typedefs=True) - type_names_resolved = [a.type_name for a in args_resolved] + args_resolved = kdb.get_kernel_arguments(name, resolve_typedefs=True) + resolved_types = [a.type_name for a in args_resolved] for alias in ("MyInt", "MyFloat", "MyDouble"): - assert alias not in type_names_resolved, ( - f"Alias {alias!r} should have been resolved, got: {type_names_resolved}" - ) + assert alias not in resolved_types - primitive_types = {"int", "float", "double", "unsigned int"} - for type_name in type_names_resolved: - base = type_name.replace("*", "").strip() - assert any(p in base for p in primitive_types), ( - f"Resolved type {type_name!r} does not look like a C primitive" - ) + primitives = {"int", "float", "double", "unsigned int"} + for t in resolved_types: + base = t.replace("*", "").strip() + assert any(p in base for p in primitives), f"{t!r} doesn't look primitive" -# --------------------------------------------------------------------------- -# Template kernel instantiations -# --------------------------------------------------------------------------- +# -- Tests: template instantiations ---------------------------------------- -def test_template_kernel_arguments(template_binary): - """scale_values must be instantiated for at least float and double.""" - kdb = KernelDB(template_binary) +@requires_rocm +def test_template_kernel_arguments(): + kdb = KernelDB(_get_binary(_TEMPLATE_SOURCE, "template")) kernels = kdb.get_kernels() scale_kernels = [k for k in kernels if "scale_values" in k] - assert len(scale_kernels) >= 2, f"Expected ≥2 template instantiations, got {scale_kernels}" + assert len(scale_kernels) >= 2 - for type_fragment in ("float", "double"): - matching = [k for k in scale_kernels if type_fragment in k] + for type_frag in ("float", "double"): + matching = [k for k in scale_kernels if type_frag in k] if not matching: - pytest.skip(f"scale_values<{type_fragment}> not found in binary") + pytest.skip(f"scale_values<{type_frag}> not found") args = kdb.get_kernel_arguments(matching[0]) - assert len(args) == 4, f"Expected 4 args for scale_values<{type_fragment}>" + assert len(args) == 4 non_int = [a for a in args if a.name != "n"] - assert all(type_fragment in a.type_name.lower() for a in non_int), ( - f"Non-count args should be {type_fragment}: {[a.type_name for a in non_int]}" - ) + assert all(type_frag in a.type_name.lower() for a in non_int) diff --git a/tests/test_imports.py b/tests/test_imports.py index 4e9a284..f02dcd4 100644 --- a/tests/test_imports.py +++ b/tests/test_imports.py @@ -1,45 +1,22 @@ # SPDX-License-Identifier: MIT # Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. -""" -Tests that verify the public Python API of kernelDB can be imported correctly. - -These tests require the kernelDB C++ extension to be compiled and installed. -They are automatically skipped when the extension is not available. -""" +"""Tests that the public Python API of kernelDB can be imported correctly.""" import pytest -# Skip the entire module if the kernelDB C++ extension is not available -kerneldb = pytest.importorskip( - "kerneldb", - reason="kernelDB C++ extension not available – build with ROCm toolchain", -) - -PUBLIC_CLASSES = [ - "KernelDB", - "Kernel", - "Instruction", - "BasicBlock", - "CDNAKernel", - "ArchDescriptor", - "HsaAgent", - "KernelArgument", -] +kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") def test_package_metadata(): - """kerneldb must expose a non-empty __version__ and a valid __all__.""" assert isinstance(kerneldb.__version__, str) and kerneldb.__version__ assert isinstance(kerneldb.__all__, list) and kerneldb.__all__ - assert all(isinstance(n, str) for n in kerneldb.__all__) def test_public_api_accessible(): - """All expected public classes must be importable from the package.""" - for name in PUBLIC_CLASSES: - obj = getattr(kerneldb, name, None) - assert obj is not None, f"kerneldb.{name} not accessible" + for name in ("KernelDB", "Kernel", "Instruction", "BasicBlock", + "CDNAKernel", "ArchDescriptor", "HsaAgent", "KernelArgument"): + assert getattr(kerneldb, name, None) is not None, f"kerneldb.{name} not accessible" def test_KernelDB_and_Kernel_are_classes(): diff --git a/tests/test_kernels.py b/tests/test_kernels.py index 31b5757..40eca91 100644 --- a/tests/test_kernels.py +++ b/tests/test_kernels.py @@ -2,39 +2,78 @@ # Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. """ -Tests for KernelDB kernel-level operations: - - Kernel discovery (get_kernels) - - Source-line mapping (get_kernel_lines) - - Instruction extraction and regex filtering (get_instructions_for_line) - - Basic block extraction - - High-level Kernel wrapper properties - -All tests in this module require a ROCm environment with hipcc and a GPU. +Tests for KernelDB kernel-level operations: discovery, source-line mapping, +instruction extraction, filtering, basic blocks, and the Kernel wrapper. + +All tests require a ROCm environment with hipcc and a GPU. """ import re +import subprocess +import tempfile +from pathlib import Path + import pytest +from conftest import requires_rocm + kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB +# -- HIP source compiled once per module ------------------------------------ + +_HIP_SOURCE = r""" +#include + +__global__ void vector_add(float* a, float* b, float* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] + b[idx]; + } +} + +__global__ void vector_multiply(float* a, float* b, float* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] * b[idx]; + } +} + +int main() { return 0; } +""" + +_binary_path = None + + +def _get_binary(): + global _binary_path + if _binary_path is not None: + return _binary_path + tmp = Path(tempfile.mkdtemp(prefix="kerneldb_test_")) + src = tmp / "simple.cpp" + exe = tmp / "simple" + src.write_text(_HIP_SOURCE) + r = subprocess.run(["hipcc", "-g", str(src), "-o", str(exe)], + capture_output=True, text=True) + if r.returncode != 0: + pytest.skip(f"hipcc compilation failed:\n{r.stderr}") + _binary_path = str(exe) + return _binary_path + def _find_kernel(kernels, fragment): - """Return the first kernel name that contains *fragment*, or skip.""" matches = [k for k in kernels if fragment in k] if not matches: - pytest.skip(f"No kernel containing {fragment!r} found in binary") + pytest.skip(f"No kernel containing {fragment!r} found") return matches[0] -# --------------------------------------------------------------------------- -# Kernel discovery -# --------------------------------------------------------------------------- +# -- Tests ------------------------------------------------------------------ -def test_kernel_discovery(simple_binary): - """Binary must expose two named kernels: vector_add and vector_multiply.""" - kdb = KernelDB(simple_binary) +@requires_rocm +def test_kernel_discovery(): + kdb = KernelDB(_get_binary()) kernels = kdb.get_kernels() assert isinstance(kernels, list) @@ -44,99 +83,70 @@ def test_kernel_discovery(simple_binary): assert any("vector_multiply" in k for k in kernels) -# --------------------------------------------------------------------------- -# Source-line mapping -# --------------------------------------------------------------------------- - - -def test_kernel_lines(simple_binary): - """Each kernel must map to a non-empty list of positive source lines.""" - kdb = KernelDB(simple_binary) +@requires_rocm +def test_kernel_lines(): + kdb = KernelDB(_get_binary()) kernels = kdb.get_kernels() - for kernel_name in kernels: - lines = kdb.get_kernel_lines(kernel_name) - assert isinstance(lines, list) and lines, f"No lines for {kernel_name}" + for name in kernels: + lines = kdb.get_kernel_lines(name) + assert isinstance(lines, list) and lines, f"No lines for {name}" assert all(isinstance(ln, int) and ln >= 0 for ln in lines) - add_name = _find_kernel(kernels, "vector_add") - mul_name = _find_kernel(kernels, "vector_multiply") - assert set(kdb.get_kernel_lines(add_name)) != set(kdb.get_kernel_lines(mul_name)) - - -# --------------------------------------------------------------------------- -# Instruction extraction -# --------------------------------------------------------------------------- + add = _find_kernel(kernels, "vector_add") + mul = _find_kernel(kernels, "vector_multiply") + assert set(kdb.get_kernel_lines(add)) != set(kdb.get_kernel_lines(mul)) -def test_instructions(simple_binary): - """Instructions for each source line must have the expected attributes.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") +@requires_rocm +def test_instructions(): + kdb = KernelDB(_get_binary()) + name = _find_kernel(kdb.get_kernels(), "vector_add") - for line in kdb.get_kernel_lines(kernel_name): - instructions = kdb.get_instructions_for_line(kernel_name, line) - assert isinstance(instructions, list) - for inst in instructions: + for line in kdb.get_kernel_lines(name): + for inst in kdb.get_instructions_for_line(name, line): assert isinstance(inst.disassembly, str) and inst.disassembly assert isinstance(inst.line, int) and inst.line == line assert isinstance(inst.column, int) assert hasattr(inst, "file_name") -# --------------------------------------------------------------------------- -# Instruction filtering -# --------------------------------------------------------------------------- - - -def test_instruction_filtering(simple_binary): - """Regex filtering must return a subset that each matches the pattern.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") +@requires_rocm +def test_instruction_filtering(): + kdb = KernelDB(_get_binary()) + name = _find_kernel(kdb.get_kernels(), "vector_add") pattern = ".*(load|store).*" all_insts, mem_ops = [], [] - for line in kdb.get_kernel_lines(kernel_name): - all_insts.extend(kdb.get_instructions_for_line(kernel_name, line)) - mem_ops.extend(kdb.get_instructions_for_line(kernel_name, line, pattern)) + for line in kdb.get_kernel_lines(name): + all_insts.extend(kdb.get_instructions_for_line(name, line)) + mem_ops.extend(kdb.get_instructions_for_line(name, line, pattern)) assert len(mem_ops) > 0, "vector_add must have at least one load/store" assert len(mem_ops) <= len(all_insts) - assert all(re.search(pattern, inst.disassembly, re.IGNORECASE) for inst in mem_ops) - - # A pattern that matches nothing returns empty lists - for line in kdb.get_kernel_lines(kernel_name): - assert kdb.get_instructions_for_line(kernel_name, line, "NOMATCH__xyz") == [] - - -# --------------------------------------------------------------------------- -# Basic blocks and Kernel wrapper -# --------------------------------------------------------------------------- + assert all(re.search(pattern, i.disassembly, re.IGNORECASE) for i in mem_ops) + for line in kdb.get_kernel_lines(name): + assert kdb.get_instructions_for_line(name, line, "NOMATCH__xyz") == [] -def test_basic_blocks_and_kernel_wrapper(simple_binary): - """Kernel wrapper properties and basic-block extraction must work correctly.""" - kdb = KernelDB(simple_binary) - kernel_name = _find_kernel(kdb.get_kernels(), "vector_add") - kernel = kdb.get_kernel(kernel_name) - # basic blocks - blocks = kernel.get_basic_blocks() - assert isinstance(blocks, list) and blocks +@requires_rocm +def test_basic_blocks_and_kernel_wrapper(): + kdb = KernelDB(_get_binary()) + name = _find_kernel(kdb.get_kernels(), "vector_add") + kernel = kdb.get_kernel(name) - # wrapper properties - assert kernel.name == kernel_name + assert isinstance(kernel.get_basic_blocks(), list) and kernel.get_basic_blocks() + assert kernel.name == name assert isinstance(kernel.signature, str) and kernel.signature assert isinstance(kernel.lines, list) and all(ln >= 0 for ln in kernel.lines) assert isinstance(kernel.assembly, list) and all(isinstance(s, str) for s in kernel.assembly) assert isinstance(kernel.files, list) and all(isinstance(f, str) for f in kernel.files) - # instructions via wrapper first_line = kernel.lines[0] insts = kernel.get_instructions_for_line(first_line) assert isinstance(insts, list) and insts filtered = kernel.get_instructions_for_line(first_line, ".*load.*") assert len(filtered) <= len(insts) - # file name helper - assert isinstance(kdb.get_file_name(kernel_name, 0), str) + assert isinstance(kdb.get_file_name(name, 0), str) From 7d814120e26e6628a1ee0c6950a9e56ad0ae39a6 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 23 Mar 2026 00:16:04 -0700 Subject: [PATCH 6/7] Update copyright to 2026 and remove section separator comments Co-Authored-By: Claude Opus 4.6 --- tests/conftest.py | 2 +- tests/test_arguments.py | 21 +-------------------- tests/test_imports.py | 2 +- tests/test_kernels.py | 7 +------ 4 files changed, 4 insertions(+), 28 deletions(-) diff --git a/tests/conftest.py b/tests/conftest.py index 48e409a..361b094 100644 --- a/tests/conftest.py +++ b/tests/conftest.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: MIT -# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. import shutil diff --git a/tests/test_arguments.py b/tests/test_arguments.py index e8c0a6d..d476233 100644 --- a/tests/test_arguments.py +++ b/tests/test_arguments.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: MIT -# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. """ Tests for KernelDB kernel argument extraction: basic metadata, nested structs, @@ -20,9 +20,6 @@ KernelDB = kerneldb.KernelDB -# -- Helpers ---------------------------------------------------------------- - - def _compile(source, name): tmp = Path(tempfile.mkdtemp(prefix=f"kerneldb_{name}_")) src = tmp / f"{name}.cpp" @@ -46,8 +43,6 @@ def _arg_by_name(arguments, name): return next((a for a in arguments if a.name == name), None) -# -- HIP sources ----------------------------------------------------------- - _ARGS_SOURCE = r""" #include @@ -119,8 +114,6 @@ def _arg_by_name(arguments, name): int main() { return 0; } """ -# -- Cached binaries (compiled once per process) ---------------------------- - _cache = {} @@ -130,9 +123,6 @@ def _get_binary(source, name): return _cache[name] -# -- Tests: basic arguments ------------------------------------------------ - - @requires_rocm def test_kernel_arguments(): kdb = KernelDB(_get_binary(_ARGS_SOURCE, "args")) @@ -165,9 +155,6 @@ def test_kernel_wrapper_arguments(): assert isinstance(kernel.arguments, list) and kernel.arguments -# -- Tests: nested structs ------------------------------------------------- - - @requires_rocm def test_nested_struct_arguments(): kdb = KernelDB(_get_binary(_NESTED_SOURCE, "nested")) @@ -193,9 +180,6 @@ def find_type(arg_list, fragment): assert len(point.members) == 3 -# -- Tests: typedef resolution --------------------------------------------- - - @requires_rocm def test_typedef_resolution(): kdb = KernelDB(_get_binary(_TYPEDEF_SOURCE, "typedef")) @@ -217,9 +201,6 @@ def test_typedef_resolution(): assert any(p in base for p in primitives), f"{t!r} doesn't look primitive" -# -- Tests: template instantiations ---------------------------------------- - - @requires_rocm def test_template_kernel_arguments(): kdb = KernelDB(_get_binary(_TEMPLATE_SOURCE, "template")) diff --git a/tests/test_imports.py b/tests/test_imports.py index f02dcd4..0b57196 100644 --- a/tests/test_imports.py +++ b/tests/test_imports.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: MIT -# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. """Tests that the public Python API of kernelDB can be imported correctly.""" diff --git a/tests/test_kernels.py b/tests/test_kernels.py index 40eca91..2cb73c4 100644 --- a/tests/test_kernels.py +++ b/tests/test_kernels.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: MIT -# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. """ Tests for KernelDB kernel-level operations: discovery, source-line mapping, @@ -20,8 +20,6 @@ kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") KernelDB = kerneldb.KernelDB -# -- HIP source compiled once per module ------------------------------------ - _HIP_SOURCE = r""" #include @@ -68,9 +66,6 @@ def _find_kernel(kernels, fragment): return matches[0] -# -- Tests ------------------------------------------------------------------ - - @requires_rocm def test_kernel_discovery(): kdb = KernelDB(_get_binary()) From d5dce004fef19de784d29e12dc53f374ba61c1a9 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 23 Mar 2026 00:18:46 -0700 Subject: [PATCH 7/7] Fail instead of skip when Point3D is missing from debug info We compile the binary ourselves with -g, so this should never be absent. Co-Authored-By: Claude Opus 4.6 --- tests/test_arguments.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/test_arguments.py b/tests/test_arguments.py index d476233..3121734 100644 --- a/tests/test_arguments.py +++ b/tests/test_arguments.py @@ -175,8 +175,7 @@ def find_type(arg_list, fragment): return None point = find_type(args, "Point3D") - if point is None: - pytest.skip("Point3D not found in argument tree (may be optimised away)") + assert point is not None, "Point3D not found in argument tree" assert len(point.members) == 3