From 80d83deae3e006186ce1ad2542f9b77c09a6fc71 Mon Sep 17 00:00:00 2001 From: Bryce Adelstein Lelbach Date: Fri, 24 Jul 2026 18:02:25 +0000 Subject: [PATCH] Reorder PyHPC profiling notebooks --- docs/cscs.md | 6 +-- tutorials/pyhpc/.gitignore | 4 +- tutorials/pyhpc/README.md | 38 +++++++++---------- tutorials/pyhpc/brev/requirements.txt | 8 ++-- ...__power_iteration__cupy__asynchrony.ipynb} | 0 ...ipynb => 04__copy__kernel_authoring.ipynb} | 0 ...5__book_histogram__kernel_authoring.ipynb} | 0 .../{03__mpi4py.ipynb => 06__mpi4py.ipynb} | 0 ..._swe__intro.ipynb => 07__swe__intro.ipynb} | 14 +++---- ...{05__swe__jax.ipynb => 08__swe__jax.ipynb} | 2 +- ..._swe__pyomp.ipynb => 09__swe__pyomp.ipynb} | 2 +- ...nanobind.ipynb => 10__swe__nanobind.ipynb} | 8 ++-- ...st.ipynb => 11__swe__cppjit__thrust.ipynb} | 6 +-- ...nthesis.ipynb => 12__swe__synthesis.ipynb} | 6 +-- ...eration__cupy__asynchrony__SOLUTION.ipynb} | 0 ...4__copy__kernel_authoring__SOLUTION.ipynb} | 0 ...stogram__kernel_authoring__SOLUTION.ipynb} | 0 ...ION.ipynb => 08__swe__jax__SOLUTION.ipynb} | 2 +- ...N.ipynb => 09__swe__pyomp__SOLUTION.ipynb} | 2 +- ...pynb => 10__swe__nanobind__SOLUTION.ipynb} | 8 ++-- ...> 11__swe__cppjit__thrust__SOLUTION.ipynb} | 6 +-- ...i4py_jax_cppinterop_kernels__8_hours.ipynb | 26 ++++++------- tutorials/pyhpc/test/test_notebooks.py | 28 +++++++------- tutorials/pyhpc/test/test_packages.py | 2 +- 24 files changed, 84 insertions(+), 84 deletions(-) rename tutorials/pyhpc/notebooks/{10__power_iteration__cupy__asynchrony.ipynb => 03__power_iteration__cupy__asynchrony.ipynb} (100%) rename tutorials/pyhpc/notebooks/{11__copy__kernel_authoring.ipynb => 04__copy__kernel_authoring.ipynb} (100%) rename tutorials/pyhpc/notebooks/{12__book_histogram__kernel_authoring.ipynb => 05__book_histogram__kernel_authoring.ipynb} (100%) rename tutorials/pyhpc/notebooks/{03__mpi4py.ipynb => 06__mpi4py.ipynb} (100%) rename tutorials/pyhpc/notebooks/{04__swe__intro.ipynb => 07__swe__intro.ipynb} (97%) rename tutorials/pyhpc/notebooks/{05__swe__jax.ipynb => 08__swe__jax.ipynb} (99%) rename tutorials/pyhpc/notebooks/{06__swe__pyomp.ipynb => 09__swe__pyomp.ipynb} (99%) rename tutorials/pyhpc/notebooks/{07__swe__nanobind.ipynb => 10__swe__nanobind.ipynb} (98%) rename tutorials/pyhpc/notebooks/{08__swe__cppjit__thrust.ipynb => 11__swe__cppjit__thrust.ipynb} (98%) rename tutorials/pyhpc/notebooks/{09__swe__synthesis.ipynb => 12__swe__synthesis.ipynb} (98%) rename tutorials/pyhpc/notebooks/solutions/{10__power_iteration__cupy__asynchrony__SOLUTION.ipynb => 03__power_iteration__cupy__asynchrony__SOLUTION.ipynb} (100%) rename tutorials/pyhpc/notebooks/solutions/{11__copy__kernel_authoring__SOLUTION.ipynb => 04__copy__kernel_authoring__SOLUTION.ipynb} (100%) rename tutorials/pyhpc/notebooks/solutions/{12__book_histogram__kernel_authoring__SOLUTION.ipynb => 05__book_histogram__kernel_authoring__SOLUTION.ipynb} (100%) rename tutorials/pyhpc/notebooks/solutions/{05__swe__jax__SOLUTION.ipynb => 08__swe__jax__SOLUTION.ipynb} (99%) rename tutorials/pyhpc/notebooks/solutions/{06__swe__pyomp__SOLUTION.ipynb => 09__swe__pyomp__SOLUTION.ipynb} (99%) rename tutorials/pyhpc/notebooks/solutions/{07__swe__nanobind__SOLUTION.ipynb => 10__swe__nanobind__SOLUTION.ipynb} (98%) rename tutorials/pyhpc/notebooks/solutions/{08__swe__cppjit__thrust__SOLUTION.ipynb => 11__swe__cppjit__thrust__SOLUTION.ipynb} (99%) diff --git a/docs/cscs.md b/docs/cscs.md index dbb9a0c5..e271b2ff 100644 --- a/docs/cscs.md +++ b/docs/cscs.md @@ -183,10 +183,10 @@ CSCS_EDF="${CSCS_EDF}" \ The driver runs: - package smoke tests through the normal tutorial entrypoint -- the full notebook ladder, including `03__mpi4py` +- the full notebook ladder, including `06__mpi4py` - direct `nsys` and `ncu` command-line smoke checks -The PyHPC image builds `mpi4py` against MPICH. Notebook 03 runs local ranks with +The PyHPC image builds `mpi4py` against MPICH. Notebook 06 runs local ranks with `mpirun.mpich -launcher fork`, avoiding nested use of the host `srun` launcher. For debugging, the individual commands are below. @@ -214,7 +214,7 @@ Run the profiling notebooks: ```bash srun -A "${CSCS_ACCOUNT}" -p normal -t 01:00:00 -N1 -n1 \ --environment="${CSCS_EDF}" \ - env ACH_RUN_TESTS=1 ACH_TEST_ARGS="10 or 11 or 12" \ + env ACH_RUN_TESTS=1 ACH_TEST_ARGS="03 or 04 or 05" \ /accelerated-computing-hub/brev/entrypoint.bash base ``` diff --git a/tutorials/pyhpc/.gitignore b/tutorials/pyhpc/.gitignore index ea60cc9f..5071e2ed 100644 --- a/tutorials/pyhpc/.gitignore +++ b/tutorials/pyhpc/.gitignore @@ -1,12 +1,12 @@ # Artifacts generated when the tutorial is run. # -# SWE ladder (notebooks 04-09): +# SWE ladder (notebooks 07-12): notebooks/build/ notebooks/CMakeLists.txt notebooks/swe_step.cpython-*.so notebooks/timings.json -# Kernel / asynchrony notebooks (10-12): %%writefile scripts, Nsight +# Kernel / asynchrony notebooks (03-05): %%writefile scripts, Nsight # profiler reports, the downloaded text corpus, and rendered plots. # Patterns are applied at the notebooks/ root and under solutions/, since # executing the solution notebooks (e.g. in CI) writes the same artifacts there. diff --git a/tutorials/pyhpc/README.md b/tutorials/pyhpc/README.md index 8fff68cc..774350ee 100644 --- a/tutorials/pyhpc/README.md +++ b/tutorials/pyhpc/README.md @@ -17,7 +17,7 @@ Brev Launchables of this tutorial should use: ## Notebooks -Each exercise notebook that has a paired solution carries `# TODO:` cells with `...` placeholders; the solution fills them in. The intro/reference notebook (04), the synthesis notebook (09), and the mpi4py walkthrough (03) are complete as written and have no separate solution. +Each exercise notebook that has a paired solution carries `# TODO:` cells with `...` placeholders; the solution fills them in. The intro/reference notebook (07), the synthesis notebook (12), and the mpi4py walkthrough (06) are complete as written and have no separate solution. ### Fundamentals @@ -27,32 +27,32 @@ Each exercise notebook that has a paired solution carries `# TODO:` cells with ` | 01 | CuPy | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/01__cupy.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/01__cupy__SOLUTION.ipynb) | | 02 | Power Iteration - CuPy - Memory Spaces | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/02__power_iteration__cupy__memory_spaces.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/02__power_iteration__cupy__memory_spaces__SOLUTION.ipynb) | -### Distributed +### Kernels | # | Notebook | Link | Solution | |---|----------|------|----------| -| 03 | mpi4py | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/03__mpi4py.ipynb) | | - -### Programming models and interoperability +| 03 | Power Iteration - CuPy - Asynchrony | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb) | +| 04 | Copy - Kernel Authoring | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb) | +| 05 | Book Histogram - Kernel Authoring | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb) | -Notebooks 04-09 are the Shallow Water Equations "ladder": an intro plus NumPy baseline, four solvers that each re-implement the same timestep with a different tool, and a synthesis that reads the measured timings and compares them. The synthesis notebook collects the per-tool rows from `timings.json` (written by notebooks 04 to 08) and closes with a matched-precision GPU comparison of the JAX and CppJIT device paths. +### Distributed | # | Notebook | Link | Solution | |---|----------|------|----------| -| 04 | SWE - Intro | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/04__swe__intro.ipynb) | | -| 05 | SWE - JAX | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/05__swe__jax.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/05__swe__jax__SOLUTION.ipynb) | -| 06 | SWE - PyOMP | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/06__swe__pyomp.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/06__swe__pyomp__SOLUTION.ipynb) | -| 07 | SWE - nanobind | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/07__swe__nanobind.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/07__swe__nanobind__SOLUTION.ipynb) | -| 08 | SWE - CppJIT - Thrust | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/08__swe__cppjit__thrust.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/08__swe__cppjit__thrust__SOLUTION.ipynb) | -| 09 | SWE - Synthesis | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/09__swe__synthesis.ipynb) | | +| 06 | mpi4py | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/06__mpi4py.ipynb) | | -### Kernels +### Programming models and interoperability + +Notebooks 07-12 are the Shallow Water Equations "ladder": an intro plus NumPy baseline, four solvers that each re-implement the same timestep with a different tool, and a synthesis that reads the measured timings and compares them. The synthesis notebook collects the per-tool rows from `timings.json` (written by notebooks 07 to 11) and closes with a matched-precision GPU comparison of the JAX and CppJIT device paths. | # | Notebook | Link | Solution | |---|----------|------|----------| -| 10 | Power Iteration - CuPy - Asynchrony | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/10__power_iteration__cupy__asynchrony.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/10__power_iteration__cupy__asynchrony__SOLUTION.ipynb) | -| 11 | Copy - Kernel Authoring | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/11__copy__kernel_authoring.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/11__copy__kernel_authoring__SOLUTION.ipynb) | -| 12 | Book Histogram - Kernel Authoring | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/12__book_histogram__kernel_authoring.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/12__book_histogram__kernel_authoring__SOLUTION.ipynb) | +| 07 | SWE - Intro | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/07__swe__intro.ipynb) | | +| 08 | SWE - JAX | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/08__swe__jax.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/08__swe__jax__SOLUTION.ipynb) | +| 09 | SWE - PyOMP | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/09__swe__pyomp.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/09__swe__pyomp__SOLUTION.ipynb) | +| 10 | SWE - nanobind | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/10__swe__nanobind.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/10__swe__nanobind__SOLUTION.ipynb) | +| 11 | SWE - CppJIT - Thrust | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/11__swe__cppjit__thrust.ipynb) | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/solutions/11__swe__cppjit__thrust__SOLUTION.ipynb) | +| 12 | SWE - Synthesis | [![](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/NVIDIA/accelerated-computing-hub/blob/main/tutorials/pyhpc/notebooks/12__swe__synthesis.ipynb) | | ### Python/C++ interoperability @@ -64,7 +64,7 @@ A standalone comparison of several ways to call C and C++ from Python (ctypes, c ## The Shallow Water Equations problem -A 1D shallow-water bump pulse: a small mound of water at rest splits into two outgoing wave packets. Two conserved fields (`h`, `hu`) advance under a forward-Euler step. This PDE is small enough to read in full while exhibiting nonlinearity, and a fixed number of steps from the initial condition gives a result we can compare across tools. The full specification is in [`04__swe__intro.ipynb`](./notebooks/04__swe__intro.ipynb). +A 1D shallow-water bump pulse: a small mound of water at rest splits into two outgoing wave packets. Two conserved fields (`h`, `hu`) advance under a forward-Euler step. This PDE is small enough to read in full while exhibiting nonlinearity, and a fixed number of steps from the initial condition gives a result we can compare across tools. The full specification is in [`07__swe__intro.ipynb`](./notebooks/07__swe__intro.ipynb). ## Running @@ -76,8 +76,8 @@ Locally, with an NVIDIA GPU and a CUDA-13-capable driver: docker compose -f tutorials/pyhpc/brev/docker-compose.yml up ``` -Then open JupyterLab on port 8888. The notebooks are self-contained and can be run in any order, with one exception: the Shallow Water Equations ladder writes `timings.json` as you run notebooks 04 to 08, so run those before the synthesis notebook (09). +Then open JupyterLab on port 8888. The notebooks are self-contained and can be run in any order, with one exception: the Shallow Water Equations ladder writes `timings.json` as you run notebooks 07 to 11, so run those before the synthesis notebook (12). ## CppJIT toolchain -Notebook 08 uses CppJIT to automatically bind our Python runtime with CUDA C++, using the clang-repl C++ interpreter and [CppInterOp](https://github.com/compiler-research/CppInterOp). This is currently source built in the [Docker image](./brev/dockerfile) and not a standard `pip install`, so this will only run in the tutorial image. [CppJIT](https://github.com/compiler-research/CppJIT) is the successor to the [cppyy](https://cppyy.readthedocs.io/) automatic bindings tool, and no official CppJIT release is published on PyPI yet (beta release planned for end of summer 2026). +Notebook 11 uses CppJIT to automatically bind our Python runtime with CUDA C++, using the clang-repl C++ interpreter and [CppInterOp](https://github.com/compiler-research/CppInterOp). This is currently source built in the [Docker image](./brev/dockerfile) and not a standard `pip install`, so this will only run in the tutorial image. [CppJIT](https://github.com/compiler-research/CppJIT) is the successor to the [cppyy](https://cppyy.readthedocs.io/) automatic bindings tool, and no official CppJIT release is published on PyPI yet (beta release planned for end of summer 2026). diff --git a/tutorials/pyhpc/brev/requirements.txt b/tutorials/pyhpc/brev/requirements.txt index e8175078..3e403dec 100644 --- a/tutorials/pyhpc/brev/requirements.txt +++ b/tutorials/pyhpc/brev/requirements.txt @@ -12,22 +12,22 @@ jupyter-dark-detect == 0.1.0 nbformat == 5.10.4 nbconvert == 7.17.1 -# Programming models and interop (notebooks 04-09: the SWE ladder) +# Programming models and interop (notebooks 07-12: the SWE ladder) jax[cuda13] == 0.10.0 pyomp == 0.5.1 nanobind == 2.12.0 pybind11 == 3.0.4 -# CuPy and CUDA kernels (notebooks 00-02, 10-12: ported accelerated-python exercises) +# CuPy and CUDA kernels (notebooks 00-05: ported accelerated-python exercises) cupy-cuda13x == 13.6.0 numba-cuda[cu13] == 0.28.2 cuda-cccl[test-cu13] == 0.4.5 -# NVIDIA developer tools (notebooks 10-12: Nsight Systems / Nsight Compute profiling) +# NVIDIA developer tools (notebooks 03-05: Nsight Systems / Nsight Compute profiling) nvtx == 0.2.14 nsightful[notebook] @ git+https://github.com/brycelelbach/nsightful.git@454dfa5124ebea193f64302058b327f2034ccd8e -# Distributed computing (notebook 03: mpi4py) +# Distributed computing (notebook 06: mpi4py) mpi4py == 4.1.1 # Python/C++ Interop (notebook 13) diff --git a/tutorials/pyhpc/notebooks/10__power_iteration__cupy__asynchrony.ipynb b/tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/10__power_iteration__cupy__asynchrony.ipynb rename to tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb diff --git a/tutorials/pyhpc/notebooks/11__copy__kernel_authoring.ipynb b/tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/11__copy__kernel_authoring.ipynb rename to tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb diff --git a/tutorials/pyhpc/notebooks/12__book_histogram__kernel_authoring.ipynb b/tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/12__book_histogram__kernel_authoring.ipynb rename to tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb diff --git a/tutorials/pyhpc/notebooks/03__mpi4py.ipynb b/tutorials/pyhpc/notebooks/06__mpi4py.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/03__mpi4py.ipynb rename to tutorials/pyhpc/notebooks/06__mpi4py.ipynb diff --git a/tutorials/pyhpc/notebooks/04__swe__intro.ipynb b/tutorials/pyhpc/notebooks/07__swe__intro.ipynb similarity index 97% rename from tutorials/pyhpc/notebooks/04__swe__intro.ipynb rename to tutorials/pyhpc/notebooks/07__swe__intro.ipynb index 794aca10..c572c7eb 100644 --- a/tutorials/pyhpc/notebooks/04__swe__intro.ipynb +++ b/tutorials/pyhpc/notebooks/07__swe__intro.ipynb @@ -440,12 +440,12 @@ "\n", "| NB | Tool | Programming-model hook |\n", "|---|---|---|\n", - "| 05 | **JAX** | JIT compilation + whole-loop fusion offloaded to device |\n", - "| 06 | **PyOMP** | OpenMP parallelisation model in Python |\n", - "| 07 | **nanobind** | Handwritten Python bindings for compiled C++ kernels |\n", - "| 08 | **CppJIT** | JIT compiled C++ with automatic bindings |\n", + "| 08 | **JAX** | JIT compilation + whole-loop fusion offloaded to device |\n", + "| 09 | **PyOMP** | OpenMP parallelisation model in Python |\n", + "| 10 | **nanobind** | Handwritten Python bindings for compiled C++ kernels |\n", + "| 11 | **CppJIT** | JIT compiled C++ with automatic bindings |\n", "\n", - "We end with `09__swe__synthesis.ipynb`: a decision matrix and the achieved-throughput plot across the different notebooks in `timings.json`." + "We end with `12__swe__synthesis.ipynb`: a decision matrix and the achieved-throughput plot across the different notebooks in `timings.json`." ] }, { @@ -455,7 +455,7 @@ "source": [ "### 10. How to read these notebooks\n", "\n", - "Each numbered exercise (05-08) comes as two files:\n", + "Each numbered exercise (08-11) comes as two files:\n", "\n", "- `NN__topic__subtopic.ipynb`: the primary notebook we run live during the tutorial, with `# TODO:` cells whose body is a `...` placeholder.\n", "- `solutions/NN__topic__subtopic__SOLUTION.ipynb`: notebooks complete with solution, for self-paced review.\n", @@ -465,7 +465,7 @@ "- A **Quick Docs** sub-block precedes each TODO cell, listing the key APIs you need.\n", "- Optional **`# BONUS TASK:`** cells explore deeper solutions if you have time.\n", "\n", - "With the NumPy baseline in hand, let's continue with JAX: `05__swe__jax.ipynb`" + "With the NumPy baseline in hand, let's continue with JAX: `08__swe__jax.ipynb`" ] } ], diff --git a/tutorials/pyhpc/notebooks/05__swe__jax.ipynb b/tutorials/pyhpc/notebooks/08__swe__jax.ipynb similarity index 99% rename from tutorials/pyhpc/notebooks/05__swe__jax.ipynb rename to tutorials/pyhpc/notebooks/08__swe__jax.ipynb index 122b75e8..5d511d7f 100644 --- a/tutorials/pyhpc/notebooks/05__swe__jax.ipynb +++ b/tutorials/pyhpc/notebooks/08__swe__jax.ipynb @@ -85,7 +85,7 @@ "source": [ "### 2. The step as a pure function\n", "\n", - "Let's look at the same operation shape as notebook 04, expressed in JAX:\n", + "Let's look at the same operation shape as notebook 07, expressed in JAX:\n", "\n", "- **Read** — `h[:-1]`, `h[1:]` for interface states, same as NumPy-style slicing.\n", "- **Compute** — Rusanov formula unchanged: `F_h = 0.5*(huL+huR) - 0.5*a*(hR-hL)`.\n", diff --git a/tutorials/pyhpc/notebooks/06__swe__pyomp.ipynb b/tutorials/pyhpc/notebooks/09__swe__pyomp.ipynb similarity index 99% rename from tutorials/pyhpc/notebooks/06__swe__pyomp.ipynb rename to tutorials/pyhpc/notebooks/09__swe__pyomp.ipynb index 8f82cf2b..d7cd2692 100644 --- a/tutorials/pyhpc/notebooks/06__swe__pyomp.ipynb +++ b/tutorials/pyhpc/notebooks/09__swe__pyomp.ipynb @@ -400,7 +400,7 @@ "- Thread scaling works (section 5), and the kernel handles any `N`\n", " with no recompile (section 6).\n", "\n", - "Next: `07__swe__nanobind.ipynb` drops into C++. We compile a\n", + "Next: `10__swe__nanobind.ipynb` drops into C++. We compile a\n", "hand-written kernel and call it from Python through a binding library.\n" ] } diff --git a/tutorials/pyhpc/notebooks/07__swe__nanobind.ipynb b/tutorials/pyhpc/notebooks/10__swe__nanobind.ipynb similarity index 98% rename from tutorials/pyhpc/notebooks/07__swe__nanobind.ipynb rename to tutorials/pyhpc/notebooks/10__swe__nanobind.ipynb index 403de9d7..b3ed0838 100644 --- a/tutorials/pyhpc/notebooks/07__swe__nanobind.ipynb +++ b/tutorials/pyhpc/notebooks/10__swe__nanobind.ipynb @@ -100,7 +100,7 @@ "source": [ "### 2. The C++ source\n", "\n", - "Same **Read / Compute / Update** shape as NB 05, expressed in C++ via nanobind:\n", + "Same **Read / Compute / Update** shape as NB 08, expressed in C++ via nanobind:\n", "\n", "- **Read**. Scalar `double` loads from the `double*` pointer nanobind hands us, a zero-copy view of the NumPy buffer.\n", "- **Compute**. An inline `rusanov_face()` helper does the scalar arithmetic, all in registers.\n", @@ -296,12 +296,12 @@ "\n", "This notebook's cost is everything around the kernel, not the kernel itself:\n", "\n", - "- **The bindings track the C++ API.** `NB_MODULE`, `m.def`, and `nb::arg` are a separate layer that has to be hand-edited whenever the C++ source evolves: a new function or changed signature requires updating the bindings to match. The kernel is also no longer regular C++, it carries nanobind types like `nb::ndarray`. CppJIT (NB 05) needs neither, compiling standard C++/CUDA with no bindings library on the C++ side.\n", + "- **The bindings track the C++ API.** `NB_MODULE`, `m.def`, and `nb::arg` are a separate layer that has to be hand-edited whenever the C++ source evolves: a new function or changed signature requires updating the bindings to match. The kernel is also no longer regular C++, it carries nanobind types like `nb::ndarray`. CppJIT (NB 11) needs neither, compiling standard C++/CUDA with no bindings library on the C++ side.\n", "- **Build system.** Building the module requires working with CMake and a C++ toolchain.\n", "- **Recompile-on-edit cycle.** Editing the source leads to recompilation, restarting the kernel and re-importing the module.\n", "- **FFI boundary cost.** Each call from Python into C++ pays a small fixed overhead for argument conversion. At ~1000 calls per run it stays well under a millisecond, negligible against the kernel time.\n", "\n", - "The warm time above lands close to the NumPy baseline at this size, but runs slower, and does not scale. The single-threaded scalar loop has little to exploit when the arrays are cache-resident while NumPy's element-wise ops are SIMD-vectorised. This is not a shortcoming of nanobind, but rather our kernel implementation. To outperform NumPy, we require multi-threading (as in NB 03) or device parallelism (GPU).\n", + "The warm time above lands close to the NumPy baseline at this size, but runs slower, and does not scale. The single-threaded scalar loop has little to exploit when the arrays are cache-resident while NumPy's element-wise ops are SIMD-vectorised. This is not a shortcoming of nanobind, but rather our kernel implementation. To outperform NumPy, we require multi-threading (as in NB 09) or device parallelism (GPU).\n", "\n", "---\n", "\n", @@ -311,7 +311,7 @@ "- The cost is the glue around the kernel: a bindings layer to maintain as the C++ evolves, plus build and FFI overhead.\n", "- At this size the scalar kernel is on par/slower than NumPy, but does not scale.\n", "\n", - "Next: `08__swe__cppjit__thrust.ipynb` uses on-the-fly JIT-compiled C++, with no CMake, recompilation, or separate shared library. It provides automatic Python bindings and runs the whole solve on the GPU using the thrust library." + "Next: `11__swe__cppjit__thrust.ipynb` uses on-the-fly JIT-compiled C++, with no CMake, recompilation, or separate shared library. It provides automatic Python bindings and runs the whole solve on the GPU using the thrust library." ] } ], diff --git a/tutorials/pyhpc/notebooks/08__swe__cppjit__thrust.ipynb b/tutorials/pyhpc/notebooks/11__swe__cppjit__thrust.ipynb similarity index 98% rename from tutorials/pyhpc/notebooks/08__swe__cppjit__thrust.ipynb rename to tutorials/pyhpc/notebooks/11__swe__cppjit__thrust.ipynb index 5ade13c0..3a26e803 100644 --- a/tutorials/pyhpc/notebooks/08__swe__cppjit__thrust.ipynb +++ b/tutorials/pyhpc/notebooks/11__swe__cppjit__thrust.ipynb @@ -80,7 +80,7 @@ "dx = L / N\n", "DT = swe_core.fixed_dt(H0 + AMP, dx, cfl=CFL, g=G)\n", "\n", - "# Reference solution for the correctness gate (the NB 05 float64 solver).\n", + "# Reference solution for the correctness gate (the NB 08 float64 solver).\n", "def numpy_solve(n_cells, n_steps, cell_dx, step_dt):\n", " h, hu = swe_core.bump_ic(n_cells, L=L, h0=H0, amplitude=AMP, sigma=SIG)\n", " for _ in range(n_steps):\n", @@ -98,7 +98,7 @@ "source": [ "### 2. The SWE solve on the GPU with Thrust\n", "\n", - "The same Rusanov step as notebook 04, now as a `thrust::for_each` over the\n", + "The same Rusanov step as notebook 07, now as a `thrust::for_each` over the\n", "interior cells with the `thrust::device` policy. The fields live in CUDA managed\n", "memory and the whole time loop runs on the device, boundary update and buffer\n", "swap included, so we copy back and synchronise only once at the end. Keeping the\n", @@ -250,7 +250,7 @@ "- The GPU trails NumPy on a small grid but crosses over and pulls far ahead as\n", " `N` grows.\n", "\n", - "Next: `09__swe__synthesis.ipynb` collects every row from `timings.json` and compares\n", + "Next: `12__swe__synthesis.ipynb` collects every row from `timings.json` and compares\n", "the tools on runtime, throughput, and memory footprint." ] } diff --git a/tutorials/pyhpc/notebooks/09__swe__synthesis.ipynb b/tutorials/pyhpc/notebooks/12__swe__synthesis.ipynb similarity index 98% rename from tutorials/pyhpc/notebooks/09__swe__synthesis.ipynb rename to tutorials/pyhpc/notebooks/12__swe__synthesis.ipynb index 3619b5de..3ffcb2d5 100644 --- a/tutorials/pyhpc/notebooks/09__swe__synthesis.ipynb +++ b/tutorials/pyhpc/notebooks/12__swe__synthesis.ipynb @@ -32,7 +32,7 @@ "source": [ "### 1. Read the timings\n", "\n", - "We load every row from `timings.json` (written by NB 05-09). Each row\n", + "We load every row from `timings.json` (written by NB 07-11). Each row\n", "carries the warm wall-clock (`median_s`), the first-call cost where the\n", "tool JIT-compiles (`cold_s`), the hardware and dtype it ran on, and the\n", "max-diff against the float64 NumPy reference." @@ -53,13 +53,13 @@ "rows = swe_core.load_timings()\n", "by_tool = {r[\"tool\"]: r for r in rows}\n", "\n", - "# timings.json is written by notebooks 04 to 08\n", + "# timings.json is written by notebooks 07 to 11\n", "expected = [\"numpy\", \"jax\", \"pyomp\", \"nanobind\", \"cppjit_gpu_thrust\"]\n", "missing = [t for t in expected if t not in by_tool]\n", "if missing:\n", " raise SystemExit(\n", " \"timings.json has no rows for: \" + \", \".join(missing) + \". \"\n", - " \"Run notebooks 04 through 08 first, then re-run this notebook.\"\n", + " \"Run notebooks 07 through 11 first, then re-run this notebook.\"\n", " )\n", "\n", "hdr = (f'{\"stage\":<24}{\"tool\":>18}{\"hardware\":>9}{\"dtype\":>9}'\n", diff --git a/tutorials/pyhpc/notebooks/solutions/10__power_iteration__cupy__asynchrony__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/solutions/10__power_iteration__cupy__asynchrony__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb diff --git a/tutorials/pyhpc/notebooks/solutions/11__copy__kernel_authoring__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/solutions/11__copy__kernel_authoring__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb diff --git a/tutorials/pyhpc/notebooks/solutions/12__book_histogram__kernel_authoring__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb similarity index 100% rename from tutorials/pyhpc/notebooks/solutions/12__book_histogram__kernel_authoring__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb diff --git a/tutorials/pyhpc/notebooks/solutions/05__swe__jax__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/08__swe__jax__SOLUTION.ipynb similarity index 99% rename from tutorials/pyhpc/notebooks/solutions/05__swe__jax__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/08__swe__jax__SOLUTION.ipynb index e45d1c6e..6633288e 100644 --- a/tutorials/pyhpc/notebooks/solutions/05__swe__jax__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/08__swe__jax__SOLUTION.ipynb @@ -94,7 +94,7 @@ "source": [ "### 2. The step as a pure function\n", "\n", - "Let's look at the same operation shape as notebook 04, expressed in JAX:\n", + "Let's look at the same operation shape as notebook 07, expressed in JAX:\n", "\n", "- **Read** — `h[:-1]`, `h[1:]` for interface states, same as NumPy-style slicing.\n", "- **Compute** — Rusanov formula unchanged: `F_h = 0.5*(huL+huR) - 0.5*a*(hR-hL)`.\n", diff --git a/tutorials/pyhpc/notebooks/solutions/06__swe__pyomp__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/09__swe__pyomp__SOLUTION.ipynb similarity index 99% rename from tutorials/pyhpc/notebooks/solutions/06__swe__pyomp__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/09__swe__pyomp__SOLUTION.ipynb index 88c28399..f29ef61e 100644 --- a/tutorials/pyhpc/notebooks/solutions/06__swe__pyomp__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/09__swe__pyomp__SOLUTION.ipynb @@ -521,7 +521,7 @@ "- Thread scaling works (section 5), and the kernel handles any `N`\n", " with no recompile (section 6).\n", "\n", - "Next: `07__swe__nanobind.ipynb` drops into C++. We compile a\n", + "Next: `10__swe__nanobind.ipynb` drops into C++. We compile a\n", "hand-written kernel and call it from Python through a binding library.\n" ] } diff --git a/tutorials/pyhpc/notebooks/solutions/07__swe__nanobind__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/10__swe__nanobind__SOLUTION.ipynb similarity index 98% rename from tutorials/pyhpc/notebooks/solutions/07__swe__nanobind__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/10__swe__nanobind__SOLUTION.ipynb index 7b89ab6d..a2cbca3f 100644 --- a/tutorials/pyhpc/notebooks/solutions/07__swe__nanobind__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/10__swe__nanobind__SOLUTION.ipynb @@ -109,7 +109,7 @@ "source": [ "### 2. The C++ source\n", "\n", - "Same **Read / Compute / Update** shape as NB 05, expressed in C++ via nanobind:\n", + "Same **Read / Compute / Update** shape as NB 08, expressed in C++ via nanobind:\n", "\n", "- **Read**. Scalar `double` loads from the `double*` pointer nanobind hands us, a zero-copy view of the NumPy buffer.\n", "- **Compute**. An inline `rusanov_face()` helper does the scalar arithmetic, all in registers.\n", @@ -351,12 +351,12 @@ "\n", "This notebook's cost is everything around the kernel, not the kernel itself:\n", "\n", - "- **The bindings track the C++ API.** `NB_MODULE`, `m.def`, and `nb::arg` are a separate layer that has to be hand-edited whenever the C++ source evolves: a new function or changed signature requires updating the bindings to match. The kernel is also no longer regular C++, it carries nanobind types like `nb::ndarray`. CppJIT (NB 05) needs neither, compiling standard C++/CUDA with no bindings library on the C++ side.\n", + "- **The bindings track the C++ API.** `NB_MODULE`, `m.def`, and `nb::arg` are a separate layer that has to be hand-edited whenever the C++ source evolves: a new function or changed signature requires updating the bindings to match. The kernel is also no longer regular C++, it carries nanobind types like `nb::ndarray`. CppJIT (NB 11) needs neither, compiling standard C++/CUDA with no bindings library on the C++ side.\n", "- **Build system.** Building the module requires working with CMake and a C++ toolchain.\n", "- **Recompile-on-edit cycle.** Editing the source leads to recompilation, restarting the kernel and re-importing the module.\n", "- **FFI boundary cost.** Each call from Python into C++ pays a small fixed overhead for argument conversion. At ~1000 calls per run it stays well under a millisecond, negligible against the kernel time.\n", "\n", - "The warm time above lands close to the NumPy baseline at this size, but runs slower, and does not scale. The single-threaded scalar loop has little to exploit when the arrays are cache-resident while NumPy's element-wise ops are SIMD-vectorised. This is not a shortcoming of nanobind, but rather our kernel implementation. To outperform NumPy, we require multi-threading (as in NB 03) or device parallelism (GPU).\n", + "The warm time above lands close to the NumPy baseline at this size, but runs slower, and does not scale. The single-threaded scalar loop has little to exploit when the arrays are cache-resident while NumPy's element-wise ops are SIMD-vectorised. This is not a shortcoming of nanobind, but rather our kernel implementation. To outperform NumPy, we require multi-threading (as in NB 09) or device parallelism (GPU).\n", "\n", "---\n", "\n", @@ -366,7 +366,7 @@ "- The cost is the glue around the kernel: a bindings layer to maintain as the C++ evolves, plus build and FFI overhead.\n", "- At this size the scalar kernel is on par/slower than NumPy, but does not scale.\n", "\n", - "Next: `08__swe__cppjit__thrust.ipynb` uses on-the-fly JIT-compiled C++, with no CMake, recompilation, or separate shared library. It provides automatic Python bindings and runs the whole solve on the GPU using the thrust library." + "Next: `11__swe__cppjit__thrust.ipynb` uses on-the-fly JIT-compiled C++, with no CMake, recompilation, or separate shared library. It provides automatic Python bindings and runs the whole solve on the GPU using the thrust library." ] } ], diff --git a/tutorials/pyhpc/notebooks/solutions/08__swe__cppjit__thrust__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/11__swe__cppjit__thrust__SOLUTION.ipynb similarity index 99% rename from tutorials/pyhpc/notebooks/solutions/08__swe__cppjit__thrust__SOLUTION.ipynb rename to tutorials/pyhpc/notebooks/solutions/11__swe__cppjit__thrust__SOLUTION.ipynb index 183e5fb0..11f55de6 100644 --- a/tutorials/pyhpc/notebooks/solutions/08__swe__cppjit__thrust__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/11__swe__cppjit__thrust__SOLUTION.ipynb @@ -87,7 +87,7 @@ "dx = L / N\n", "DT = swe_core.fixed_dt(H0 + AMP, dx, cfl=CFL, g=G)\n", "\n", - "# Reference solution for the correctness gate (the NB 05 float64 solver).\n", + "# Reference solution for the correctness gate (the NB 08 float64 solver).\n", "def numpy_solve(n_cells, n_steps, cell_dx, step_dt):\n", " h, hu = swe_core.bump_ic(n_cells, L=L, h0=H0, amplitude=AMP, sigma=SIG)\n", " for _ in range(n_steps):\n", @@ -105,7 +105,7 @@ "source": [ "### 2. The SWE solve on the GPU with Thrust\n", "\n", - "The same Rusanov step as notebook 04, now as a `thrust::for_each` over the\n", + "The same Rusanov step as notebook 07, now as a `thrust::for_each` over the\n", "interior cells with the `thrust::device` policy. The fields live in CUDA managed\n", "memory and the whole time loop runs on the device, boundary update and buffer\n", "swap included, so we copy back and synchronise only once at the end. Keeping the\n", @@ -325,7 +325,7 @@ "- The GPU trails NumPy on a small grid but crosses over and pulls far ahead as\n", " `N` grows.\n", "\n", - "Next: `09__swe__synthesis.ipynb` collects every row from `timings.json` and compares\n", + "Next: `12__swe__synthesis.ipynb` collects every row from `timings.json` and compares\n", "the tools on runtime, throughput, and memory footprint." ] } diff --git a/tutorials/pyhpc/notebooks/syllabi/pyhpc__cupy_mpi4py_jax_cppinterop_kernels__8_hours.ipynb b/tutorials/pyhpc/notebooks/syllabi/pyhpc__cupy_mpi4py_jax_cppinterop_kernels__8_hours.ipynb index 7c697ddc..facb96bc 100644 --- a/tutorials/pyhpc/notebooks/syllabi/pyhpc__cupy_mpi4py_jax_cppinterop_kernels__8_hours.ipynb +++ b/tutorials/pyhpc/notebooks/syllabi/pyhpc__cupy_mpi4py_jax_cppinterop_kernels__8_hours.ipynb @@ -14,16 +14,16 @@ "| 00 | [NumPy](../00__numpy.ipynb) | [Click](../solutions/00__numpy__SOLUTION.ipynb) | NumPy |\n", "| 01 | [CuPy](../01__cupy.ipynb) | [Click](../solutions/01__cupy__SOLUTION.ipynb) | CuPy |\n", "| 02 | [Power Iteration - CuPy - Memory Spaces](../02__power_iteration__cupy__memory_spaces.ipynb) | [Click](../solutions/02__power_iteration__cupy__memory_spaces__SOLUTION.ipynb) | CuPy |\n", - "| 03 | [mpi4py](../03__mpi4py.ipynb) | - | mpi4py |\n", - "| 04 | [SWE - Intro](../04__swe__intro.ipynb) | - | NumPy |\n", - "| 05 | [SWE - JAX](../05__swe__jax.ipynb) | [Click](../solutions/05__swe__jax__SOLUTION.ipynb) | JAX |\n", - "| 06 | [SWE - PyOMP](../06__swe__pyomp.ipynb) | [Click](../solutions/06__swe__pyomp__SOLUTION.ipynb) | PyOMP |\n", - "| 07 | [SWE - nanobind](../07__swe__nanobind.ipynb) | [Click](../solutions/07__swe__nanobind__SOLUTION.ipynb) | nanobind |\n", - "| 08 | [SWE - CppJIT - Thrust](../08__swe__cppjit__thrust.ipynb) | [Click](../solutions/08__swe__cppjit__thrust__SOLUTION.ipynb) | CppJIT, CUDA |\n", - "| 09 | [SWE - Synthesis](../09__swe__synthesis.ipynb) | - | - |\n", - "| 10 | [Power Iteration - CuPy - Asynchrony](../10__power_iteration__cupy__asynchrony.ipynb) | [Click](../solutions/10__power_iteration__cupy__asynchrony__SOLUTION.ipynb) | CuPy, Nsight Systems |\n", - "| 11 | [Copy - Kernel Authoring](../11__copy__kernel_authoring.ipynb) | [Click](../solutions/11__copy__kernel_authoring__SOLUTION.ipynb) | Numba CUDA, Nsight Compute |\n", - "| 12 | [Book Histogram - Kernel Authoring](../12__book_histogram__kernel_authoring.ipynb) | [Click](../solutions/12__book_histogram__kernel_authoring__SOLUTION.ipynb) | Numba CUDA, Nsight Compute |\n", + "| 03 | [Power Iteration - CuPy - Asynchrony](../03__power_iteration__cupy__asynchrony.ipynb) | [Click](../solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb) | CuPy, Nsight Systems |\n", + "| 04 | [Copy - Kernel Authoring](../04__copy__kernel_authoring.ipynb) | [Click](../solutions/04__copy__kernel_authoring__SOLUTION.ipynb) | Numba CUDA, Nsight Compute |\n", + "| 05 | [Book Histogram - Kernel Authoring](../05__book_histogram__kernel_authoring.ipynb) | [Click](../solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb) | Numba CUDA, Nsight Compute |\n", + "| 06 | [mpi4py](../06__mpi4py.ipynb) | - | mpi4py |\n", + "| 07 | [SWE - Intro](../07__swe__intro.ipynb) | - | NumPy |\n", + "| 08 | [SWE - JAX](../08__swe__jax.ipynb) | [Click](../solutions/08__swe__jax__SOLUTION.ipynb) | JAX |\n", + "| 09 | [SWE - PyOMP](../09__swe__pyomp.ipynb) | [Click](../solutions/09__swe__pyomp__SOLUTION.ipynb) | PyOMP |\n", + "| 10 | [SWE - nanobind](../10__swe__nanobind.ipynb) | [Click](../solutions/10__swe__nanobind__SOLUTION.ipynb) | nanobind |\n", + "| 11 | [SWE - CppJIT - Thrust](../11__swe__cppjit__thrust.ipynb) | [Click](../solutions/11__swe__cppjit__thrust__SOLUTION.ipynb) | CppJIT, CUDA |\n", + "| 12 | [SWE - Synthesis](../12__swe__synthesis.ipynb) | - | - |\n", "| 13 | [C++ Interop](../13__cpp_interop.ipynb) | [Click](../solutions/13__cpp_interop__SOLUTION.ipynb) | ctypes, cffi, nanobind, CppJIT |\n", "\n", "### Materials\n", @@ -34,16 +34,16 @@ "\n", "### Introduction\n", "\n", - "This tutorial tours the high-performance Python landscape, from array programming to custom GPU kernels and multi-process parallelism. We start with the NumPy and CuPy array model, move data deliberately between host and device, and distribute work across processes with mpi4py.\n", + "This tutorial tours the high-performance Python landscape, from array programming to custom GPU kernels and multi-process parallelism. We start with the NumPy and CuPy array model, move data deliberately between host and device, profile asynchronous execution, and author custom CUDA kernels.\n", "\n", - "We then solve one problem -- the 1D Shallow Water Equations -- five different ways, trading a NumPy baseline for JAX, PyOMP, nanobind, and CppJIT, and measure each against the others. We close by authoring and profiling our own CUDA kernels with Numba CUDA, Nsight Systems, and Nsight Compute.\n", + "We then distribute work across processes with mpi4py and solve one problem -- the 1D Shallow Water Equations -- five different ways, trading a NumPy baseline for JAX, PyOMP, nanobind, and CppJIT, and measure each against the others.\n", "\n", "In this tutorial we will cover:\n", "- NumPy fundamentals and the `ndarray` memory model.\n", "- Accelerating array workflows on the GPU with CuPy, including memory spaces and asynchrony.\n", + "- Authoring and profiling custom CUDA kernels with Numba CUDA, Nsight Systems, and Nsight Compute.\n", "- Distributed computing with mpi4py.\n", "- Programming models and Python/C++ interoperability: JAX, PyOMP, nanobind, and CppJIT.\n", - "- Authoring and profiling custom CUDA kernels with Numba CUDA, Nsight Systems, and Nsight Compute.\n", "\n", "Attendees are expected to have general knowledge of HPC and Python programming. Some familiarity with NumPy and C/C++ is helpful but not required." ] diff --git a/tutorials/pyhpc/test/test_notebooks.py b/tutorials/pyhpc/test/test_notebooks.py index 78da88c2..1dc8ab48 100644 --- a/tutorials/pyhpc/test/test_notebooks.py +++ b/tutorials/pyhpc/test/test_notebooks.py @@ -3,12 +3,12 @@ The notebooks run as an ordered ladder (00 to 13). For each rung we prefer the filled-in solution notebook when one exists, and otherwise fall back to -the exercise notebook (the intro/reference notebook 04, the synthesis -notebook 09, and the mpi4py walkthrough 03 have no separate solution and are +the exercise notebook (the intro/reference notebook 07, the synthesis +notebook 12, and the mpi4py walkthrough 06 have no separate solution and are complete as written). -Ordering matters for the SWE sub-ladder: notebooks 04 to 08 each append a -row to timings.json and 09 reads them, so 04-08 must run before 09. pytest +Ordering matters for the SWE sub-ladder: notebooks 07 to 11 each append a +row to timings.json and 12 reads them, so 07-11 must run before 12. pytest executes the parametrized cases in list order, so listing the rungs 00..13 in order is sufficient. """ @@ -39,16 +39,16 @@ def _runnable_notebook(stem): "00__numpy", "01__cupy", "02__power_iteration__cupy__memory_spaces", - "03__mpi4py", - "04__swe__intro", - "05__swe__jax", - "06__swe__pyomp", - "07__swe__nanobind", - "08__swe__cppjit__thrust", - "09__swe__synthesis", - "10__power_iteration__cupy__asynchrony", - "11__copy__kernel_authoring", - "12__book_histogram__kernel_authoring", + "03__power_iteration__cupy__asynchrony", + "04__copy__kernel_authoring", + "05__book_histogram__kernel_authoring", + "06__mpi4py", + "07__swe__intro", + "08__swe__jax", + "09__swe__pyomp", + "10__swe__nanobind", + "11__swe__cppjit__thrust", + "12__swe__synthesis", "13__cpp_interop", ] ladder = [(stem, _runnable_notebook(stem)) for stem in LADDER_STEMS] diff --git a/tutorials/pyhpc/test/test_packages.py b/tutorials/pyhpc/test/test_packages.py index f4584526..c5b1e1d9 100644 --- a/tutorials/pyhpc/test/test_packages.py +++ b/tutorials/pyhpc/test/test_packages.py @@ -45,7 +45,7 @@ def add_one(x): def test_cuda_cooperative(): - """cuda.cooperative builds a block-load algorithm (used by notebook 12).""" + """cuda.cooperative builds a block-load algorithm (used by notebook 05).""" import cuda.coop as coop block_load = coop.block.load(np.uint8, 128, 4, "striped")