Address PyHPC profiler notebook review

tutorials/pyhpc/brev/dockerfile

@@ -163,7 +163,6 @@ RUN set -eux; \
     rm -rf /var/lib/apt/lists/*
 
 # Install profiler-backed kernels for in-place notebook cell profiling.
-# The JupyterLab Nsight extension remains installed for the streamer tiles.
 RUN nsightful-ncu install --sys-prefix '--profiler-args=--set full --clock-control none' \
  && nsightful-nsys install --sys-prefix '--profiler-args=--trace=cuda,nvtx,osrt'
 

tutorials/pyhpc/brev/requirements.txt

@@ -24,7 +24,7 @@ cuda-cccl[test-cu13] == 0.4.5
 
 # NVIDIA developer tools (notebooks 03-05: Nsight Systems / Nsight Compute profiling)
 nvtx == 0.2.14
-nsightful[notebook] @ git+https://github.com/brycelelbach/nsightful.git@fa9ee4d81441ac62379c5306f2f2d8b0894d06ec
+nsightful[notebook] @ git+https://github.com/brycelelbach/nsightful.git@a41989403430168e02ac3cfdc4060bca4ebb8040
 
 # Distributed computing (notebooks 06 and 13: mpi4py)
 mpi4py == 4.1.1

tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb

@@ -25,7 +25,7 @@
     "\n",
     "We will revisit the Power Iteration algorithm. Our goal is to take a standard implementation, profile it to identify bottlenecks caused by implicit synchronization, and then optimize it using CUDA streams and asynchronous memory transfers.\n",
     "\n",
-    "First, we need to ensure the Nsight Systems profiler (nsys), Nsightful, and NVTX are installed and available."
+    "First, we need to ensure NVIDIA's developer tools are installed and available and do all of our imports."
    ]
   },
   {
@@ -155,12 +155,31 @@
     "      y = A_gpu @ x\n",
     "      x = y / cp.linalg.norm(y)\n",
     "\n",
-    "  return cp.asnumpy((x.T @ (A_gpu @ x)) / (x.T @ x))\n",
+    "  return cp.asnumpy((x.T @ (A_gpu @ x)) / (x.T @ x))"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "id": "3bcee9e1",
+   "metadata": {},
+   "source": [
+    "Now let's make sure it works:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "id": "a3c12c49",
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "lam_est_baseline = estimate_device_baseline(A_device)\n",
+    "\n",
+    "assert isinstance(lam_est_baseline, (np.ndarray, np.generic)), \"Must return a NumPy array or NumPy scalar\"\n",
+    "np.testing.assert_allclose(lam_est_baseline, 1, atol=1e-4)\n",
     "\n",
-    "estimate_device_baseline(\n",
-    "  A_device,\n",
-    "  cfg=PowerIterationConfig(max_steps=1, check_frequency=1, progress=False),\n",
-    ")"
+    "print()\n",
+    "print(\"Dominant Eigenvalue:\", lam_est_baseline)"
    ]
   },
   {
@@ -172,7 +191,7 @@
    "source": [
     "### 4. Profiling the Baseline\n",
     "\n",
-    "This notebook uses the **Python 3 (Nsight Systems)** kernel so we can profile the baseline in place. The `%%nsys` magic collects only this cell and displays the resulting timeline without restarting the kernel."
+    "The `%%nsys` cell magic profiles the code in its cell with Nsight Systems and saves the native report to the path given by `-o`."
    ]
   },
   {
@@ -185,8 +204,7 @@
    "outputs": [],
    "source": [
     "%%nsys -o power_iteration__baseline.nsys-rep\n",
-    "lam_est_baseline = estimate_device_baseline(A_device)\n",
-    "np.testing.assert_allclose(lam_est_baseline, 1, atol=1e-4)"
+    "lam_est_baseline = estimate_device_baseline(A_device)"
    ]
   },
   {
@@ -196,21 +214,7 @@
     "id": "IlGIAIEPe3SV"
    },
    "source": [
-    "The timeline is displayed below the profiled cell. Explore what's going on in the program.\n",
-    "\n",
-    "**EXTRA CREDIT:** Download the Nsight Systems GUI and open the report in it to see even more information."
-   ]
-  },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "id": "c6028322",
-   "metadata": {
-    "id": "s6VVOnGQR3Ph"
-   },
-   "outputs": [],
-   "source": [
-    "# The native report is saved as power_iteration__baseline.nsys-rep."
+    "Explore the profile in Perfetto by clicking the button below the profiled cell. For richer analysis, click the **+** in the JupyterLab tab bar and open Nsight Systems. You can also install the Nsight Systems GUI on your local system, download the `.nsys-rep` report, and open it there."
    ]
   },
   {
@@ -285,7 +289,12 @@
    "outputs": [],
    "source": [
     "lam_est_async = estimate_device_async(A_device)\n",
-    "np.testing.assert_allclose(lam_est_async, 1, atol=1e-4)"
+    "\n",
+    "assert isinstance(lam_est_async, (np.ndarray, np.generic)), \"Must return a NumPy array or NumPy scalar\"\n",
+    "np.testing.assert_allclose(lam_est_async, 1, atol=1e-4)\n",
+    "\n",
+    "print()\n",
+    "print(\"Dominant Eigenvalue:\", lam_est_baseline)"
    ]
   },
   {
@@ -343,8 +352,7 @@
    "outputs": [],
    "source": [
     "%%nsys -o power_iteration__async.nsys-rep\n",
-    "lam_est_async = estimate_device_async(A_device)\n",
-    "np.testing.assert_allclose(lam_est_async, 1, atol=1e-4)"
+    "lam_est_async = estimate_device_async(A_device)"
    ]
   },
   {
@@ -357,18 +365,6 @@
     "Finally, let's look at the profile in Perfetto and confirm we've gotten rid of the idling."
    ]
   },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "id": "55e6d51f",
-   "metadata": {
-    "id": "mWXBvi-hFGhU"
-   },
-   "outputs": [],
-   "source": [
-    "# The native report is saved as power_iteration__async.nsys-rep."
-   ]
-  },
   {
    "cell_type": "markdown",
    "id": "b8a12d4e",

tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb

@@ -24,9 +24,7 @@
     "\n",
     "In this exercise, we'll learn how to analyze and reason about the performance of CUDA kernels using the NVIDIA Nsight Compute profiler. We'll look at a few different ways of writing a simple kernel that copies items from one array to another.\n",
     "\n",
-    "First, we need to make sure the Nsight Compute profiler, Nsightful, Numba CUDA, and CuPy are available in our notebook:\n",
-    "\n",
-    "The profiling sections require the **Python 3 (Nsight Compute)** custom kernel provided by the ACH environment. Google Colab cannot select this kernel."
+    "First, we need to ensure NVIDIA's developer tools are installed and available and do all of our imports."
    ]
   },
   {
@@ -51,7 +49,7 @@
     "  print(\"Uninstalling PIP packages.\")\n",
     "  !pip uninstall \"cuda-python\" --yes > /dev/null\n",
     "  print(\"Installing PIP packages.\")\n",
-    "  !pip install \"numba-cuda\" \"cuda-cccl[test-cu12]\" \"nvtx\" \"nsightful[notebook] @ git+https://github.com/brycelelbach/nsightful.git@fa9ee4d81441ac62379c5306f2f2d8b0894d06ec\" > /dev/null 2>&1\n",
+    "  !pip install \"numba-cuda\" \"cuda-cccl[test-cu12]\" \"nvtx\" \"nsightful[notebook] @ git+https://github.com/brycelelbach/nsightful.git@a41989403430168e02ac3cfdc4060bca4ebb8040\" > /dev/null 2>&1\n",
     "  open(\"/accelerated-computing-hub-installed\", \"a\").close()\n",
     "  print(\"All packages installed.\")\n",
     "\n",
@@ -99,11 +97,7 @@
     "def copy_blocked(src, dst, items_per_thread):\n",
     " base = cuda.grid(1) * items_per_thread\n",
     " for i in range(items_per_thread):\n",
-    "   dst[base + i] = src[base + i]\n",
-    "\n",
-    "\n",
-    "def launch_blocked():\n",
-    "  copy_blocked[blocks, threads_per_block](src, dst, items_per_thread)\n"
+    "   dst[base + i] = src[base + i]"
    ]
   },
   {
@@ -123,9 +117,10 @@
    "metadata": {},
    "outputs": [],
    "source": [
-    "launch_blocked()\n",
+    "dst[:] = 0\n",
+    "copy_blocked[blocks, threads_per_block](src, dst, items_per_thread)\n",
     "cp.testing.assert_array_equal(src, dst)\n",
-    "print(f\"Problem size: {total_items * src.dtype.itemsize / 2**30:.2f} GB, dtype: {src.dtype}\")\n"
+    "print(f\"Problem size: {total_items * src.dtype.itemsize / 2**30:.2f} GB, dtype: {src.dtype}\")"
    ]
   },
   {
@@ -151,7 +146,8 @@
    },
    "outputs": [],
    "source": [
-    "%%ncu -o copy_blocked.ncu-rep\n",
+    "%%ncu -o copy_blocked.ncu-rep --kernel-name regex:copy_blocked\n",
+    "dst[:] = 0\n",
     "copy_blocked[blocks, threads_per_block](src, dst, items_per_thread)"
    ]
   },
@@ -166,7 +162,7 @@
     "\n",
     "**TODO:** Spend a few minutes reviewing the report. What stands out to you? Based on the information in the report, how can the kernel be improved?\n",
     "\n",
-    "**EXTRA CREDIT:** Download the [Nsight Compute GUI](https://developer.nvidia.com/nsight-compute) and open the report in it to see even more information."
+    "The Nsight Compute GUI provides richer information, charts, and diagrams. Click the **+** in the JupyterLab tab bar and open Nsight Compute, or install the GUI on your local system and download the `.ncu-rep` report."
    ]
   },
   {
@@ -212,11 +208,7 @@
    "source": [
     "@cuda.jit\n",
     "def copy_optimized(src, dst, items_per_thread):\n",
-    "  TODO() # TODO: You need to implement this kernel! DELETE THIS LINE.\n",
-    "\n",
-    "\n",
-    "def launch_optimized():\n",
-    "  copy_optimized[blocks, threads_per_block](src, dst, items_per_thread)\n"
+    "  TODO() # TODO: You need to implement this kernel! DELETE THIS LINE."
    ]
   },
   {
@@ -240,8 +232,9 @@
    },
    "outputs": [],
    "source": [
-    "launch_optimized()\n",
-    "cp.testing.assert_array_equal(src, dst)\n"
+    "dst[:] = 0\n",
+    "copy_optimized[blocks, threads_per_block](src, dst, items_per_thread)\n",
+    "cp.testing.assert_array_equal(src, dst)"
    ]
   },
   {
@@ -263,15 +256,15 @@
    },
    "outputs": [],
    "source": [
-    "blocked_times = cpx.profiler.benchmark(launch_blocked, n_repeat=15, n_warmup=1).gpu_times[0]\n",
-    "optimized_times = cpx.profiler.benchmark(launch_optimized, n_repeat=15, n_warmup=1).gpu_times[0]\n",
+    "blocked_times = cpx.profiler.benchmark(lambda: copy_blocked[blocks, threads_per_block](src, dst, items_per_thread), n_repeat=15, n_warmup=1).gpu_times[0]\n",
+    "optimized_times = cpx.profiler.benchmark(lambda: copy_optimized[blocks, threads_per_block](src, dst, items_per_thread), n_repeat=15, n_warmup=1).gpu_times[0]\n",
     "copy_blocked_duration = blocked_times.mean() * 1000\n",
     "copy_optimized_duration = optimized_times.mean() * 1000\n",
     "speedup = copy_blocked_duration / copy_optimized_duration\n",
     "\n",
     "print(f\"copy_blocked:   {copy_blocked_duration:.3g} ms\")\n",
     "print(f\"copy_optimized: {copy_optimized_duration:.3g} ms\")\n",
-    "print(f\"copy_optimized speedup over copy_blocked: {speedup:.2f}\")\n"
+    "print(f\"copy_optimized speedup over copy_blocked: {speedup:.2f}\")"
    ]
   },
   {
@@ -295,7 +288,8 @@
    },
    "outputs": [],
    "source": [
-    "%%ncu -o copy_optimized.ncu-rep\n",
+    "%%ncu -o copy_optimized.ncu-rep --kernel-name regex:copy_optimized\n",
+    "dst[:] = 0\n",
     "copy_optimized[blocks, threads_per_block](src, dst, items_per_thread)"
    ]
   },
@@ -332,6 +326,18 @@
     "\n",
     "**EXTRA CREDIT:** Experiment with different problem sizes, threads per block, and items per thread by changing the configuration variables above. If you're feeling really ambitious, do a parameter sweep to study the impact these knobs have on performance."
    ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "id": "a845f694",
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# Try changing total_items, threads_per_block, and items_per_thread above.\n",
+    "# Rerun the kernel definitions, correctness checks, and benchmarks after each change.\n",
+    "# For a deeper study, sweep several values and plot runtime or memory throughput."
+   ]
   }
  ],
  "metadata": {