Build the kernels with cunumpy's KernelPackage

CHANGELOG.md

@@ -117,6 +117,9 @@ format follows [Keep a Changelog](https://keepachangelog.com/en/1.1.0/).
 
 ### Changed
 
+- The kernels are built with cunumpy's `KernelPackage`
+  (`simgeometry.kernels.KERNEL_PACKAGE`): `kernels/_build.py`, `kernels/_metal.py` and
+  `kernels/compile.py` are gone, and the build cache is still `SIMGEOMETRY_KERNEL_CACHE`.
 - `__version__` is read from the installed package metadata (`importlib.metadata`), not hardcoded,
   so it can no longer disagree with `pyproject.toml` (the 0.1.1 wheel reported 0.1.0).
 - One out-of-range rule for owners: per axis, positions below the lower bound belong to the

src/simgeometry/kernels/__init__.py

@@ -1 +1,23 @@
-"""Compiled containment kernels for geometric objects."""
+"""simgeometry's compiled kernels, one folder per kernel.
+
+Every kernel folder builds its kernel with :data:`KERNEL_PACKAGE`::
+
+    kernel = KERNEL_PACKAGE.from_folder(__name__)
+
+with the defaults of :class:`cunumpy.kernels.KernelPackage`: the Pyccel version is
+compiled on its first host call and kept under ``SIMGEOMETRY_KERNEL_CACHE`` (default
+``~/.cache/simgeometry/kernels``; ``SIMGEOMETRY_OPENMP=0`` builds serial), the CUDA version runs
+for CuPy arguments and the host version for NumPy ones, one CUDA thread per entry of
+the last axis of the first array. Importing a kernel needs neither Pyccel nor CuPy.
+
+The Metal versions (``<name>/<name>_metal.py``) wrap a kernel built by
+``KERNEL_PACKAGE.metal_kernel`` in a function with the arguments of the Pyccel
+version. The Apple GPU has no float64, so they compute in float32: expect a relative
+error of about 1e-5 against the CPU versions. Import one only where
+:func:`cunumpy.kernels.metal_available` is true.
+"""
+
+import cunumpy as xp
+
+#: The build cache and launch options of simgeometry's kernels.
+KERNEL_PACKAGE = xp.kernels.KernelPackage("simgeometry")

src/simgeometry/kernels/absorb_polygons_2d/__init__.py

@@ -6,15 +6,11 @@
 version runs one thread per entry of ``rows`` (the third argument).
 """
 
-import cunumpy as xp
-
-from simgeometry.kernels._build import KERNEL_OPTIONS
+from simgeometry.kernels import KERNEL_PACKAGE
 
 
 def _row_count(args: tuple) -> int:
     return int(args[2].shape[0])
 
 
-kernel = xp.kernels.Kernel.from_folder(
-    __name__, **{**KERNEL_OPTIONS, "n_threads_from": _row_count}
-)
+kernel = KERNEL_PACKAGE.from_folder(__name__, n_threads_from=_row_count)

src/simgeometry/kernels/contains_ball/__init__.py

@@ -5,8 +5,6 @@
 (see ``cunumpy.kernels.set_host_kernel_implementation`` to choose).
 """
 
-import cunumpy as xp
+from simgeometry.kernels import KERNEL_PACKAGE
 
-from simgeometry.kernels._build import KERNEL_OPTIONS
-
-kernel = xp.kernels.Kernel.from_folder(__name__, **KERNEL_OPTIONS)
+kernel = KERNEL_PACKAGE.from_folder(__name__)

src/simgeometry/kernels/contains_ball/contains_ball_metal.py

@@ -1,7 +1,7 @@
 """The Metal version (Apple-silicon GPU, float32) of the ``contains_ball`` kernel.
 
 ``contains_ball_pyccel.py`` next to this file is the CPU version with the same
-arguments. One thread per marker (``positions`` is ``(ndim, N)``); see :mod:`simgeometry.kernels._metal`. The test is
+arguments. One thread per marker (``positions`` is ``(ndim, N)``); see :mod:`simgeometry.kernels`. The test is
 done in float32, so a marker within about 1e-7 (relative) of the boundary may be
 classified differently from the CPU version.
 """
@@ -11,7 +11,7 @@
 import numpy as np
 
 from simgeometry._types import Array
-from simgeometry.kernels._metal import make_kernel
+from simgeometry.kernels import KERNEL_PACKAGE
 
 _SOURCE = """
     uint marker = thread_position_in_grid.x;
@@ -28,7 +28,7 @@
     inside_flag[marker] = inside ? 1 : 0;
 """
 
-_KERNEL = make_kernel(
+_KERNEL = KERNEL_PACKAGE.metal_kernel(
     "simgeometry_contains_ball",
     _SOURCE,
     inputs=["positions", "midpoint", "params"],

src/simgeometry/kernels/contains_box/__init__.py

@@ -5,8 +5,6 @@
 (see ``cunumpy.kernels.set_host_kernel_implementation`` to choose).
 """
 
-import cunumpy as xp
+from simgeometry.kernels import KERNEL_PACKAGE
 
-from simgeometry.kernels._build import KERNEL_OPTIONS
-
-kernel = xp.kernels.Kernel.from_folder(__name__, **KERNEL_OPTIONS)
+kernel = KERNEL_PACKAGE.from_folder(__name__)

src/simgeometry/kernels/contains_box/contains_box_metal.py

@@ -1,7 +1,7 @@
 """The Metal version (Apple-silicon GPU, float32) of the ``contains_box`` kernel.
 
 ``contains_box_pyccel.py`` next to this file is the CPU version with the same
-arguments. One thread per marker (``positions`` is ``(ndim, N)``); see :mod:`simgeometry.kernels._metal`. The test is
+arguments. One thread per marker (``positions`` is ``(ndim, N)``); see :mod:`simgeometry.kernels`. The test is
 done in float32, so a marker within about 1e-7 (relative) of the boundary may be
 classified differently from the CPU version.
 """
@@ -11,7 +11,7 @@
 import numpy as np
 
 from simgeometry._types import Array
-from simgeometry.kernels._metal import make_kernel
+from simgeometry.kernels import KERNEL_PACKAGE
 
 _SOURCE = """
     uint marker = thread_position_in_grid.x;
@@ -27,7 +27,7 @@
     inside_flag[marker] = inside ? 1 : 0;
 """
 
-_KERNEL = make_kernel(
+_KERNEL = KERNEL_PACKAGE.metal_kernel(
     "simgeometry_contains_box",
     _SOURCE,
     inputs=["positions", "lower", "upper", "flags"],

src/simgeometry/kernels/contains_polygon_2d/__init__.py

@@ -5,8 +5,6 @@
 (see ``cunumpy.kernels.set_host_kernel_implementation`` to choose).
 """
 
-import cunumpy as xp
+from simgeometry.kernels import KERNEL_PACKAGE
 
-from simgeometry.kernels._build import KERNEL_OPTIONS
-
-kernel = xp.kernels.Kernel.from_folder(__name__, **KERNEL_OPTIONS)
+kernel = KERNEL_PACKAGE.from_folder(__name__)

src/simgeometry/kernels/contains_polygon_2d/contains_polygon_2d_metal.py

@@ -1,7 +1,7 @@
 """The Metal version (Apple-silicon GPU, float32) of the ``contains_polygon_2d`` kernel.
 
 ``contains_polygon_2d_pyccel.py`` next to this file is the CPU version with the same
-arguments. One thread per marker (``positions`` is ``(2, N)``); see :mod:`simgeometry.kernels._metal`. The test is
+arguments. One thread per marker (``positions`` is ``(2, N)``); see :mod:`simgeometry.kernels`. The test is
 done in float32, so a marker within about 1e-7 (relative) of the boundary may be
 classified differently from the CPU version.
 """
@@ -11,7 +11,7 @@
 import numpy as np
 
 from simgeometry._types import Array
-from simgeometry.kernels._metal import make_kernel
+from simgeometry.kernels import KERNEL_PACKAGE
 
 _SOURCE = """
     uint marker = thread_position_in_grid.x;
@@ -36,7 +36,7 @@
     inside_flag[marker] = inside ? 1 : 0;
 """
 
-_KERNEL = make_kernel(
+_KERNEL = KERNEL_PACKAGE.metal_kernel(
     "simgeometry_contains_polygon_2d",
     _SOURCE,
     inputs=["positions", "vertices", "flags"],