diff --git a/tutorials/pyhpc/brev/dockerfile b/tutorials/pyhpc/brev/dockerfile index 1c43c253..d012dc81 100644 --- a/tutorials/pyhpc/brev/dockerfile +++ b/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' diff --git a/tutorials/pyhpc/brev/requirements.txt b/tutorials/pyhpc/brev/requirements.txt index d1cd0cb2..4e6616cf 100644 --- a/tutorials/pyhpc/brev/requirements.txt +++ b/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 diff --git a/tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb b/tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb index 73b16954..c89a87c9 100644 --- a/tutorials/pyhpc/notebooks/03__power_iteration__cupy__asynchrony.ipynb +++ b/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", diff --git a/tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb b/tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb index cc3f7eec..62af4a5d 100644 --- a/tutorials/pyhpc/notebooks/04__copy__kernel_authoring.ipynb +++ b/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": { diff --git a/tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb b/tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb index a90a78bd..9de6e752 100644 --- a/tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb +++ b/tutorials/pyhpc/notebooks/05__book_histogram__kernel_authoring.ipynb @@ -15,7 +15,7 @@ "2. [First Attempt: Global Memory Histogram](#2.-First-Attempt:-Global-Memory-Histogram)\n", "3. [Fixing Data Races with Atomics](#3.-Fixing-Data-Races-with-Atomics)\n", "4. [Profiling the Naive Solution](#4.-Profiling-the-Naive-Solution)\n", - "5. [Optimization: Shared Memory & Cooperative Groups](#5.-Optimization:-Shared-Memory-&-Cooperative-Groups)\n", + "5. [Optimization: Shared Memory](#5.-Optimization:-Shared-Memory)\n", "6. [Performance Comparison](#6.-Performance-Comparison)\n", "\n", "### 1. Environment Setup & Data Download\n", @@ -46,7 +46,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", @@ -55,7 +55,7 @@ "import matplotlib.pyplot as plt\n", "from numba import cuda\n", "import cupy as cp\n", - "import cupyx as cpx\n" + "import cupyx as cpx" ] }, { @@ -96,8 +96,10 @@ "outputs": [], "source": [ "bins = 256\n", + "\n", "values = cp.fromfile(\"books__15m.txt\", dtype=cp.uint8)\n", "histogram = cp.zeros(bins, dtype=cp.int32)\n", + "\n", "threads_per_block = 512\n", "items_per_thread = 8\n", "items_per_block = threads_per_block * items_per_thread\n", @@ -110,12 +112,7 @@ " value = values[cuda.grid(1) * items_per_thread + i]\n", " old_count = histogram[value]\n", " new_count = old_count + 1\n", - " histogram[value] = new_count\n", - "\n", - "\n", - "def launch_global():\n", - " histogram[:] = 0\n", - " histogram_global[blocks, threads_per_block](values, histogram)\n" + " histogram[value] = new_count" ] }, { @@ -137,19 +134,10 @@ }, "outputs": [], "source": [ - "launch_global()\n", - "cp.cuda.runtime.deviceSynchronize()\n" - ] - }, - { - "cell_type": "code", - "execution_count": null, - "id": "815cb072-b3a0-47aa-af6c-dd66c626a440", - "metadata": { - "id": "815cb072-b3a0-47aa-af6c-dd66c626a440" - }, - "outputs": [], - "source": [ + "histogram[:] = 0\n", + "histogram_global[blocks, threads_per_block](values, histogram)\n", + "assert cp.sum(histogram) < len(values)\n", + "\n", "histogram_host = cp.asnumpy(histogram)\n", "\n", "# Print most frequently occurring characters.\n", @@ -161,7 +149,7 @@ "plt.title(\"Top 20 Bins\")\n", "plt.show()\n", "\n", - "print(f\"Characters in dataset: {values.size / 1e6:.1f} MB\")\n" + "print(f\"Characters in dataset: {values.size / 1e6:.1f} MB\")" ] }, { @@ -212,10 +200,9 @@ }, "outputs": [], "source": [ - "%%ncu -o histogram_global.ncu-rep\n", + "%%ncu -o histogram_global.ncu-rep --kernel-name regex:histogram_global\n", "histogram[:] = 0\n", - "histogram_global[blocks, threads_per_block](values, histogram)\n", - "cp.cuda.runtime.deviceSynchronize()" + "histogram_global[blocks, threads_per_block](values, histogram)" ] }, { @@ -227,7 +214,7 @@ }, "outputs": [], "source": [ - "assert cp.sum(histogram) < len(values), \"The intentionally racy kernel should lose updates.\"" + "assert cp.sum(histogram) < len(values)" ] }, { @@ -237,7 +224,7 @@ "id": "e1f72831-780f-4cf5-8ff1-2092ecb193d9" }, "source": [ - "### 5. Optimization: Shared Memory & Cooperative Groups\n", + "### 5. Optimization: Shared Memory\n", "\n", "Looking at the profile trace, it seems like our code is quite slow - look at the memory workload tab and see how low the throughput is!\n", "\n", @@ -265,23 +252,21 @@ }, "outputs": [], "source": [ - "localized_items_per_thread = 8\n", - "localized_items_per_block = threads_per_block * localized_items_per_thread\n", - "localized_blocks = len(values) // localized_items_per_block\n", - "assert values.size % localized_items_per_block == 0\n", + "items_per_thread = 8\n", + "items_per_block = threads_per_block * items_per_thread\n", + "blocks = len(values) // items_per_block\n", + "assert values.size % items_per_block == 0\n", "\n", "@cuda.jit\n", "def histogram_localized(values, histogram):\n", - " for i in range(localized_items_per_thread):\n", - " value = values[cuda.grid(1) * localized_items_per_thread + i]\n", + " for i in range(items_per_thread):\n", + " value = values[cuda.grid(1) * items_per_thread + i]\n", " old_count = histogram[value]\n", " new_count = old_count + 1\n", " histogram[value] = new_count\n", "\n", - "\n", "def launch_localized():\n", - " histogram[:] = 0\n", - " histogram_localized[localized_blocks, threads_per_block](values, histogram)\n" + " histogram_localized[blocks, threads_per_block](values, histogram)" ] }, { @@ -303,20 +288,10 @@ }, "outputs": [], "source": [ + "histogram[:] = 0\n", "launch_localized()\n", - "cp.cuda.runtime.deviceSynchronize()\n", - "assert cp.sum(histogram) == len(values)\n" - ] - }, - { - "cell_type": "code", - "execution_count": null, - "id": "73f0c3cd-349b-490f-b6bd-7afbeb442fff", - "metadata": { - "id": "73f0c3cd-349b-490f-b6bd-7afbeb442fff" - }, - "outputs": [], - "source": [ + "assert cp.sum(histogram) == len(values)\n", + "\n", "histogram_host = cp.asnumpy(histogram)\n", "pairs = sorted(((i, c) for i, c in enumerate(histogram_host) if c), key=lambda x: x[1], reverse=True)[:20]\n", "labels = [('SPACE' if i == 32 else chr(i)) if 32 <= i <= 126 else f'0x{i:02X}' for i, _ in pairs]\n", @@ -324,7 +299,7 @@ "plt.xlabel('count')\n", "plt.tight_layout()\n", "plt.title(\"Top 20 Bins\")\n", - "plt.show()\n" + "plt.show()" ] }, { @@ -336,10 +311,9 @@ }, "outputs": [], "source": [ - "%%ncu -o histogram_localized.ncu-rep\n", + "%%ncu -o histogram_localized.ncu-rep --kernel-name regex:histogram_localized\n", "histogram[:] = 0\n", - "histogram_localized[localized_blocks, threads_per_block](values, histogram)\n", - "cp.cuda.runtime.deviceSynchronize()" + "histogram_localized[blocks, threads_per_block](values, histogram)" ] }, { @@ -373,7 +347,7 @@ }, "outputs": [], "source": [ - "global_times = cpx.profiler.benchmark(launch_global, n_repeat=15, n_warmup=4).gpu_times[0]\n", + "global_times = cpx.profiler.benchmark(lambda: histogram_global[blocks, threads_per_block](values, histogram), n_repeat=15, n_warmup=4).gpu_times[0]\n", "localized_times = cpx.profiler.benchmark(launch_localized, n_repeat=15, n_warmup=4).gpu_times[0]\n", "histogram_global_duration = global_times.mean() * 1000\n", "histogram_localized_duration = localized_times.mean() * 1000\n", @@ -381,7 +355,7 @@ "\n", "print(f\"histogram_global: {histogram_global_duration:.3g} ms\")\n", "print(f\"histogram_localized: {histogram_localized_duration:.3g} ms\")\n", - "print(f\"histogram_localized speedup over histogram_global: {speedup:.2f}\")\n" + "print(f\"histogram_localized speedup over histogram_global: {speedup:.2f}\")" ] } ], diff --git a/tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb index f039c1ca..5b818df9 100644 --- a/tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/03__power_iteration__cupy__asynchrony__SOLUTION.ipynb @@ -23,7 +23,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." ] }, { @@ -163,54 +163,60 @@ " 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", - "\n", - "estimate_device_baseline(\n", - " A_device,\n", - " cfg=PowerIterationConfig(max_steps=1, check_frequency=1, progress=False),\n", - ")" + " return cp.asnumpy((x.T @ (A_gpu @ x)) / (x.T @ x))" ] }, { "cell_type": "markdown", - "id": "2f3e8e69", + "id": "2ca78d38", "metadata": {}, "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." + "Now let's make sure it works:" ] }, { "cell_type": "code", "execution_count": null, - "id": "4fbea5ca", + "id": "fbda9e32", "metadata": {}, "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)" + "\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", + "print()\n", + "print(\"Dominant Eigenvalue:\", lam_est_baseline)" ] }, { "cell_type": "markdown", - "id": "2b09a36e", + "id": "2f3e8e69", "metadata": {}, "source": [ - "The timeline is displayed below the profiled cell. Explore what's going on in the program.\n", + "### 4. Profiling the Baseline\n", "\n", - "**EXTRA CREDIT:** Download the Nsight Systems GUI and open the report in it to see even more information." + "The `%%nsys` cell magic profiles the code in its cell with Nsight Systems and saves the native report to the path given by `-o`." ] }, { "cell_type": "code", "execution_count": null, - "id": "71ac2641", + "id": "4fbea5ca", "metadata": {}, "outputs": [], "source": [ - "# The native report is saved as power_iteration__baseline.nsys-rep." + "%%nsys -o power_iteration__baseline.nsys-rep\n", + "lam_est_baseline = estimate_device_baseline(A_device)" + ] + }, + { + "cell_type": "markdown", + "id": "2b09a36e", + "metadata": {}, + "source": [ + "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." ] }, { @@ -332,7 +338,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)" ] }, { @@ -382,8 +393,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)" ] }, { @@ -394,16 +404,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": "65666ee6", - "metadata": {}, - "outputs": [], - "source": [ - "# The native report is saved as power_iteration__async.nsys-rep." - ] - }, { "cell_type": "markdown", "id": "a47d9c2b", diff --git a/tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb index eb77a81d..8fc5fb5e 100644 --- a/tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/04__copy__kernel_authoring__SOLUTION.ipynb @@ -23,12 +23,12 @@ "\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:" + "First, we need to ensure NVIDIA's developer tools are installed and available and do all of our imports." ] }, { "cell_type": "code", - "execution_count": 1, + "execution_count": null, "id": "82b8f596", "metadata": { "execution": { @@ -54,7 +54,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", @@ -83,7 +83,7 @@ }, { "cell_type": "code", - "execution_count": 2, + "execution_count": null, "id": "f724215c", "metadata": { "colab": { @@ -99,15 +99,7 @@ "id": "I9Tz2hG-_tBj", "outputId": "e52f41cb-f70b-4792-ea76-e78a554956f4" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Writing copy_blocked.py\n" - ] - } - ], + "outputs": [], "source": [ "threads_per_block = 256\n", "items_per_thread = 64\n", @@ -121,11 +113,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]" ] }, { @@ -140,7 +128,7 @@ }, { "cell_type": "code", - "execution_count": 3, + "execution_count": null, "id": "baacfd0d", "metadata": { "execution": { @@ -151,19 +139,12 @@ "shell.execute_reply.started": "2026-03-09T19:43:42.408663Z" } }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Problem size: 2.00 GB, dtype: int64\n" - ] - } - ], + "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}\")" ] }, { @@ -182,7 +163,7 @@ }, { "cell_type": "code", - "execution_count": 4, + "execution_count": null, "id": "54aea2c6", "metadata": { "colab": { @@ -198,22 +179,11 @@ "id": "5pyHvJtxVnDB", "outputId": "87ad5a72-9244-4b21-9776-ecd93262f274" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "==PROF== Connected to process 644 (/usr/bin/python3.12)\n", - "==PROF== Profiling \"copy_blocked[abi:v1,cw51cXTLSUwv1sDUaKthoaNgqamjgOR3W3Cw6igA9duC0hkwnNGiHEkYkrqQBFBTDEwCyQe21eqwcFW3UoDGjzpQEsgzrtUEAA_3d_3d]\": 0%....50%....100% - 44 passes\n", - "==PROF== Disconnected from process 644\n", - "==PROF== Report: /accelerated-computing-hub/tutorials/accelerated-python/notebooks/kernels/solutions/copy_blocked.ncu-rep\n" - ] - } - ], + "outputs": [], "source": [ - "%%ncu -o copy_blocked.ncu-rep\n", - "copy_blocked[blocks, threads_per_block](src, dst, items_per_thread)\n", - "cp.cuda.runtime.deviceSynchronize()" + "%%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)" ] }, { @@ -225,12 +195,12 @@ "source": [ "Let's take a look at the profiling report on the kernel. When you run the next cell, a number of tabs will be displayed. The first tab will have a summary of all of the Nsight recommendations and advisories. Subsequent tabs will have more detailed information on a particular area.\n", "\n", - "Remember, you can see even more information in the [Nsight Compute GUI](https://developer.nvidia.com/nsight-compute)." + "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." ] }, { "cell_type": "code", - "execution_count": 5, + "execution_count": null, "id": "4e99c6b9", "metadata": { "colab": { @@ -280,87 +250,7 @@ "id": "40w07iG5k6Vl", "outputId": "f5d0becc-0285-4cb7-a78e-427cdc105454" }, - "outputs": [ - { - "data": { - "text/html": [ - "\n", - " \n", - " " - ], - "text/plain": [ - "" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "ef5ef643588149639d9499dde26dd7bf", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Dropdown(description='Kernel:', layout=Layout(width='400px'), options=('copy_blocked',), style=DescriptionStyl…" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "8daf9c2eec544287875436afe1817f0b", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Output()" - ] - }, - "metadata": {}, - "output_type": "display_data" - } - ], + "outputs": [], "source": [ "cp.testing.assert_array_equal(src, dst)" ] @@ -397,7 +287,7 @@ }, { "cell_type": "code", - "execution_count": 6, + "execution_count": null, "id": "d7b514e9", "metadata": { "colab": { @@ -413,15 +303,7 @@ "id": "B5PBpaY2HnE0", "outputId": "9bf8af36-f011-4eaa-ed28-c39abde2c315" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Writing copy_optimized.py\n" - ] - } - ], + "outputs": [], "source": [ "@cuda.jit\n", "def copy_optimized(src, dst, items_per_thread):\n", @@ -432,11 +314,7 @@ "\n", " base = tx + bx * items_per_block\n", " for i in range(0, items_per_block, bd):\n", - " dst[base + i] = src[base + i]\n", - "\n", - "\n", - "def launch_optimized():\n", - " copy_optimized[blocks, threads_per_block](src, dst, items_per_thread)\n" + " dst[base + i] = src[base + i]" ] }, { @@ -453,7 +331,7 @@ }, { "cell_type": "code", - "execution_count": 7, + "execution_count": null, "id": "20fca1a6", "metadata": { "execution": { @@ -464,18 +342,11 @@ "shell.execute_reply.started": "2026-03-09T19:44:09.977991Z" } }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Problem size: 2.00 GB, dtype: int64\n" - ] - } - ], + "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)" ] }, { @@ -488,7 +359,7 @@ }, { "cell_type": "code", - "execution_count": 8, + "execution_count": null, "id": "0921fb67", "metadata": { "colab": { @@ -504,27 +375,17 @@ "id": "kJ7viF-i06qd", "outputId": "e2d84395-6324-4e55-c144-bc34399cc515" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "copy_blocked: 38.1 ms ± 0.43% (mean ± relative stdev of 15 runs)\n", - "copy_optimized: 18.3 ms ± 0.28% (mean ± relative stdev of 15 runs)\n", - "copy_optimized speedup over copy_blocked: 2.08\n" - ] - } - ], + "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}\")" ] }, { @@ -541,7 +402,7 @@ }, { "cell_type": "code", - "execution_count": 9, + "execution_count": null, "id": "79bc5620", "metadata": { "colab": { @@ -557,22 +418,11 @@ "id": "zO_y6ObXV_wX", "outputId": "8e643e49-6fbf-4b1e-ff58-d391700451ab" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "==PROF== Connected to process 785 (/usr/bin/python3.12)\n", - "==PROF== Profiling \"copy_optimized[abi:v1,cw51cXTLSUwv1sDUaKthoaNgqamjgOR3W3Cw6igA9duC0hkwnNGiHEkYkrqQBFBTDEwCyQe21eqwcFW3UoDGjzpQEsgzrtUEAA_3d_3d]\": 0%....50%....100% - 44 passes\n", - "==PROF== Disconnected from process 785\n", - "==PROF== Report: /accelerated-computing-hub/tutorials/accelerated-python/notebooks/kernels/solutions/copy_optimized.ncu-rep\n" - ] - } - ], + "outputs": [], "source": [ - "%%ncu -o copy_optimized.ncu-rep\n", - "copy_optimized[blocks, threads_per_block](src, dst, items_per_thread)\n", - "cp.cuda.runtime.deviceSynchronize()" + "%%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)" ] }, { @@ -587,7 +437,7 @@ }, { "cell_type": "code", - "execution_count": 10, + "execution_count": null, "id": "8f83e8fe", "metadata": { "colab": { @@ -637,90 +487,32 @@ "id": "KjE0Vgu_zgs3", "outputId": "632a4728-519d-48fa-938e-7a9777d52c89" }, - "outputs": [ - { - "data": { - "text/html": [ - "\n", - " \n", - " " - ], - "text/plain": [ - "" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "2ef5798ed5d84b439f29fbf7da7c0020", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Dropdown(description='Kernel:', layout=Layout(width='400px'), options=('copy_optimized',), style=DescriptionSt…" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "51106bf176b2450c99d5dac3ba9dccb1", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Output()" - ] - }, - "metadata": {}, - "output_type": "display_data" - } - ], + "outputs": [], "source": [ "cp.testing.assert_array_equal(src, dst)" ] + }, + { + "cell_type": "markdown", + "id": "a5f46116", + "metadata": {}, + "source": [ + "### 7. Further Exploration\n", + "\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": "90d2bf86", + "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": { diff --git a/tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb b/tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb index 1737d7b8..b2e2da48 100644 --- a/tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb +++ b/tutorials/pyhpc/notebooks/solutions/05__book_histogram__kernel_authoring__SOLUTION.ipynb @@ -15,7 +15,7 @@ "2. [First Attempt: Global Memory Histogram](#2.-First-Attempt:-Global-Memory-Histogram)\n", "3. [Fixing Data Races with Atomics](#3.-Fixing-Data-Races-with-Atomics)\n", "4. [Profiling the Naive Solution](#4.-Profiling-the-Naive-Solution)\n", - "5. [Solution: Shared Memory & Cooperative Groups](#5.-Solution:-Shared-Memory-&-Cooperative-Groups)\n", + "5. [Solution: Shared Memory](#5.-Solution:-Shared-Memory)\n", "6. [Performance Comparison](#6.-Performance-Comparison)\n", "\n", "### 1. Environment Setup & Data Download\n", @@ -27,7 +27,7 @@ }, { "cell_type": "code", - "execution_count": 1, + "execution_count": null, "id": "ce42d5e5-db1e-46da-a64a-831d0f3d59ff", "metadata": { "execution": { @@ -53,7 +53,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", @@ -62,12 +62,12 @@ "import matplotlib.pyplot as plt\n", "from numba import cuda\n", "import cupy as cp\n", - "import cupyx as cpx\n" + "import cupyx as cpx" ] }, { "cell_type": "code", - "execution_count": 2, + "execution_count": null, "id": "6e1ae9dc-39f1-4c93-b923-85a96b45a057", "metadata": { "colab": { @@ -83,18 +83,7 @@ "id": "6e1ae9dc-39f1-4c93-b923-85a96b45a057", "outputId": "ee334037-7f39-4a91-ad60-f0408a56b4de" }, - "outputs": [ - { - "data": { - "text/plain": [ - "('books__15m.txt', )" - ] - }, - "execution_count": 2, - "metadata": {}, - "output_type": "execute_result" - } - ], + "outputs": [], "source": [ "urllib.request.urlretrieve(\n", " \"https://drive.usercontent.google.com/download?id=1MW1lPgkTq3YG9ikuq6u3d9sfpt-wKQZ0&export=download\",\n", @@ -117,7 +106,7 @@ }, { "cell_type": "code", - "execution_count": 3, + "execution_count": null, "id": "61c12795-b14a-4447-9dcf-9748616cc453", "metadata": { "colab": { @@ -133,19 +122,13 @@ "id": "61c12795-b14a-4447-9dcf-9748616cc453", "outputId": "141b29b8-bbf7-4224-c85a-23e49ee7abaf" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Writing histogram_global.py\n" - ] - } - ], + "outputs": [], "source": [ "bins = 256\n", + "\n", "values = cp.fromfile(\"books__15m.txt\", dtype=cp.uint8)\n", "histogram = cp.zeros(bins, dtype=cp.int32)\n", + "\n", "threads_per_block = 512\n", "items_per_thread = 8\n", "items_per_block = threads_per_block * items_per_thread\n", @@ -156,12 +139,7 @@ "def histogram_global(values, histogram):\n", " for i in range(items_per_thread):\n", " value = values[cuda.grid(1) * items_per_thread + i]\n", - " cuda.atomic.add(histogram, value, 1)\n", - "\n", - "\n", - "def launch_global():\n", - " histogram[:] = 0\n", - " histogram_global[blocks, threads_per_block](values, histogram)\n" + " cuda.atomic.add(histogram, value, 1)" ] }, { @@ -176,7 +154,7 @@ }, { "cell_type": "code", - "execution_count": 4, + "execution_count": null, "id": "b3f32240-ad7b-4717-98b3-82b0298a099a", "metadata": { "colab": { @@ -192,59 +170,12 @@ "id": "b3f32240-ad7b-4717-98b3-82b0298a099a", "outputId": "cdf42faf-b6e6-45ad-85e2-d3b0d4c87ad5" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "4.19 ms ± 1.13% (mean ± relative stdev of 15 runs)\n" - ] - } - ], - "source": [ - "launch_global()\n", - "cp.cuda.runtime.deviceSynchronize()\n" - ] - }, - { - "cell_type": "code", - "execution_count": 5, - "id": "815cb072-b3a0-47aa-af6c-dd66c626a440", - "metadata": { - "colab": { - "base_uri": "https://localhost:8080/", - "height": 487 - }, - "execution": { - "iopub.execute_input": "2026-03-09T19:42:47.997273Z", - "iopub.status.busy": "2026-03-09T19:42:47.996885Z", - "iopub.status.idle": "2026-03-09T19:42:50.354006Z", - "shell.execute_reply": "2026-03-09T19:42:50.352914Z", - "shell.execute_reply.started": "2026-03-09T19:42:47.997232Z" - }, - "id": "815cb072-b3a0-47aa-af6c-dd66c626a440", - "outputId": "d7a3d2e2-1df5-4681-c6a3-923cfb921c52" - }, - "outputs": [ - { - "data": { - "image/png": "iVBORw0KGgoAAAANSUhEUgAAAnUAAAHsCAYAAAC0dltCAAAAOnRFWHRTb2Z0d2FyZQBNYXRwbG90bGliIHZlcnNpb24zLjEwLjgsIGh0dHBzOi8vbWF0cGxvdGxpYi5vcmcvwVt1zgAAAAlwSFlzAAAPYQAAD2EBqD+naQAAOUVJREFUeJzt3Xt0FPXh/vFnE5JJAuxCuCTQBiK3CAbCRaEISFQgIKXFiiItELBKS6USoaDRn8RQ6yKKoBVL5RYQWxVRRKkgIkFECgpBuShoJBCVO7jLRRdM5veHX7ZGEiBxs5OdvF/nzDnu7Gd2n2FOso+fmdk4TNM0BQAAgJAWZnUAAAAA/HSUOgAAABug1AEAANgApQ4AAMAGKHUAAAA2QKkDAACwAUodAACADVDqAAAAbIBSBwAAYAOUOgAIATk5OXI4HCooKLA6CoAqilIHoEpwOByXtOTm5lZqjsLCQmVnZ6tz586qW7eu6tevr9TUVL311luljv/66681atQoNWjQQDVr1tS1116rLVu2XNJ7paamlti3yMhIXXbZZRo1apQKCwsDuVsAqgEHf/sVQFWwaNGiEo8XLlyoVatW6dlnny2xvnfv3oqLi6u0HE899ZQmTpyogQMHqlu3bvruu++0cOFCbdmyRfPmzdPIkSP9Y4uLi9WjRw99+OGHmjBhgurXr6+nn35ahYWF2rx5s1q2bHnB90pNTVV+fr7cbrck6cyZM9q5c6dmzZqlevXq6eOPP1ZMTIwkqaioSGfPnpVhGHI4HJW2/wBCF6UOQJU0ZswYzZw5U8H+FbVjxw7FxcWpfv36/nU+n0/t27fXyZMnS8ygvfjiixo8eLAWL16sQYMGSZIOHz6sVq1aqV+/fvrXv/51wfdKTU3VkSNHtH379hLrZ86cqTFjxujNN99U7969A7h3AOyM068AQsapU6c0fvx4JSQkyDAMJSUl6bHHHjuv+DkcDo0ZM0bPPfeckpKSFBUVpU6dOumdd9656HtcccUVJQqdJBmGoRtuuEFffPGFTpw44V//0ksvKS4uTr/5zW/86xo0aKBbbrlFr776qnw+X4X2Mz4+XpJUo0YN/7rSrqlLTEzUL3/5S7377rvq3LmzoqKi1KxZMy1cuLDE6509e1bZ2dlq2bKloqKiVK9ePXXv3l2rVq2qUD4AVROlDkBIME1Tv/rVrzR9+nT17dtXjz/+uJKSkjRhwgSNGzfuvPFr165VRkaGhg4dqsmTJ+vo0aPq27fvebNil+rAgQOKiYnxnw6VpLy8PHXs2FFhYSV/lXbu3FmnT5/W7t27L/q6RUVFOnLkiI4cOaL9+/fr7bffVlZWllq0aKFu3bpddPvPPvtMgwYNUu/evTVt2jTVrVtXI0aM0I4dO/xjHnzwQWVnZ+vaa6/VU089pfvvv19NmjS55Gv/AIQIEwCqoDvvvNP84a+opUuXmpLMhx56qMS4QYMGmQ6Hw/zss8/86ySZkswPPvjAv27v3r1mVFSUeeONN5Y7y6effmpGRUWZw4YNK7G+Zs2a5m233Xbe+OXLl5uSzBUrVlzwdXv27OnP+sOldevW5ueff15i7Pz5801J5p49e/zrmjZtakoy33nnHf+6Q4cOmYZhmOPHj/evS0lJMfv371+eXQYQgpipAxAS/vOf/yg8PFx33XVXifXjx4+XaZp64403Sqzv2rWrOnXq5H/cpEkT/frXv9bKlStVVFR0ye97+vRp3XzzzYqOjtaUKVNKPPfNN9/IMIzztomKivI/fzGJiYlatWqVVq1apTfeeEMzZsyQx+NRv379dPjw4Ytu36ZNG/Xo0cP/uEGDBkpKStLnn3/uX1enTh3t2LFDn3766UVfD0DootQBCAl79+5V48aNVbt27RLrW7du7X/+h0q787RVq1Y6ffr0JZUl6ftTo7feeqt27typl156SY0bNy7xfHR0dKnXzX377bf+5y+mZs2a6tWrl3r16qW+fftq7NixWrZsmXbt2nVeiSxNkyZNzltXt25dHT9+3P948uTJ+vrrr9WqVSu1bdtWEyZM0EcffXTR1wYQWih1AFCGO+64Q6+//rpycnJ03XXXnfd8o0aNtH///vPWn1v34xJ4qTp16iSXy3VJN3aEh4eXut78wc0j11xzjfLz8zVv3jwlJydrzpw56tixo+bMmVOhfACqJkodgJDQtGlTffXVVyXuPpWkTz75xP/8D5V2qnH37t2KiYlRgwYNLvp+EyZM0Pz58zV9+nQNGTKk1DHt27fXli1bVFxcXGL9xo0bFRMTo1atWl30fcpSVFSkkydPVnj7H4uNjdXIkSP173//W4WFhWrXrp0efPDBgL0+AOtR6gCEhBtuuEFFRUV66qmnSqyfPn26HA6H+vXrV2L9hg0bStzdWVhYqFdffVV9+vQpc3brnEcffVSPPfaY7rvvPo0dO7bMcYMGDdLBgwf18ssv+9cdOXJEixcv1oABA0q93u5SrFmzRidPnlRKSkqFtv+xo0ePlnhcq1YttWjRosJfuQKgaqpx8SEAYL0BAwbo2muv1f3336+CggKlpKTozTff1KuvvqqMjAw1b968xPjk5GSlpaXprrvukmEYevrppyVJ2dnZF3yfV155RRMnTlTLli3VunXr8/7SxQ//osWgQYP0i1/8QiNHjtTOnTv9f1GiqKjoou9zjsfj8b/Hd999p127dukf//iHoqOjde+9917Sa1xMmzZtlJqaqk6dOik2NlYffPCBXnrpJY0ZMyYgrw+gaqDUAQgJYWFhWrZsmSZNmqQXXnhB8+fPV2Jioh599FGNHz/+vPE9e/ZU165dlZ2drX379qlNmzbKyclRu3btLvg+H374oaTvT98OGzbsvOfXrFnjL3Xh4eH6z3/+owkTJujJJ5/UN998o6uuuko5OTlKSkq6pP364osv/O/jcDhUt25d9ezZU1lZWWrfvv0lvcbF3HXXXVq2bJnefPNN+Xw+NW3aVA899JAmTJgQkNcHUDXwZ8IA2I7D4dCdd9553qlaALAzrqkDAACwAUodAACADVDqAAAAbIAbJQDYDpcKA6iOmKkDAACwAUodAACADVTr06/FxcX66quvVLt2bTkcDqvjAACAasw0TZ04cUKNGzdWWFj5592qdan76quvlJCQYHUMAAAAv8LCQv385z8v93bVutTVrl1b0vf/eE6n0+I0AACgOvN6vUpISPD3k/Kq1qXu3ClXp9NJqQMAAFVCRS8J40YJAAAAG6DUAQAA2AClDgAAwAYodQAAADZAqQMAALABSh0AAIANUOoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2AClDgAAwAYodQAAADZAqQMAALABSh0AAIANUOoAAABsoIbVAaqC5KyVCjNirI4BAACqoIIp/a2OcEmYqQMAALABSh0AAIANUOoAAABsoNyl7vDhwxo9erSaNGkiwzAUHx+vtLQ0rV+/XpKUmJgoh8Mhh8OhmjVrqmPHjlq8eHGJ1/jmm28UGxur+vXry+fzlfo+S5YsUWpqqlwul2rVqqV27dpp8uTJOnbsmCQpJyfH/z4/XKKiosq7SwAAACGv3KXupptuUl5enhYsWKDdu3dr2bJlSk1N1dGjR/1jJk+erP379ysvL09XXXWVBg8erPfee8///JIlS3TFFVfo8ssv19KlS897j/vvv1+DBw/WVVddpTfeeEPbt2/XtGnT9OGHH+rZZ5/1j3M6ndq/f3+JZe/eveXdJQAAgJBXrrtfv/76a61bt065ubnq2bOnJKlp06bq3LlziXG1a9dWfHy84uPjNXPmTC1atEivvfaarr76aknS3LlzNXToUJmmqblz52rw4MH+bTdt2qSHH35YM2bM0NixY/3rExMT1bt3b3399df+dQ6HQ/Hx8eXeaQAAALsp10xdrVq1VKtWLS1durTM06Y/VqNGDUVEROjMmTOSpPz8fG3YsEG33HKLbrnlFq1bt67E7Npzzz2nWrVq6U9/+lOpr1enTp3yRC7B5/PJ6/WWWAAAAOygXKWuRo0aysnJ0YIFC1SnTh1169ZN9913nz766KNSx585c0Zut1sej0fXXXedJGnevHnq16+f6tatq9jYWKWlpWn+/Pn+bT799FM1a9ZMERERF83j8Xj8RfPc0q9fvzLHu91uuVwu/5KQkFCe3QcAAKiyKnRN3VdffaVly5apb9++ys3NVceOHZWTk+Mfc88996hWrVqKiYnRI488oilTpqh///4qKirSggULNHToUP/YoUOHKicnR8XFxZIk0zQvOUvt2rW1devWEsucOXPKHJ+ZmSmPx+NfCgsLy7v7AAAAVVKF/qJEVFSUevfurd69e+uBBx7Q7bffrqysLI0YMUKSNGHCBI0YMUK1atVSXFycHA6HJGnlypX68ssvS1xDJ0lFRUVavXq1evfurVatWundd9/V2bNnLzpbFxYWphYtWlxybsMwZBhG+XYWAAAgBATke+ratGmjU6dO+R/Xr19fLVq0UHx8vL/QSd/fIHHrrbeeN7t26623au7cuZKk3/72tzp58qSefvrpUt/rhzdKAAAA4Hvlmqk7evSobr75Zt12221q166dateurQ8++EBTp07Vr3/96wtue/jwYb322mtatmyZkpOTSzw3fPhw3XjjjTp27Ji6dOmiiRMnavz48fryyy914403qnHjxvrss880a9Ysde/e3X9XrGmaOnDgwHnv1bBhQ4WF8b3KAACg+ihXqatVq5a6dOmi6dOnKz8/X2fPnlVCQoLuuOMO3XfffRfcduHChapZs6auv/768567/vrrFR0drUWLFumuu+7SI488ok6dOmnmzJmaNWuWiouL1bx5cw0aNEjp6en+7bxerxo1anTe6+3fv5+vOgEAANWKwyzPnQk24/V6v78LNuNFhRkxVscBAABVUMGU/kF5n3O9xOPxyOl0lnv7Ct0oYTfbs9Mq9I8HAABQVXDhGQAAgA1Q6gAAAGyAUgcAAGADXFMnKTlrJTdKAKh0wbrYGkD1xEwdAACADVDqAAAAbIBSBwAAYAOUOgAAABsI6VJXXFwst9utyy67TNHR0UpJSdFLL71kdSwAAICgC+m7X91utxYtWqRZs2apZcuWeueddzR06FA1aNBAPXv2tDoeAABA0IRsqfP5fHr44Yf11ltvqWvXrpKkZs2a6d1339U///nPUkudz+eTz+fzP/Z6vUHLCwAAUJlCttR99tlnOn36tHr37l1i/ZkzZ9ShQ4dSt3G73crOzg5GPAAAgKAK2VJ38uRJSdLy5cv1s5/9rMRzhmGUuk1mZqbGjRvnf+z1epWQkFB5IQEAAIIkZEtdmzZtZBiG9u3bd8nXzxmGUWbhAwAACGUhW+pq166tv/zlL7r77rtVXFys7t27y+PxaP369XI6nUpPT7c6IgAAQNCEbKmTpL/+9a9q0KCB3G63Pv/8c9WpU0cdO3bUfffdZ3U0AACAoArpUudwODR27FiNHTvW6igAAACWCukvHwYAAMD3QnqmLlC2Z6fJ6XRaHQMAAKDCmKkDAACwAUodAACADVDqAAAAbIBr6iQlZ61UmBFjdQwgpBRM6W91BADADzBTBwAAYAOUOgAAABuwTalLTU1VRkaG1TEAAAAsYZtSBwAAUJ3ZotSNGDFCa9eu1RNPPCGHwyGHw6GCggKrYwEAAASNLe5+feKJJ7R7924lJydr8uTJkqQGDRqcN87n88nn8/kfe73eoGUEAACoTLaYqXO5XIqMjFRMTIzi4+MVHx+v8PDw88a53W65XC7/kpCQYEFaAACAwLNFqbtUmZmZ8ng8/qWwsNDqSAAAAAFhi9Ovl8owDBmGYXUMAACAgLPNTF1kZKSKioqsjgEAAGAJ25S6xMREbdy4UQUFBTpy5IiKi4utjgQAABA0til1f/nLXxQeHq42bdqoQYMG2rdvn9WRAAAAgsY219S1atVKGzZssDoGAACAJWwzUwcAAFCd2Wam7qfYnp0mp9NpdQwAAIAKY6YOAADABih1AAAANkCpAwAAsAGuqZOUnLVSYUaM1TGACiuY0t/qCAAAizFTBwAAYAOUOgAAABug1AEAANgApQ4AAMAGQrrU+Xw+3XXXXWrYsKGioqLUvXt3vf/++1bHAgAACLqQLnUTJ07UkiVLtGDBAm3ZskUtWrRQWlqajh07Vup4n88nr9dbYgEAALCDkC11p06d0j/+8Q89+uij6tevn9q0aaPZs2crOjpac+fOLXUbt9stl8vlXxISEoKcGgAAoHKEbKnLz8/X2bNn1a1bN/+6iIgIde7cWR9//HGp22RmZsrj8fiXwsLCYMUFAACoVNXqy4cNw5BhGFbHAAAACLiQnalr3ry5IiMjtX79ev+6s2fP6v3331ebNm0sTAYAABB8ITtTV7NmTY0ePVoTJkxQbGysmjRpoqlTp+r06dP6/e9/b3U8AACAoArZUidJU6ZMUXFxsYYNG6YTJ07oyiuv1MqVK1W3bl2rowEAAARVSJe6qKgoPfnkk3ryySetjgIAAGCpkC51gbI9O01Op9PqGAAAABUWsjdKAAAA4H8odQAAADZAqQMAALABrqmTlJy1UmFGjNUxgHIrmNLf6ggAgCqCmToAAAAboNQBAADYAKUOAADABih1AAAANhDSpW7FihXq3r276tSpo3r16umXv/yl8vPzrY4FAAAQdCFd6k6dOqVx48bpgw8+0OrVqxUWFqYbb7xRxcXFVkcDAAAIqpD+SpObbrqpxON58+apQYMG2rlzp5KTk88b7/P55PP5/I+9Xm+lZwQAAAiGkJ6p+/TTTzVkyBA1a9ZMTqdTiYmJkqR9+/aVOt7tdsvlcvmXhISEIKYFAACoPCFd6gYMGKBjx45p9uzZ2rhxozZu3ChJOnPmTKnjMzMz5fF4/EthYWEw4wIAAFSakD39evToUe3atUuzZ89Wjx49JEnvvvvuBbcxDEOGYQQjHgAAQFCFbKmrW7eu6tWrp2eeeUaNGjXSvn37dO+991odCwAAwBIhe/o1LCxMzz//vDZv3qzk5GTdfffdevTRR62OBQAAYImQnamTpF69emnnzp0l1pmmaVEaAAAA64TsTB0AAAD+J6Rn6gJle3aanE6n1TEAAAAqjJk6AAAAG6DUAQAA2AClDgAAwAa4pk5SctZKhRkxVsdANVUwpb/VEQAANsBMHQAAgA2EbKlLTU1VRkaG1TEAAACqhJAtdQAAAPgfSh0AAIANhHSpKy4u1sSJExUbG6v4+Hg9+OCDVkcCAACwREiXugULFqhmzZrauHGjpk6dqsmTJ2vVqlVljvf5fPJ6vSUWAAAAOwjpUteuXTtlZWWpZcuWGj58uK688kqtXr26zPFut1sul8u/JCQkBDEtAABA5Qn5UvdDjRo10qFDh8ocn5mZKY/H418KCwsrOyIAAEBQhPSXD0dERJR47HA4VFxcXOZ4wzBkGEZlxwIAAAi6kJ6pAwAAwPcodQAAADZAqQMAALCBkL2mLjc397x1S5cuDXoOAACAqoCZOgAAABsI2Zm6QNqenSan02l1DAAAgApjpg4AAMAGKHUAAAA2QKkDAACwAa6pk5SctVJhRozVMWBjBVP6Wx0BAGBzzNQBAADYAKUOAADABih1AAAANkCpAwAAsIGQvVEiNTVV7dq1U1RUlObMmaPIyEj98Y9/1IMPPmh1NAAAgKAL6Zm6BQsWqGbNmtq4caOmTp2qyZMna9WqVWWO9/l88nq9JRYAAAA7COlS165dO2VlZally5YaPny4rrzySq1evbrM8W63Wy6Xy78kJCQEMS0AAEDlCflS90ONGjXSoUOHyhyfmZkpj8fjXwoLCys7IgAAQFCE7DV1khQREVHiscPhUHFxcZnjDcOQYRiVHQsAACDoQnqmDgAAAN+j1AEAANgApQ4AAMAGQvaautzc3PPWLV26NOg5AAAAqoKQLXWBtD07TU6n0+oYAAAAFcbpVwAAABug1AEAANgApQ4AAMAGuKZOUnLWSoUZMVbHQBVWMKW/1REAALggZuoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2AClDgAAwAZsUepeeukltW3bVtHR0apXr5569eqlU6dOWR0LAAAgaEL+K03279+vIUOGaOrUqbrxxht14sQJrVu3TqZpnjfW5/PJ5/P5H3u93mBGBQAAqDS2KHXfffedfvOb36hp06aSpLZt25Y61u12Kzs7O5jxAAAAgiLkT7+mpKTo+uuvV9u2bXXzzTdr9uzZOn78eKljMzMz5fF4/EthYWGQ0wIAAFSOkC914eHhWrVqld544w21adNGf//735WUlKQ9e/acN9YwDDmdzhILAACAHYR8qZMkh8Ohbt26KTs7W3l5eYqMjNQrr7xidSwAAICgCflr6jZu3KjVq1erT58+atiwoTZu3KjDhw+rdevWVkcDAAAImpAvdU6nU++8845mzJghr9erpk2batq0aerXr5/V0QAAAIIm5Etd69attWLFCqtjAAAAWCrkS10gbM9O46YJAAAQ0mxxowQAAEB1R6kDAACwAUodAACADXBNnaTkrJUKM2KsjoEqpmBKf6sjAABwyZipAwAAsAFblbrU1FRlZGRYHQMAACDobHX69eWXX1ZERITVMQAAAILOVqUuNjbW6ggAAACW4PQrAACADdhqpu5ifD6ffD6f/7HX67UwDQAAQODYaqbuYtxut1wul39JSEiwOhIAAEBAVKtSl5mZKY/H418KCwutjgQAABAQ1er0q2EYMgzD6hgAAAABV61m6gAAAOyKUgcAAGADlDoAAAAbsNU1dbm5uVZHAAAAsAQzdQAAADZgq5m6itqenSan02l1DAAAgApjpg4AAMAGKHUAAAA2QKkDAACwAa6pk5SctVJhRozVMWChgin9rY4AAMBPwkwdAACADVDqAAAAbIBSBwAAYAO2LXVnzpyxOgIAAEDQ2OZGidTUVCUnJ6tGjRpatGiR2rZtqzVr1lgdCwAAIChsU+okacGCBRo9erTWr19f6vM+n08+n8//2Ov1BisaAABApbJVqWvZsqWmTp1a5vNut1vZ2dlBTAQAABActrqmrlOnThd8PjMzUx6Px78UFhYGKRkAAEDlstVMXc2aNS/4vGEYMgwjSGkAAACCx1YzdQAAANUVpQ4AAMAGKHUAAAA2YJtr6nJzc62OAAAAYBnblLqfYnt2mpxOp9UxAAAAKozTrwAAADZAqQMAALABSh0AAIANcE2dpOSslQozYqyOgf9TMKW/1REAAAg5zNQBAADYAKUOAADABmxX6lJTU5WRkWF1DAAAgKCyXakDAACojih1AAAANhDSpe7UqVMaPny4atWqpUaNGmnatGlWRwIAALBESJe6CRMmaO3atXr11Vf15ptvKjc3V1u2bClzvM/nk9frLbEAAADYQciWupMnT2ru3Ll67LHHdP3116tt27ZasGCBvvvuuzK3cbvdcrlc/iUhISGIiQEAACpPyJa6/Px8nTlzRl26dPGvi42NVVJSUpnbZGZmyuPx+JfCwsJgRAUAAKh01eovShiGIcMwrI4BAAAQcCE7U9e8eXNFRERo48aN/nXHjx/X7t27LUwFAABgjZCdqatVq5Z+//vfa8KECapXr54aNmyo+++/X2FhIdtTAQAAKixkS50kPfroozp58qQGDBig2rVra/z48fJ4PFbHAgAACDqHaZqm1SGs4vV6v78LNuNFhRkxVsfB/ymY0t/qCAAABN25XuLxeOR0Osu9fUjP1AXK9uy0Cv3jAQAAVBVcgAYAAGADlDoAAAAboNQBAADYANfUSUrOWsmNEkHCTRAAAFQOZuoAAABswFalLjU1VRkZGVbHAAAACDpblToAAIDqilIHAABgA5Q6AAAAG6hWd7/6fD75fD7/Y6/Xa2EaAACAwKlWM3Vut1sul8u/JCQkWB0JAAAgIKpVqcvMzJTH4/EvhYWFVkcCAAAIiGp1+tUwDBmGYXUMAACAgKtWM3UAAAB2RakDAACwAUodAACADdjqmrrc3FyrIwAAAFiCmToAAAAbsNVMXUVtz06T0+m0OgYAAECFMVMHAABgA5Q6AAAAG6DUAQAA2ADX1ElKzlqpMCPG6hi2UTClv9URAACodpipAwAAsAFKHQAAgA1UWqmbOXOmEhMTFRUVpS5dumjTpk2Vsr3b7VZ4eLgeffTRQMQGAAAISZVS6l544QWNGzdOWVlZ2rJli1JSUpSWlqZDhw4FfPt58+Zp4sSJmjdvXqB3AwAAIGRUSql7/PHHdccdd2jkyJFq06aNZs2apZiYGM2bN0+5ubmKjIzUunXr/OOnTp2qhg0b6uDBgxfd/ofWrl2rb775RpMnT5bX69V7771XGbsDAABQ5QW81J05c0abN29Wr169/vcmYWHq1auXNmzYoNTUVGVkZGjYsGHyeDzKy8vTAw88oDlz5iguLu6i2//Q3LlzNWTIEEVERGjIkCGaO3fuBbP5fD55vd4SCwAAgB0EvNQdOXJERUVFiouLK7E+Li5OBw4ckCQ99NBDqlu3rkaNGqWhQ4cqPT1dv/rVry55e0nyer166aWXNHToUEnS0KFD9eKLL+rkyZNlZnO73XK5XP4lISEhIPsMAABgNUvufo2MjNRzzz2nJUuW6Ntvv9X06dPL/Rr//ve/1bx5c6WkpEiS2rdvr6ZNm+qFF14oc5vMzEx5PB7/UlhYWOF9AAAAqEoCXurq16+v8PBw//Vx5xw8eFDx8fH+x+eufzt27JiOHTtW7u3nzp2rHTt2qEaNGv5l586dF7xhwjAMOZ3OEgsAAIAdBLzURUZGqlOnTlq9erV/XXFxsVavXq2uXbtKkvLz83X33Xdr9uzZ6tKli9LT01VcXHzJ22/btk0ffPCBcnNztXXrVv+Sm5urDRs26JNPPgn0bgEAAFRplfJnwsaNG6f09HRdeeWV6ty5s2bMmKFTp05p5MiRKioq0tChQ5WWlqaRI0eqb9++atu2raZNm6YJEyZcdHvp+1m6zp0765prrjnvva+66irNnTuX760DAADVSqWUusGDB+vw4cOaNGmSDhw4oPbt22vFihWKi4vT5MmTtXfvXr3++uuSpEaNGumZZ57RkCFD1KdPH6WkpFxw+zNnzmjRokW65557Sn3vm266SdOmTdPDDz+siIiIytg9AACAKsdhmqZpdQireL3e7++CzXhRYUaM1XFso2BKf6sjAAAQcs71Eo/HU6Hr/vnbrwAAADZQKadfQ8327DTuhAUAACGNmToAAAAboNQBAADYAKdfJSVnreRGiQDiRgkAAIKPmToAAAAboNQBAADYAKUOAADABgJe6mbOnKnExERFRUWpS5cu2rRpU0C3T0xMlMPhkMPhUHR0tBITE3XLLbfo7bffDuRuAAAAhJSAlroXXnhB48aNU1ZWlrZs2aKUlBSlpaXp0KFDAd1+8uTJ2r9/v3bt2qWFCxeqTp066tWrl/72t78FcncAAABCRkBL3eOPP6477rhDI0eOVJs2bTRr1izFxMRo3rx5ys3NVWRkpNatW+cfP3XqVDVs2FAHDx686PY/VLt2bcXHx6tJkya65ppr9Mwzz+iBBx7QpEmTtGvXrkDuEgAAQEgIWKk7c+aMNm/erF69ev3vxcPC1KtXL23YsEGpqanKyMjQsGHD5PF4lJeXpwceeEBz5sxRXFzcRbe/mLFjx8o0Tb366qtljvH5fPJ6vSUWAAAAOwhYqTty5IiKiooUFxdXYn1cXJwOHDggSXrooYdUt25djRo1SkOHDlV6erp+9atfXfL2FxIbG6uGDRuqoKCgzDFut1sul8u/JCQklHMvAQAAqqag3v0aGRmp5557TkuWLNG3336r6dOnB/T1TdOUw+Eo8/nMzEx5PB7/UlhYGND3BwAAsErA/qJE/fr1FR4e7r8+7pyDBw8qPj7e//i9996TJB07dkzHjh1TzZo1y7V9WY4eParDhw/rsssuK3OMYRgyDOOS9wkAACBUBGymLjIyUp06ddLq1av964qLi7V69Wp17dpVkpSfn6+7775bs2fPVpcuXZSenq7i4uJL3v5CnnjiCYWFhWngwIGB2iUAAICQEdC//Tpu3Dilp6fryiuvVOfOnTVjxgydOnVKI0eOVFFRkYYOHaq0tDSNHDlSffv2Vdu2bTVt2jRNmDDhotv/0IkTJ3TgwAGdPXtWe/bs0aJFizRnzhy53W61aNEikLsEAAAQEgJa6gYPHqzDhw9r0qRJOnDggNq3b68VK1YoLi5OkydP1t69e/X6669Lkho1aqRnnnlGQ4YMUZ8+fZSSknLB7X9o0qRJmjRpkiIjIxUfH69f/OIXWr16ta699tpA7g4AAEDIcJimaVodwiper/f7u2AzXlSYEWN1HNsomNLf6ggAAIScc73E4/HI6XSWe/uAztSFqu3ZaRX6xwMAAKgqgvqVJgAAAKgclDoAAAAboNQBAADYANfUSUrOWsmNEv+HmxwAAAhNzNQBAADYAKUOAADABih1AAAANkCpAwAAsIGQLXWJiYmaMWNGiXXt27fXgw8+aEkeAAAAK1Wru199Pp98Pp//sdfrtTANAABA4ITsTF1FuN1uuVwu/5KQkGB1JAAAgICoVqUuMzNTHo/HvxQWFlodCQAAICBC9vRrWFiYTNMsse7s2bMX3MYwDBmGUZmxAAAALBGyM3UNGjTQ/v37/Y+9Xq/27NljYSIAAADrhGypu+666/Tss89q3bp12rZtm9LT0xUeHm51LAAAAEuE7OnXzMxM7dmzR7/85S/lcrn017/+lZk6AABQbYVsqXM6nXr++edLrEtPT7coDQAAgLVC9vQrAAAA/idkZ+oCaXt2mpxOp9UxAAAAKoyZOgAAABug1AEAANgApQ4AAMAGuKZOUnLWSoUZMVbHqDQFU/pbHQEAAFQyZuoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2EDIlLrU1FT9+c9/VkZGhurWrau4uDjNnj1bp06d0siRI1W7dm21aNFCb7zxhtVRAQAAgi5kSp0kLViwQPXr19emTZv05z//WaNHj9bNN9+sq6++Wlu2bFGfPn00bNgwnT59utTtfT6fvF5viQUAAMAOQqrUpaSk6P/9v/+nli1bKjMzU1FRUapfv77uuOMOtWzZUpMmTdLRo0f10Ucflbq92+2Wy+XyLwkJCUHeAwAAgMoRUqWuXbt2/v8ODw9XvXr11LZtW/+6uLg4SdKhQ4dK3T4zM1Mej8e/FBYWVm5gAACAIAmpLx+OiIgo8djhcJRY53A4JEnFxcWlbm8YhgzDqLyAAAAAFgmpmToAAACUjlIHAABgA5Q6AAAAGwiZa+pyc3PPW1dQUHDeOtM0Kz8MAABAFcNMHQAAgA2EzExdZdqenSan02l1DAAAgApjpg4AAMAGKHUAAAA2wOlXSclZKxVmxFgd4ycpmNLf6ggAAMBCzNQBAADYAKUOAADABmxT6kzT1KhRoxQbGyuHw6GtW7daHQkAACBobHNN3YoVK5STk6Pc3Fw1a9ZM9evXtzoSAABA0Nim1OXn56tRo0a6+uqrrY4CAAAQdLYodSNGjNCCBQskSQ6HQ02bNi31T4gBAADYlS1K3RNPPKHmzZvrmWee0fvvv6/w8PBSx/l8Pvl8Pv9jr9cbrIgAAACVyhY3SrhcLtWuXVvh4eGKj49XgwYNSh3ndrvlcrn8S0JCQpCTAgAAVA5blLpLlZmZKY/H418KCwutjgQAABAQtjj9eqkMw5BhGFbHAAAACLhqNVMHAABgV5Q6AAAAG6DUAQAA2IBtSl1GRgbfTQcAAKqtanWjRFm2Z6fJ6XRaHQMAAKDCbDNTBwAAUJ1R6gAAAGyAUgcAAGADXFMnKTlrpcKMGKtj/CQFU/pbHQEAAFiImToAAAAboNQBAADYAKUOAADABih1AAAANkCpAwAAsIFqdferz+eTz+fzP/Z6vRamAQAACJxqNVPndrvlcrn8S0JCgtWRAAAAAqJalbrMzEx5PB7/UlhYaHUkAACAgKhWp18Nw5BhGFbHAAAACLhqNVMHAABgV5Q6AAAAG7BVqcvJyZHD4bA6BgAAQNDZqtTt2bNHPXv2tDoGAABA0NnqRok33nhDTz31lNUxAAAAgs5hmqZpdQireL1euVwueTweOZ1Oq+MAAIBq7Kf2EludfgUAAKiuKHUAAAA2QKkDAACwAVvdKFFRyVkrFWbEWB2jwgqm9Lc6AgAAsBgzdQAAADZAqQMAALABSh0AAIANUOoAAABsIKRK3euvv646deqoqKhIkrR161Y5HA7de++9/jG33367hg4dalVEAAAAS4RUqevRo4dOnDihvLw8SdLatWtVv3595ebm+sesXbtWqamppW7v8/nk9XpLLAAAAHYQUqXO5XKpffv2/hKXm5uru+++W3l5eTp58qS+/PJLffbZZ+rZs2ep27vdbrlcLv+SkJAQxPQAAACVJ6RKnST17NlTubm5Mk1T69at029+8xu1bt1a7777rtauXavGjRurZcuWpW6bmZkpj8fjXwoLC4OcHgAAoHKE3JcPp6amat68efrwww8VERGhyy+/XKmpqcrNzdXx48fLnKWTJMMwZBhGENMCAAAER8jN1J27rm769On+Aneu1OXm5pZ5PR0AAICdhVypq1u3rtq1a6fnnnvOX+CuueYabdmyRbt3777gTB0AAIBdhVypk76/rq6oqMhf6mJjY9WmTRvFx8crKSnJ2nAAAAAWCMlSN2PGDJmmqcsvv9y/buvWrdq/f7+FqQAAAKwTkqUOAAAAJYXc3a+VYXt2mpxOp9UxAAAAKoyZOgAAABug1AEAANgAp18lJWetVJgRY3WMCiuY0t/qCAAAwGLM1AEAANgApQ4AAMAGKHUAAAA2QKkDAACwAUodAACADYRsqVu4cKHq1asnn89XYv3AgQM1bNgwi1IBAABYI2RL3c0336yioiItW7bMv+7QoUNavny5brvttlK38fl88nq9JRYAAAA7CNlSFx0drd/+9reaP3++f92iRYvUpEkTpaamlrqN2+2Wy+XyLwkJCUFKCwAAULlCttRJ0h133KE333xTX375pSQpJydHI0aMkMPhKHV8ZmamPB6PfyksLAxmXAAAgEoT0n9RokOHDkpJSdHChQvVp08f7dixQ8uXLy9zvGEYMgwjiAkBAACCI6RLnSTdfvvtmjFjhr788kv16tWLU6oAAKBaCunTr5L029/+Vl988YVmz55d5g0SAAAAdhfypc7lcummm25SrVq1NHDgQKvjAAAAWCLkS50kffnll/rd737H9XIAAKDaCulr6o4fP67c3Fzl5ubq6aefrvDrbM9Ok9PpDGAyAACA4ArpUtehQwcdP35cjzzyiJKSkqyOAwAAYJmQLnUFBQVWRwAAAKgSbHFNHQAAQHVHqQMAALABSh0AAIANUOoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2AClDgAAwAYodQAAADZAqQMAALABSh0AAIANUOoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2EANqwNYyTRNSZLX67U4CQAAqO7O9ZFz/aS8qnWpO3r0qCQpISHB4iQAAADfO3HihFwuV7m3q9alLjY2VpK0b9++Cv3jIfi8Xq8SEhJUWFgop9NpdRxcAo5Z6OGYhR6OWegp7ZiZpqkTJ06ocePGFXrNal3qwsK+v6TQ5XLxQxBinE4nxyzEcMxCD8cs9HDMQs+Pj9lPmWTiRgkAAAAboNQBAADYQLUudYZhKCsrS4ZhWB0Fl4hjFno4ZqGHYxZ6OGahpzKOmcOs6H2zAAAAqDKq9UwdAACAXVDqAAAAbIBSBwAAYAOUOgAAABuwfambOXOmEhMTFRUVpS5dumjTpk0XHL948WJdfvnlioqKUtu2bfWf//wnSElxTnmOWU5OjhwOR4klKioqiGmrt3feeUcDBgxQ48aN5XA4tHTp0otuk5ubq44dO8owDLVo0UI5OTmVnhP/U95jlpube97PmMPh0IEDB4ITGHK73brqqqtUu3ZtNWzYUAMHDtSuXbsuuh2fZ9apyDELxOeZrUvdCy+8oHHjxikrK0tbtmxRSkqK0tLSdOjQoVLHv/feexoyZIh+//vfKy8vTwMHDtTAgQO1ffv2ICevvsp7zKTvv417//79/mXv3r1BTFy9nTp1SikpKZo5c+Yljd+zZ4/69++va6+9Vlu3blVGRoZuv/12rVy5spKT4pzyHrNzdu3aVeLnrGHDhpWUED+2du1a3Xnnnfrvf/+rVatW6ezZs+rTp49OnTpV5jZ8nlmrIsdMCsDnmWljnTt3Nu+8807/46KiIrNx48am2+0udfwtt9xi9u/fv8S6Ll26mH/4wx8qNSf+p7zHbP78+abL5QpSOlyIJPOVV1654JiJEyeaV1xxRYl1gwcPNtPS0ioxGcpyKcdszZo1piTz+PHjQcmEizt06JApyVy7dm2ZY/g8q1ou5ZgF4vPMtjN1Z86c0ebNm9WrVy//urCwMPXq1UsbNmwodZsNGzaUGC9JaWlpZY5HYFXkmEnSyZMn1bRpUyUkJOjXv/61duzYEYy4qAB+xkJX+/bt1ahRI/Xu3Vvr16+3Ok615vF4JEmxsbFljuFnrWq5lGMm/fTPM9uWuiNHjqioqEhxcXEl1sfFxZV5LciBAwfKNR6BVZFjlpSUpHnz5unVV1/VokWLVFxcrKuvvlpffPFFMCKjnMr6GfN6vfrmm28sSoULadSokWbNmqUlS5ZoyZIlSkhIUGpqqrZs2WJ1tGqpuLhYGRkZ6tatm5KTk8scx+dZ1XGpxywQn2c1AhEYsErXrl3VtWtX/+Orr75arVu31j//+U/99a9/tTAZYA9JSUlKSkryP7766quVn5+v6dOn69lnn7UwWfV05513avv27Xr33XetjoJLdKnHLBCfZ7adqatfv77Cw8N18ODBEusPHjyo+Pj4UreJj48v13gEVkWO2Y9FRESoQ4cO+uyzzyojIn6isn7GnE6noqOjLUqF8urcuTM/YxYYM2aMXn/9da1Zs0Y///nPLziWz7OqoTzH7Mcq8nlm21IXGRmpTp06afXq1f51xcXFWr16dYkm/ENdu3YtMV6SVq1aVeZ4BFZFjtmPFRUVadu2bWrUqFFlxcRPwM+YPWzdupWfsSAyTVNjxozRK6+8orfffluXXXbZRbfhZ81aFTlmP1ahz7OfdJtFFff888+bhmGYOTk55s6dO81Ro0aZderUMQ8cOGCapmkOGzbMvPfee/3j169fb9aoUcN87LHHzI8//tjMysoyIyIizG3btlm1C9VOeY9Zdna2uXLlSjM/P9/cvHmzeeutt5pRUVHmjh07rNqFauXEiRNmXl6emZeXZ0oyH3/8cTMvL8/cu3evaZqmee+995rDhg3zj//888/NmJgYc8KECebHH39szpw50wwPDzdXrFhh1S5UO+U9ZtOnTzeXLl1qfvrpp+a2bdvMsWPHmmFhYeZbb71l1S5UO6NHjzZdLpeZm5tr7t+/37+cPn3aP4bPs6qlIscsEJ9nti51pmmaf//7380mTZqYkZGRZufOnc3//ve//ud69uxppqenlxj/4osvmq1atTIjIyPNK664wly+fHmQE6M8xywjI8M/Ni4uzrzhhhvMLVu2WJC6ejr3dRc/Xs4do/T0dLNnz57nbdO+fXszMjLSbNasmTl//vyg567OynvMHnnkEbN58+ZmVFSUGRsba6ampppvv/22NeGrqdKOl6QSPzt8nlUtFTlmgfg8c/zfmwMAACCE2faaOgAAgOqEUgcAAGADlDoAAAAboNQBAADYAKUOAADABih1AAAANkCpAwAAsAFKHQAAwCV45513NGDAADVu3FgOh0NLly4t92uYpqnHHntMrVq1kmEY+tnPfqa//e1vAclHqQMACxUUFMjhcGjr1q1WRwFwEadOnVJKSopmzpxZ4dcYO3as5syZo8cee0yffPKJli1bps6dOwckX42AvAoAAIDN9evXT/369SvzeZ/Pp/vvv1///ve/9fXXXys5OVmPPPKIUlNTJUkff/yx/vGPf2j79u1KSkqSJF122WUBy8dMHYBqrbi4WFOnTlWLFi1kGIaaNGniPxWybds2XXfddYqOjla9evU0atQonTx50r9tamqqMjIySrzewIEDNWLECP/jxMREPfzww7rttttUu3ZtNWnSRM8884z/+XO/0Dt06CCHw+H/5Q8g9IwZM0YbNmzQ888/r48++kg333yz+vbtq08//VSS9Nprr6lZs2Z6/fXXddlllykxMVG33367jh07FpD3p9QBqNYyMzM1ZcoUPfDAA9q5c6f+9a9/KS4uTqdOnVJaWprq1q2r999/X4sXL9Zbb72lMWPGlPs9pk2bpiuvvFJ5eXn605/+pNGjR2vXrl2SpE2bNkmS3nrrLe3fv18vv/xyQPcPQHDs27dP8+fP1+LFi9WjRw81b95cf/nLX9S9e3fNnz9fkvT5559r7969Wrx4sRYuXKicnBxt3rxZgwYNCkgGTr8CqLZOnDihJ554Qk899ZTS09MlSc2bN1f37t01e/Zsffvtt1q4cKFq1qwpSXrqqac0YMAAPfLII4qLi7vk97nhhhv0pz/9SZJ0zz33aPr06VqzZo2SkpLUoEEDSVK9evUUHx8f4D0EECzbtm1TUVGRWrVqVWK9z+dTvXr1JH1/ZsDn82nhwoX+cXPnzlWnTp20a9cu/ynZiqLUAai2Pv74Y/l8Pl1//fWlPpeSkuIvdJLUrVs3FRcXa9euXeUqde3atfP/t8PhUHx8vA4dOvTTwgOoUk6ePKnw8HBt3rxZ4eHhJZ6rVauWJKlRo0aqUaNGieLXunVrSd/P9FHqAKCCoqOjf9L2YWFhMk2zxLqzZ8+eNy4iIqLEY4fDoeLi4p/03gCqlg4dOqioqEiHDh1Sjx49Sh3TrVs3fffdd8rPz1fz5s0lSbt375YkNW3a9Cdn4Jo6ANVWy5YtFR0drdWrV5/3XOvWrfXhhx/q1KlT/nXr169XWFiY//+mGzRooP379/ufLyoq0vbt28uVITIy0r8tgKrt5MmT2rp1q/8riPbs2aOtW7dq3759atWqlX73u99p+PDhevnll7Vnzx5t2rRJbrdby5cvlyT16tVLHTt21G233aa8vDxt3rxZf/jDH9S7d+/zTttWBKUOQLUVFRWle+65RxMnTtTChQuVn5+v//73v5o7d65+97vfKSoqSunp6dq+fbvWrFmjP//5zxo2bJj/1Ot1112n5cuXa/ny5frkk080evRoff311+XK0LBhQ0VHR2vFihU6ePCgPB5PJewpgED44IMP1KFDB3Xo0EGSNG7cOHXo0EGTJk2SJM2fP1/Dhw/X+PHjlZSUpIEDB+r9999XkyZNJH0/u//aa6+pfv36uuaaa9S/f3+1bt1azz//fEDycfoVQLX2wAMPqEaNGpo0aZK++uorNWrUSH/84x8VExOjlStXauzYsbrqqqsUExOjm266SY8//rh/29tuu00ffvihhg8frho1aujuu+/WtddeW673r1Gjhp588klNnjxZkyZNUo8ePZSbmxvgvQQQCKmpqeddcvFDERERys7OVnZ2dpljGjdurCVLllRGPDnMC6UDAABASOD0KwAAgA1Q6gAAAGyAUgcAAGADlDoAAAAboNQBAADYAKUOAADABih1AAAANkCpAwAAsAFKHQAAgA1Q6gAAAGyAUgcAAGADlDoAAAAb+P9EYpjpexBJsAAAAABJRU5ErkJggg==", - "text/plain": [ - "
" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Characters in dataset: 15.0 MB\n" - ] - } - ], + "outputs": [], "source": [ + "histogram[:] = 0\n", + "histogram_global[blocks, threads_per_block](values, histogram)\n", + "assert cp.sum(histogram) == len(values)\n", + "\n", "histogram_host = cp.asnumpy(histogram)\n", "\n", "# Print most frequently occurring characters.\n", @@ -256,7 +187,7 @@ "plt.title(\"Top 20 Bins\")\n", "plt.show()\n", "\n", - "print(f\"Characters in dataset: {values.size / 1e6:.1f} MB\")\n" + "print(f\"Characters in dataset: {values.size / 1e6:.1f} MB\")" ] }, { @@ -298,7 +229,7 @@ }, { "cell_type": "code", - "execution_count": 6, + "execution_count": null, "id": "8dbd226c-66f2-43df-868a-6b024b1de24c", "metadata": { "colab": { @@ -314,28 +245,16 @@ "id": "8dbd226c-66f2-43df-868a-6b024b1de24c", "outputId": "a082d13a-8e86-436d-c3de-8351f337d824" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "==PROF== Connected to process 311 (/usr/bin/python3.12)\n", - "==PROF== Profiling \"histogram_global[abi:v1,cw51cXTLSUwv1sDUaKthoaNgqamjgOR3W3Cw6igA9duC0hkwnNGiHEkYkrqQBFBTDEwCyQe21eqwcFW3UoDGjzpQEsgzrtUEAA_3d_3d]\": 0%....50%....100% - 44 passes\n", - "==PROF== Disconnected from process 311\n", - "==PROF== Report: /accelerated-computing-hub/tutorials/accelerated-python/notebooks/kernels/solutions/histogram_global.ncu-rep\n" - ] - } - ], + "outputs": [], "source": [ - "%%ncu -o histogram_global.ncu-rep\n", + "%%ncu -o histogram_global.ncu-rep --kernel-name regex:histogram_global\n", "histogram[:] = 0\n", - "histogram_global[blocks, threads_per_block](values, histogram)\n", - "cp.cuda.runtime.deviceSynchronize()" + "histogram_global[blocks, threads_per_block](values, histogram)" ] }, { "cell_type": "code", - "execution_count": 7, + "execution_count": null, "id": "ad12380e-253b-4410-ab34-9479411fdf81", "metadata": { "colab": { @@ -387,87 +306,7 @@ "id": "ad12380e-253b-4410-ab34-9479411fdf81", "outputId": "f2326771-7c13-46fa-a489-bf494d76cc1e" }, - "outputs": [ - { - "data": { - "text/html": [ - "\n", - " \n", - " " - ], - "text/plain": [ - "" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "049fca414b834d99a98ff0a9ebe39597", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Dropdown(description='Kernel:', layout=Layout(width='400px'), options=('histogram_global',), style=Description…" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "6f775d2efc07462f873cac04802d2a23", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Output()" - ] - }, - "metadata": {}, - "output_type": "display_data" - } - ], + "outputs": [], "source": [ "assert cp.sum(histogram) == len(values)" ] @@ -479,7 +318,7 @@ "id": "e1f72831-780f-4cf5-8ff1-2092ecb193d9" }, "source": [ - "### 5. Solution: Shared Memory & Cooperative Groups\n", + "### 5. Solution: Shared Memory\n", "\n", "We improved the code by separating loading values from the histogram update and to perform striped loads (also known as coalesced access) using [cuda.cooperative](https://nvidia.github.io/cccl/unstable/python/coop.html)'s block load instead of doing the I/O by hand.\n", "\n", @@ -492,7 +331,7 @@ }, { "cell_type": "code", - "execution_count": 8, + "execution_count": null, "id": "cf7c9865-646a-4bbd-9b41-61cadfc5484c", "metadata": { "colab": { @@ -508,31 +347,23 @@ "id": "cf7c9865-646a-4bbd-9b41-61cadfc5484c", "outputId": "851b231f-4431-4cf3-b857-c5a966a7162f" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Writing histogram_localized.py\n" - ] - } - ], + "outputs": [], "source": [ "import cuda.coop as coop\n", "\n", - "localized_items_per_thread = 32\n", - "localized_items_per_block = threads_per_block * localized_items_per_thread\n", - "localized_blocks = len(values) // localized_items_per_block\n", - "assert values.size % localized_items_per_block == 0\n", - "block_load = coop.block.load(cp.uint8, threads_per_block, localized_items_per_thread, 'striped')\n", + "items_per_thread = 8\n", + "items_per_block = threads_per_block * items_per_thread\n", + "blocks = len(values) // items_per_block\n", + "assert values.size % items_per_block == 0\n", + "block_load = coop.block.load(cp.uint8, threads_per_block, items_per_thread, 'striped')\n", "\n", "@cuda.jit(link=block_load.files)\n", "def histogram_localized(values, histogram):\n", - " items = cuda.local.array(localized_items_per_thread, dtype=values.dtype)\n", + " items = cuda.local.array(items_per_thread, dtype=values.dtype)\n", "\n", - " base = cuda.blockIdx.x * localized_items_per_block\n", + " base = cuda.blockIdx.x * items_per_block\n", "\n", - " block_load(values[base : base + localized_items_per_block], items)\n", + " block_load(values[base : base + items_per_block], items)\n", "\n", " local_histogram = cuda.shared.array(bins, dtype=histogram.dtype)\n", "\n", @@ -543,7 +374,7 @@ "\n", " cuda.syncthreads()\n", "\n", - " for i in range(localized_items_per_thread):\n", + " for i in range(items_per_thread):\n", " cuda.atomic.add(local_histogram, items[i], 1)\n", "\n", " cuda.syncthreads()\n", @@ -553,10 +384,8 @@ " if bin < histogram.size:\n", " cuda.atomic.add(histogram, bin, local_histogram[bin])\n", "\n", - "\n", "def launch_localized():\n", - " histogram[:] = 0\n", - " histogram_localized[localized_blocks, threads_per_block](values, histogram)\n" + " histogram_localized[blocks, threads_per_block](values, histogram)" ] }, { @@ -571,7 +400,7 @@ }, { "cell_type": "code", - "execution_count": 9, + "execution_count": null, "id": "1b30e9b3-5a4c-4181-b642-b7def5e9f258", "metadata": { "colab": { @@ -587,60 +416,12 @@ "id": "1b30e9b3-5a4c-4181-b642-b7def5e9f258", "outputId": "3a04d1d0-3409-441c-af2a-3138e9dec324" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "0.16 ms ± 8.47% (mean ± relative stdev of 15 runs)\n" - ] - } - ], + "outputs": [], "source": [ + "histogram[:] = 0\n", "launch_localized()\n", - "cp.cuda.runtime.deviceSynchronize()\n", - "assert cp.sum(histogram) == len(values)\n" - ] - }, - { - "cell_type": "code", - "execution_count": 10, - "id": "73f0c3cd-349b-490f-b6bd-7afbeb442fff", - "metadata": { - "colab": { - "base_uri": "https://localhost:8080/", - "height": 487 - }, - "execution": { - "iopub.execute_input": "2026-03-09T19:43:04.086115Z", - "iopub.status.busy": "2026-03-09T19:43:04.085849Z", - "iopub.status.idle": "2026-03-09T19:43:08.196008Z", - "shell.execute_reply": "2026-03-09T19:43:08.194609Z", - "shell.execute_reply.started": "2026-03-09T19:43:04.086089Z" - }, - "id": "73f0c3cd-349b-490f-b6bd-7afbeb442fff", - "outputId": "01324ca9-72c3-4317-a5be-f0bf144b7b7c" - }, - "outputs": [ - { - "data": { - "image/png": "iVBORw0KGgoAAAANSUhEUgAAAnUAAAHsCAYAAAC0dltCAAAAOnRFWHRTb2Z0d2FyZQBNYXRwbG90bGliIHZlcnNpb24zLjEwLjgsIGh0dHBzOi8vbWF0cGxvdGxpYi5vcmcvwVt1zgAAAAlwSFlzAAAPYQAAD2EBqD+naQAAOUVJREFUeJzt3Xt0FPXh/vFnE5JJAuxCuCTQBiK3CAbCRaEISFQgIKXFiiItELBKS6USoaDRn8RQ6yKKoBVL5RYQWxVRRKkgIkFECgpBuShoJBCVO7jLRRdM5veHX7ZGEiBxs5OdvF/nzDnu7Gd2n2FOso+fmdk4TNM0BQAAgJAWZnUAAAAA/HSUOgAAABug1AEAANgApQ4AAMAGKHUAAAA2QKkDAACwAUodAACADVDqAAAAbIBSBwAAYAOUOgAIATk5OXI4HCooKLA6CoAqilIHoEpwOByXtOTm5lZqjsLCQmVnZ6tz586qW7eu6tevr9TUVL311luljv/66681atQoNWjQQDVr1tS1116rLVu2XNJ7paamlti3yMhIXXbZZRo1apQKCwsDuVsAqgEHf/sVQFWwaNGiEo8XLlyoVatW6dlnny2xvnfv3oqLi6u0HE899ZQmTpyogQMHqlu3bvruu++0cOFCbdmyRfPmzdPIkSP9Y4uLi9WjRw99+OGHmjBhgurXr6+nn35ahYWF2rx5s1q2bHnB90pNTVV+fr7cbrck6cyZM9q5c6dmzZqlevXq6eOPP1ZMTIwkqaioSGfPnpVhGHI4HJW2/wBCF6UOQJU0ZswYzZw5U8H+FbVjxw7FxcWpfv36/nU+n0/t27fXyZMnS8ygvfjiixo8eLAWL16sQYMGSZIOHz6sVq1aqV+/fvrXv/51wfdKTU3VkSNHtH379hLrZ86cqTFjxujNN99U7969A7h3AOyM068AQsapU6c0fvx4JSQkyDAMJSUl6bHHHjuv+DkcDo0ZM0bPPfeckpKSFBUVpU6dOumdd9656HtcccUVJQqdJBmGoRtuuEFffPGFTpw44V//0ksvKS4uTr/5zW/86xo0aKBbbrlFr776qnw+X4X2Mz4+XpJUo0YN/7rSrqlLTEzUL3/5S7377rvq3LmzoqKi1KxZMy1cuLDE6509e1bZ2dlq2bKloqKiVK9ePXXv3l2rVq2qUD4AVROlDkBIME1Tv/rVrzR9+nT17dtXjz/+uJKSkjRhwgSNGzfuvPFr165VRkaGhg4dqsmTJ+vo0aPq27fvebNil+rAgQOKiYnxnw6VpLy8PHXs2FFhYSV/lXbu3FmnT5/W7t27L/q6RUVFOnLkiI4cOaL9+/fr7bffVlZWllq0aKFu3bpddPvPPvtMgwYNUu/evTVt2jTVrVtXI0aM0I4dO/xjHnzwQWVnZ+vaa6/VU089pfvvv19NmjS55Gv/AIQIEwCqoDvvvNP84a+opUuXmpLMhx56qMS4QYMGmQ6Hw/zss8/86ySZkswPPvjAv27v3r1mVFSUeeONN5Y7y6effmpGRUWZw4YNK7G+Zs2a5m233Xbe+OXLl5uSzBUrVlzwdXv27OnP+sOldevW5ueff15i7Pz5801J5p49e/zrmjZtakoy33nnHf+6Q4cOmYZhmOPHj/evS0lJMfv371+eXQYQgpipAxAS/vOf/yg8PFx33XVXifXjx4+XaZp64403Sqzv2rWrOnXq5H/cpEkT/frXv9bKlStVVFR0ye97+vRp3XzzzYqOjtaUKVNKPPfNN9/IMIzztomKivI/fzGJiYlatWqVVq1apTfeeEMzZsyQx+NRv379dPjw4Ytu36ZNG/Xo0cP/uEGDBkpKStLnn3/uX1enTh3t2LFDn3766UVfD0DootQBCAl79+5V48aNVbt27RLrW7du7X/+h0q787RVq1Y6ffr0JZUl6ftTo7feeqt27typl156SY0bNy7xfHR0dKnXzX377bf+5y+mZs2a6tWrl3r16qW+fftq7NixWrZsmXbt2nVeiSxNkyZNzltXt25dHT9+3P948uTJ+vrrr9WqVSu1bdtWEyZM0EcffXTR1wYQWih1AFCGO+64Q6+//rpycnJ03XXXnfd8o0aNtH///vPWn1v34xJ4qTp16iSXy3VJN3aEh4eXut78wc0j11xzjfLz8zVv3jwlJydrzpw56tixo+bMmVOhfACqJkodgJDQtGlTffXVVyXuPpWkTz75xP/8D5V2qnH37t2KiYlRgwYNLvp+EyZM0Pz58zV9+nQNGTKk1DHt27fXli1bVFxcXGL9xo0bFRMTo1atWl30fcpSVFSkkydPVnj7H4uNjdXIkSP173//W4WFhWrXrp0efPDBgL0+AOtR6gCEhBtuuEFFRUV66qmnSqyfPn26HA6H+vXrV2L9hg0bStzdWVhYqFdffVV9+vQpc3brnEcffVSPPfaY7rvvPo0dO7bMcYMGDdLBgwf18ssv+9cdOXJEixcv1oABA0q93u5SrFmzRidPnlRKSkqFtv+xo0ePlnhcq1YttWjRosJfuQKgaqpx8SEAYL0BAwbo2muv1f3336+CggKlpKTozTff1KuvvqqMjAw1b968xPjk5GSlpaXprrvukmEYevrppyVJ2dnZF3yfV155RRMnTlTLli3VunXr8/7SxQ//osWgQYP0i1/8QiNHjtTOnTv9f1GiqKjoou9zjsfj8b/Hd999p127dukf//iHoqOjde+9917Sa1xMmzZtlJqaqk6dOik2NlYffPCBXnrpJY0ZMyYgrw+gaqDUAQgJYWFhWrZsmSZNmqQXXnhB8+fPV2Jioh599FGNHz/+vPE9e/ZU165dlZ2drX379qlNmzbKyclRu3btLvg+H374oaTvT98OGzbsvOfXrFnjL3Xh4eH6z3/+owkTJujJJ5/UN998o6uuuko5OTlKSkq6pP364osv/O/jcDhUt25d9ezZU1lZWWrfvv0lvcbF3HXXXVq2bJnefPNN+Xw+NW3aVA899JAmTJgQkNcHUDXwZ8IA2I7D4dCdd9553qlaALAzrqkDAACwAUodAACADVDqAAAAbIAbJQDYDpcKA6iOmKkDAACwAUodAACADVTr06/FxcX66quvVLt2bTkcDqvjAACAasw0TZ04cUKNGzdWWFj5592qdan76quvlJCQYHUMAAAAv8LCQv385z8v93bVutTVrl1b0vf/eE6n0+I0AACgOvN6vUpISPD3k/Kq1qXu3ClXp9NJqQMAAFVCRS8J40YJAAAAG6DUAQAA2AClDgAAwAYodQAAADZAqQMAALABSh0AAIANUOoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2AClDgAAwAYodQAAADZAqQMAALABSh0AAIANUOoAAABsoIbVAaqC5KyVCjNirI4BAACqoIIp/a2OcEmYqQMAALABSh0AAIANUOoAAABsoNyl7vDhwxo9erSaNGkiwzAUHx+vtLQ0rV+/XpKUmJgoh8Mhh8OhmjVrqmPHjlq8eHGJ1/jmm28UGxur+vXry+fzlfo+S5YsUWpqqlwul2rVqqV27dpp8uTJOnbsmCQpJyfH/z4/XKKiosq7SwAAACGv3KXupptuUl5enhYsWKDdu3dr2bJlSk1N1dGjR/1jJk+erP379ysvL09XXXWVBg8erPfee8///JIlS3TFFVfo8ssv19KlS897j/vvv1+DBw/WVVddpTfeeEPbt2/XtGnT9OGHH+rZZ5/1j3M6ndq/f3+JZe/eveXdJQAAgJBXrrtfv/76a61bt065ubnq2bOnJKlp06bq3LlziXG1a9dWfHy84uPjNXPmTC1atEivvfaarr76aknS3LlzNXToUJmmqblz52rw4MH+bTdt2qSHH35YM2bM0NixY/3rExMT1bt3b3399df+dQ6HQ/Hx8eXeaQAAALsp10xdrVq1VKtWLS1durTM06Y/VqNGDUVEROjMmTOSpPz8fG3YsEG33HKLbrnlFq1bt67E7Npzzz2nWrVq6U9/+lOpr1enTp3yRC7B5/PJ6/WWWAAAAOygXKWuRo0aysnJ0YIFC1SnTh1169ZN9913nz766KNSx585c0Zut1sej0fXXXedJGnevHnq16+f6tatq9jYWKWlpWn+/Pn+bT799FM1a9ZMERERF83j8Xj8RfPc0q9fvzLHu91uuVwu/5KQkFCe3QcAAKiyKnRN3VdffaVly5apb9++ys3NVceOHZWTk+Mfc88996hWrVqKiYnRI488oilTpqh///4qKirSggULNHToUP/YoUOHKicnR8XFxZIk0zQvOUvt2rW1devWEsucOXPKHJ+ZmSmPx+NfCgsLy7v7AAAAVVKF/qJEVFSUevfurd69e+uBBx7Q7bffrqysLI0YMUKSNGHCBI0YMUK1atVSXFycHA6HJGnlypX68ssvS1xDJ0lFRUVavXq1evfurVatWundd9/V2bNnLzpbFxYWphYtWlxybsMwZBhG+XYWAAAgBATke+ratGmjU6dO+R/Xr19fLVq0UHx8vL/QSd/fIHHrrbeeN7t26623au7cuZKk3/72tzp58qSefvrpUt/rhzdKAAAA4Hvlmqk7evSobr75Zt12221q166dateurQ8++EBTp07Vr3/96wtue/jwYb322mtatmyZkpOTSzw3fPhw3XjjjTp27Ji6dOmiiRMnavz48fryyy914403qnHjxvrss880a9Ysde/e3X9XrGmaOnDgwHnv1bBhQ4WF8b3KAACg+ihXqatVq5a6dOmi6dOnKz8/X2fPnlVCQoLuuOMO3XfffRfcduHChapZs6auv/768567/vrrFR0drUWLFumuu+7SI488ok6dOmnmzJmaNWuWiouL1bx5cw0aNEjp6en+7bxerxo1anTe6+3fv5+vOgEAANWKwyzPnQk24/V6v78LNuNFhRkxVscBAABVUMGU/kF5n3O9xOPxyOl0lnv7Ct0oYTfbs9Mq9I8HAABQVXDhGQAAgA1Q6gAAAGyAUgcAAGADXFMnKTlrJTdKAKh0wbrYGkD1xEwdAACADVDqAAAAbIBSBwAAYAOUOgAAABsI6VJXXFwst9utyy67TNHR0UpJSdFLL71kdSwAAICgC+m7X91utxYtWqRZs2apZcuWeueddzR06FA1aNBAPXv2tDoeAABA0IRsqfP5fHr44Yf11ltvqWvXrpKkZs2a6d1339U///nPUkudz+eTz+fzP/Z6vUHLCwAAUJlCttR99tlnOn36tHr37l1i/ZkzZ9ShQ4dSt3G73crOzg5GPAAAgKAK2VJ38uRJSdLy5cv1s5/9rMRzhmGUuk1mZqbGjRvnf+z1epWQkFB5IQEAAIIkZEtdmzZtZBiG9u3bd8nXzxmGUWbhAwAACGUhW+pq166tv/zlL7r77rtVXFys7t27y+PxaP369XI6nUpPT7c6IgAAQNCEbKmTpL/+9a9q0KCB3G63Pv/8c9WpU0cdO3bUfffdZ3U0AACAoArpUudwODR27FiNHTvW6igAAACWCukvHwYAAMD3QnqmLlC2Z6fJ6XRaHQMAAKDCmKkDAACwAUodAACADVDqAAAAbIBr6iQlZ61UmBFjdQwgpBRM6W91BADADzBTBwAAYAOUOgAAABuwTalLTU1VRkaG1TEAAAAsYZtSBwAAUJ3ZotSNGDFCa9eu1RNPPCGHwyGHw6GCggKrYwEAAASNLe5+feKJJ7R7924lJydr8uTJkqQGDRqcN87n88nn8/kfe73eoGUEAACoTLaYqXO5XIqMjFRMTIzi4+MVHx+v8PDw88a53W65XC7/kpCQYEFaAACAwLNFqbtUmZmZ8ng8/qWwsNDqSAAAAAFhi9Ovl8owDBmGYXUMAACAgLPNTF1kZKSKioqsjgEAAGAJ25S6xMREbdy4UQUFBTpy5IiKi4utjgQAABA0til1f/nLXxQeHq42bdqoQYMG2rdvn9WRAAAAgsY219S1atVKGzZssDoGAACAJWwzUwcAAFCd2Wam7qfYnp0mp9NpdQwAAIAKY6YOAADABih1AAAANkCpAwAAsAGuqZOUnLVSYUaM1TGACiuY0t/qCAAAizFTBwAAYAOUOgAAABug1AEAANgApQ4AAMAGQrrU+Xw+3XXXXWrYsKGioqLUvXt3vf/++1bHAgAACLqQLnUTJ07UkiVLtGDBAm3ZskUtWrRQWlqajh07Vup4n88nr9dbYgEAALCDkC11p06d0j/+8Q89+uij6tevn9q0aaPZs2crOjpac+fOLXUbt9stl8vlXxISEoKcGgAAoHKEbKnLz8/X2bNn1a1bN/+6iIgIde7cWR9//HGp22RmZsrj8fiXwsLCYMUFAACoVNXqy4cNw5BhGFbHAAAACLiQnalr3ry5IiMjtX79ev+6s2fP6v3331ebNm0sTAYAABB8ITtTV7NmTY0ePVoTJkxQbGysmjRpoqlTp+r06dP6/e9/b3U8AACAoArZUidJU6ZMUXFxsYYNG6YTJ07oyiuv1MqVK1W3bl2rowEAAARVSJe6qKgoPfnkk3ryySetjgIAAGCpkC51gbI9O01Op9PqGAAAABUWsjdKAAAA4H8odQAAADZAqQMAALABrqmTlJy1UmFGjNUxgHIrmNLf6ggAgCqCmToAAAAboNQBAADYAKUOAADABih1AAAANhDSpW7FihXq3r276tSpo3r16umXv/yl8vPzrY4FAAAQdCFd6k6dOqVx48bpgw8+0OrVqxUWFqYbb7xRxcXFVkcDAAAIqpD+SpObbrqpxON58+apQYMG2rlzp5KTk88b7/P55PP5/I+9Xm+lZwQAAAiGkJ6p+/TTTzVkyBA1a9ZMTqdTiYmJkqR9+/aVOt7tdsvlcvmXhISEIKYFAACoPCFd6gYMGKBjx45p9uzZ2rhxozZu3ChJOnPmTKnjMzMz5fF4/EthYWEw4wIAAFSakD39evToUe3atUuzZ89Wjx49JEnvvvvuBbcxDEOGYQQjHgAAQFCFbKmrW7eu6tWrp2eeeUaNGjXSvn37dO+991odCwAAwBIhe/o1LCxMzz//vDZv3qzk5GTdfffdevTRR62OBQAAYImQnamTpF69emnnzp0l1pmmaVEaAAAA64TsTB0AAAD+J6Rn6gJle3aanE6n1TEAAAAqjJk6AAAAG6DUAQAA2AClDgAAwAa4pk5SctZKhRkxVsdANVUwpb/VEQAANsBMHQAAgA2EbKlLTU1VRkaG1TEAAACqhJAtdQAAAPgfSh0AAIANhHSpKy4u1sSJExUbG6v4+Hg9+OCDVkcCAACwREiXugULFqhmzZrauHGjpk6dqsmTJ2vVqlVljvf5fPJ6vSUWAAAAOwjpUteuXTtlZWWpZcuWGj58uK688kqtXr26zPFut1sul8u/JCQkBDEtAABA5Qn5UvdDjRo10qFDh8ocn5mZKY/H418KCwsrOyIAAEBQhPSXD0dERJR47HA4VFxcXOZ4wzBkGEZlxwIAAAi6kJ6pAwAAwPcodQAAADZAqQMAALCBkL2mLjc397x1S5cuDXoOAACAqoCZOgAAABsI2Zm6QNqenSan02l1DAAAgApjpg4AAMAGKHUAAAA2QKkDAACwAa6pk5SctVJhRozVMWBjBVP6Wx0BAGBzzNQBAADYAKUOAADABih1AAAANkCpAwAAsIGQvVEiNTVV7dq1U1RUlObMmaPIyEj98Y9/1IMPPmh1NAAAgKAL6Zm6BQsWqGbNmtq4caOmTp2qyZMna9WqVWWO9/l88nq9JRYAAAA7COlS165dO2VlZally5YaPny4rrzySq1evbrM8W63Wy6Xy78kJCQEMS0AAEDlCflS90ONGjXSoUOHyhyfmZkpj8fjXwoLCys7IgAAQFCE7DV1khQREVHiscPhUHFxcZnjDcOQYRiVHQsAACDoQnqmDgAAAN+j1AEAANgApQ4AAMAGQvaautzc3PPWLV26NOg5AAAAqoKQLXWBtD07TU6n0+oYAAAAFcbpVwAAABug1AEAANgApQ4AAMAGuKZOUnLWSoUZMVbHQBVWMKW/1REAALggZuoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2AClDgAAwAZsUepeeukltW3bVtHR0apXr5569eqlU6dOWR0LAAAgaEL+K03279+vIUOGaOrUqbrxxht14sQJrVu3TqZpnjfW5/PJ5/P5H3u93mBGBQAAqDS2KHXfffedfvOb36hp06aSpLZt25Y61u12Kzs7O5jxAAAAgiLkT7+mpKTo+uuvV9u2bXXzzTdr9uzZOn78eKljMzMz5fF4/EthYWGQ0wIAAFSOkC914eHhWrVqld544w21adNGf//735WUlKQ9e/acN9YwDDmdzhILAACAHYR8qZMkh8Ohbt26KTs7W3l5eYqMjNQrr7xidSwAAICgCflr6jZu3KjVq1erT58+atiwoTZu3KjDhw+rdevWVkcDAAAImpAvdU6nU++8845mzJghr9erpk2batq0aerXr5/V0QAAAIIm5Etd69attWLFCqtjAAAAWCrkS10gbM9O46YJAAAQ0mxxowQAAEB1R6kDAACwAUodAACADXBNnaTkrJUKM2KsjoEqpmBKf6sjAABwyZipAwAAsAFblbrU1FRlZGRYHQMAACDobHX69eWXX1ZERITVMQAAAILOVqUuNjbW6ggAAACW4PQrAACADdhqpu5ifD6ffD6f/7HX67UwDQAAQODYaqbuYtxut1wul39JSEiwOhIAAEBAVKtSl5mZKY/H418KCwutjgQAABAQ1er0q2EYMgzD6hgAAAABV61m6gAAAOyKUgcAAGADlDoAAAAbsNU1dbm5uVZHAAAAsAQzdQAAADZgq5m6itqenSan02l1DAAAgApjpg4AAMAGKHUAAAA2QKkDAACwAa6pk5SctVJhRozVMWChgin9rY4AAMBPwkwdAACADVDqAAAAbIBSBwAAYAO2LXVnzpyxOgIAAEDQ2OZGidTUVCUnJ6tGjRpatGiR2rZtqzVr1lgdCwAAIChsU+okacGCBRo9erTWr19f6vM+n08+n8//2Ov1BisaAABApbJVqWvZsqWmTp1a5vNut1vZ2dlBTAQAABActrqmrlOnThd8PjMzUx6Px78UFhYGKRkAAEDlstVMXc2aNS/4vGEYMgwjSGkAAACCx1YzdQAAANUVpQ4AAMAGKHUAAAA2YJtr6nJzc62OAAAAYBnblLqfYnt2mpxOp9UxAAAAKozTrwAAADZAqQMAALABSh0AAIANcE2dpOSslQozYqyOgf9TMKW/1REAAAg5zNQBAADYAKUOAADABmxX6lJTU5WRkWF1DAAAgKCyXakDAACojih1AAAANhDSpe7UqVMaPny4atWqpUaNGmnatGlWRwIAALBESJe6CRMmaO3atXr11Vf15ptvKjc3V1u2bClzvM/nk9frLbEAAADYQciWupMnT2ru3Ll67LHHdP3116tt27ZasGCBvvvuuzK3cbvdcrlc/iUhISGIiQEAACpPyJa6/Px8nTlzRl26dPGvi42NVVJSUpnbZGZmyuPx+JfCwsJgRAUAAKh01eovShiGIcMwrI4BAAAQcCE7U9e8eXNFRERo48aN/nXHjx/X7t27LUwFAABgjZCdqatVq5Z+//vfa8KECapXr54aNmyo+++/X2FhIdtTAQAAKixkS50kPfroozp58qQGDBig2rVra/z48fJ4PFbHAgAACDqHaZqm1SGs4vV6v78LNuNFhRkxVsfB/ymY0t/qCAAABN25XuLxeOR0Osu9fUjP1AXK9uy0Cv3jAQAAVBVcgAYAAGADlDoAAAAboNQBAADYANfUSUrOWsmNEkHCTRAAAFQOZuoAAABswFalLjU1VRkZGVbHAAAACDpblToAAIDqilIHAABgA5Q6AAAAG6hWd7/6fD75fD7/Y6/Xa2EaAACAwKlWM3Vut1sul8u/JCQkWB0JAAAgIKpVqcvMzJTH4/EvhYWFVkcCAAAIiGp1+tUwDBmGYXUMAACAgKtWM3UAAAB2RakDAACwAUodAACADdjqmrrc3FyrIwAAAFiCmToAAAAbsNVMXUVtz06T0+m0OgYAAECFMVMHAABgA5Q6AAAAG6DUAQAA2ADX1ElKzlqpMCPG6hi2UTClv9URAACodpipAwAAsAFKHQAAgA1UWqmbOXOmEhMTFRUVpS5dumjTpk2Vsr3b7VZ4eLgeffTRQMQGAAAISZVS6l544QWNGzdOWVlZ2rJli1JSUpSWlqZDhw4FfPt58+Zp4sSJmjdvXqB3AwAAIGRUSql7/PHHdccdd2jkyJFq06aNZs2apZiYGM2bN0+5ubmKjIzUunXr/OOnTp2qhg0b6uDBgxfd/ofWrl2rb775RpMnT5bX69V7771XGbsDAABQ5QW81J05c0abN29Wr169/vcmYWHq1auXNmzYoNTUVGVkZGjYsGHyeDzKy8vTAw88oDlz5iguLu6i2//Q3LlzNWTIEEVERGjIkCGaO3fuBbP5fD55vd4SCwAAgB0EvNQdOXJERUVFiouLK7E+Li5OBw4ckCQ99NBDqlu3rkaNGqWhQ4cqPT1dv/rVry55e0nyer166aWXNHToUEnS0KFD9eKLL+rkyZNlZnO73XK5XP4lISEhIPsMAABgNUvufo2MjNRzzz2nJUuW6Ntvv9X06dPL/Rr//ve/1bx5c6WkpEiS2rdvr6ZNm+qFF14oc5vMzEx5PB7/UlhYWOF9AAAAqEoCXurq16+v8PBw//Vx5xw8eFDx8fH+x+eufzt27JiOHTtW7u3nzp2rHTt2qEaNGv5l586dF7xhwjAMOZ3OEgsAAIAdBLzURUZGqlOnTlq9erV/XXFxsVavXq2uXbtKkvLz83X33Xdr9uzZ6tKli9LT01VcXHzJ22/btk0ffPCBcnNztXXrVv+Sm5urDRs26JNPPgn0bgEAAFRplfJnwsaNG6f09HRdeeWV6ty5s2bMmKFTp05p5MiRKioq0tChQ5WWlqaRI0eqb9++atu2raZNm6YJEyZcdHvp+1m6zp0765prrjnvva+66irNnTuX760DAADVSqWUusGDB+vw4cOaNGmSDhw4oPbt22vFihWKi4vT5MmTtXfvXr3++uuSpEaNGumZZ57RkCFD1KdPH6WkpFxw+zNnzmjRokW65557Sn3vm266SdOmTdPDDz+siIiIytg9AACAKsdhmqZpdQireL3e7++CzXhRYUaM1XFso2BKf6sjAAAQcs71Eo/HU6Hr/vnbrwAAADZQKadfQ8327DTuhAUAACGNmToAAAAboNQBAADYAKdfJSVnreRGiQDiRgkAAIKPmToAAAAboNQBAADYAKUOAADABgJe6mbOnKnExERFRUWpS5cu2rRpU0C3T0xMlMPhkMPhUHR0tBITE3XLLbfo7bffDuRuAAAAhJSAlroXXnhB48aNU1ZWlrZs2aKUlBSlpaXp0KFDAd1+8uTJ2r9/v3bt2qWFCxeqTp066tWrl/72t78FcncAAABCRkBL3eOPP6477rhDI0eOVJs2bTRr1izFxMRo3rx5ys3NVWRkpNatW+cfP3XqVDVs2FAHDx686PY/VLt2bcXHx6tJkya65ppr9Mwzz+iBBx7QpEmTtGvXrkDuEgAAQEgIWKk7c+aMNm/erF69ev3vxcPC1KtXL23YsEGpqanKyMjQsGHD5PF4lJeXpwceeEBz5sxRXFzcRbe/mLFjx8o0Tb366qtljvH5fPJ6vSUWAAAAOwhYqTty5IiKiooUFxdXYn1cXJwOHDggSXrooYdUt25djRo1SkOHDlV6erp+9atfXfL2FxIbG6uGDRuqoKCgzDFut1sul8u/JCQklHMvAQAAqqag3v0aGRmp5557TkuWLNG3336r6dOnB/T1TdOUw+Eo8/nMzEx5PB7/UlhYGND3BwAAsErA/qJE/fr1FR4e7r8+7pyDBw8qPj7e//i9996TJB07dkzHjh1TzZo1y7V9WY4eParDhw/rsssuK3OMYRgyDOOS9wkAACBUBGymLjIyUp06ddLq1av964qLi7V69Wp17dpVkpSfn6+7775bs2fPVpcuXZSenq7i4uJL3v5CnnjiCYWFhWngwIGB2iUAAICQEdC//Tpu3Dilp6fryiuvVOfOnTVjxgydOnVKI0eOVFFRkYYOHaq0tDSNHDlSffv2Vdu2bTVt2jRNmDDhotv/0IkTJ3TgwAGdPXtWe/bs0aJFizRnzhy53W61aNEikLsEAAAQEgJa6gYPHqzDhw9r0qRJOnDggNq3b68VK1YoLi5OkydP1t69e/X6669Lkho1aqRnnnlGQ4YMUZ8+fZSSknLB7X9o0qRJmjRpkiIjIxUfH69f/OIXWr16ta699tpA7g4AAEDIcJimaVodwiper/f7u2AzXlSYEWN1HNsomNLf6ggAAIScc73E4/HI6XSWe/uAztSFqu3ZaRX6xwMAAKgqgvqVJgAAAKgclDoAAAAboNQBAADYANfUSUrOWsmNEv+HmxwAAAhNzNQBAADYAKUOAADABih1AAAANkCpAwAAsIGQLXWJiYmaMWNGiXXt27fXgw8+aEkeAAAAK1Wru199Pp98Pp//sdfrtTANAABA4ITsTF1FuN1uuVwu/5KQkGB1JAAAgICoVqUuMzNTHo/HvxQWFlodCQAAICBC9vRrWFiYTNMsse7s2bMX3MYwDBmGUZmxAAAALBGyM3UNGjTQ/v37/Y+9Xq/27NljYSIAAADrhGypu+666/Tss89q3bp12rZtm9LT0xUeHm51LAAAAEuE7OnXzMxM7dmzR7/85S/lcrn017/+lZk6AABQbYVsqXM6nXr++edLrEtPT7coDQAAgLVC9vQrAAAA/idkZ+oCaXt2mpxOp9UxAAAAKoyZOgAAABug1AEAANgApQ4AAMAGuKZOUnLWSoUZMVbHqDQFU/pbHQEAAFQyZuoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2EDIlLrU1FT9+c9/VkZGhurWrau4uDjNnj1bp06d0siRI1W7dm21aNFCb7zxhtVRAQAAgi5kSp0kLViwQPXr19emTZv05z//WaNHj9bNN9+sq6++Wlu2bFGfPn00bNgwnT59utTtfT6fvF5viQUAAMAOQqrUpaSk6P/9v/+nli1bKjMzU1FRUapfv77uuOMOtWzZUpMmTdLRo0f10Ucflbq92+2Wy+XyLwkJCUHeAwAAgMoRUqWuXbt2/v8ODw9XvXr11LZtW/+6uLg4SdKhQ4dK3T4zM1Mej8e/FBYWVm5gAACAIAmpLx+OiIgo8djhcJRY53A4JEnFxcWlbm8YhgzDqLyAAAAAFgmpmToAAACUjlIHAABgA5Q6AAAAGwiZa+pyc3PPW1dQUHDeOtM0Kz8MAABAFcNMHQAAgA2EzExdZdqenSan02l1DAAAgApjpg4AAMAGKHUAAAA2wOlXSclZKxVmxFgd4ycpmNLf6ggAAMBCzNQBAADYAKUOAADABmxT6kzT1KhRoxQbGyuHw6GtW7daHQkAACBobHNN3YoVK5STk6Pc3Fw1a9ZM9evXtzoSAABA0Nim1OXn56tRo0a6+uqrrY4CAAAQdLYodSNGjNCCBQskSQ6HQ02bNi31T4gBAADYlS1K3RNPPKHmzZvrmWee0fvvv6/w8PBSx/l8Pvl8Pv9jr9cbrIgAAACVyhY3SrhcLtWuXVvh4eGKj49XgwYNSh3ndrvlcrn8S0JCQpCTAgAAVA5blLpLlZmZKY/H418KCwutjgQAABAQtjj9eqkMw5BhGFbHAAAACLhqNVMHAABgV5Q6AAAAG6DUAQAA2IBtSl1GRgbfTQcAAKqtanWjRFm2Z6fJ6XRaHQMAAKDCbDNTBwAAUJ1R6gAAAGyAUgcAAGADXFMnKTlrpcKMGKtj/CQFU/pbHQEAAFiImToAAAAboNQBAADYAKUOAADABih1AAAANkCpAwAAsIFqdferz+eTz+fzP/Z6vRamAQAACJxqNVPndrvlcrn8S0JCgtWRAAAAAqJalbrMzEx5PB7/UlhYaHUkAACAgKhWp18Nw5BhGFbHAAAACLhqNVMHAABgV5Q6AAAAG7BVqcvJyZHD4bA6BgAAQNDZqtTt2bNHPXv2tDoGAABA0NnqRok33nhDTz31lNUxAAAAgs5hmqZpdQireL1euVwueTweOZ1Oq+MAAIBq7Kf2EludfgUAAKiuKHUAAAA2QKkDAACwAVvdKFFRyVkrFWbEWB2jwgqm9Lc6AgAAsBgzdQAAADZAqQMAALABSh0AAIANUOoAAABsIKRK3euvv646deqoqKhIkrR161Y5HA7de++9/jG33367hg4dalVEAAAAS4RUqevRo4dOnDihvLw8SdLatWtVv3595ebm+sesXbtWqamppW7v8/nk9XpLLAAAAHYQUqXO5XKpffv2/hKXm5uru+++W3l5eTp58qS+/PJLffbZZ+rZs2ep27vdbrlcLv+SkJAQxPQAAACVJ6RKnST17NlTubm5Mk1T69at029+8xu1bt1a7777rtauXavGjRurZcuWpW6bmZkpj8fjXwoLC4OcHgAAoHKE3JcPp6amat68efrwww8VERGhyy+/XKmpqcrNzdXx48fLnKWTJMMwZBhGENMCAAAER8jN1J27rm769On+Aneu1OXm5pZ5PR0AAICdhVypq1u3rtq1a6fnnnvOX+CuueYabdmyRbt3777gTB0AAIBdhVypk76/rq6oqMhf6mJjY9WmTRvFx8crKSnJ2nAAAAAWCMlSN2PGDJmmqcsvv9y/buvWrdq/f7+FqQAAAKwTkqUOAAAAJYXc3a+VYXt2mpxOp9UxAAAAKoyZOgAAABug1AEAANgAp18lJWetVJgRY3WMCiuY0t/qCAAAwGLM1AEAANgApQ4AAMAGKHUAAAA2QKkDAACwAUodAACADYRsqVu4cKHq1asnn89XYv3AgQM1bNgwi1IBAABYI2RL3c0336yioiItW7bMv+7QoUNavny5brvttlK38fl88nq9JRYAAAA7CNlSFx0drd/+9reaP3++f92iRYvUpEkTpaamlrqN2+2Wy+XyLwkJCUFKCwAAULlCttRJ0h133KE333xTX375pSQpJydHI0aMkMPhKHV8ZmamPB6PfyksLAxmXAAAgEoT0n9RokOHDkpJSdHChQvVp08f7dixQ8uXLy9zvGEYMgwjiAkBAACCI6RLnSTdfvvtmjFjhr788kv16tWLU6oAAKBaCunTr5L029/+Vl988YVmz55d5g0SAAAAdhfypc7lcummm25SrVq1NHDgQKvjAAAAWCLkS50kffnll/rd737H9XIAAKDaCulr6o4fP67c3Fzl5ubq6aefrvDrbM9Ok9PpDGAyAACA4ArpUtehQwcdP35cjzzyiJKSkqyOAwAAYJmQLnUFBQVWRwAAAKgSbHFNHQAAQHVHqQMAALABSh0AAIANUOoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2AClDgAAwAYodQAAADZAqQMAALABSh0AAIANUOoAAABsgFIHAABgA5Q6AAAAG6DUAQAA2EANqwNYyTRNSZLX67U4CQAAqO7O9ZFz/aS8qnWpO3r0qCQpISHB4iQAAADfO3HihFwuV7m3q9alLjY2VpK0b9++Cv3jIfi8Xq8SEhJUWFgop9NpdRxcAo5Z6OGYhR6OWegp7ZiZpqkTJ06ocePGFXrNal3qwsK+v6TQ5XLxQxBinE4nxyzEcMxCD8cs9HDMQs+Pj9lPmWTiRgkAAAAboNQBAADYQLUudYZhKCsrS4ZhWB0Fl4hjFno4ZqGHYxZ6OGahpzKOmcOs6H2zAAAAqDKq9UwdAACAXVDqAAAAbIBSBwAAYAOUOgAAABuwfambOXOmEhMTFRUVpS5dumjTpk0XHL948WJdfvnlioqKUtu2bfWf//wnSElxTnmOWU5OjhwOR4klKioqiGmrt3feeUcDBgxQ48aN5XA4tHTp0otuk5ubq44dO8owDLVo0UI5OTmVnhP/U95jlpube97PmMPh0IEDB4ITGHK73brqqqtUu3ZtNWzYUAMHDtSuXbsuuh2fZ9apyDELxOeZrUvdCy+8oHHjxikrK0tbtmxRSkqK0tLSdOjQoVLHv/feexoyZIh+//vfKy8vTwMHDtTAgQO1ffv2ICevvsp7zKTvv417//79/mXv3r1BTFy9nTp1SikpKZo5c+Yljd+zZ4/69++va6+9Vlu3blVGRoZuv/12rVy5spKT4pzyHrNzdu3aVeLnrGHDhpWUED+2du1a3Xnnnfrvf/+rVatW6ezZs+rTp49OnTpV5jZ8nlmrIsdMCsDnmWljnTt3Nu+8807/46KiIrNx48am2+0udfwtt9xi9u/fv8S6Ll26mH/4wx8qNSf+p7zHbP78+abL5QpSOlyIJPOVV1654JiJEyeaV1xxRYl1gwcPNtPS0ioxGcpyKcdszZo1piTz+PHjQcmEizt06JApyVy7dm2ZY/g8q1ou5ZgF4vPMtjN1Z86c0ebNm9WrVy//urCwMPXq1UsbNmwodZsNGzaUGC9JaWlpZY5HYFXkmEnSyZMn1bRpUyUkJOjXv/61duzYEYy4qAB+xkJX+/bt1ahRI/Xu3Vvr16+3Ok615vF4JEmxsbFljuFnrWq5lGMm/fTPM9uWuiNHjqioqEhxcXEl1sfFxZV5LciBAwfKNR6BVZFjlpSUpHnz5unVV1/VokWLVFxcrKuvvlpffPFFMCKjnMr6GfN6vfrmm28sSoULadSokWbNmqUlS5ZoyZIlSkhIUGpqqrZs2WJ1tGqpuLhYGRkZ6tatm5KTk8scx+dZ1XGpxywQn2c1AhEYsErXrl3VtWtX/+Orr75arVu31j//+U/99a9/tTAZYA9JSUlKSkryP7766quVn5+v6dOn69lnn7UwWfV05513avv27Xr33XetjoJLdKnHLBCfZ7adqatfv77Cw8N18ODBEusPHjyo+Pj4UreJj48v13gEVkWO2Y9FRESoQ4cO+uyzzyojIn6isn7GnE6noqOjLUqF8urcuTM/YxYYM2aMXn/9da1Zs0Y///nPLziWz7OqoTzH7Mcq8nlm21IXGRmpTp06afXq1f51xcXFWr16dYkm/ENdu3YtMV6SVq1aVeZ4BFZFjtmPFRUVadu2bWrUqFFlxcRPwM+YPWzdupWfsSAyTVNjxozRK6+8orfffluXXXbZRbfhZ81aFTlmP1ahz7OfdJtFFff888+bhmGYOTk55s6dO81Ro0aZderUMQ8cOGCapmkOGzbMvPfee/3j169fb9aoUcN87LHHzI8//tjMysoyIyIizG3btlm1C9VOeY9Zdna2uXLlSjM/P9/cvHmzeeutt5pRUVHmjh07rNqFauXEiRNmXl6emZeXZ0oyH3/8cTMvL8/cu3evaZqmee+995rDhg3zj//888/NmJgYc8KECebHH39szpw50wwPDzdXrFhh1S5UO+U9ZtOnTzeXLl1qfvrpp+a2bdvMsWPHmmFhYeZbb71l1S5UO6NHjzZdLpeZm5tr7t+/37+cPn3aP4bPs6qlIscsEJ9nti51pmmaf//7380mTZqYkZGRZufOnc3//ve//ud69uxppqenlxj/4osvmq1atTIjIyPNK664wly+fHmQE6M8xywjI8M/Ni4uzrzhhhvMLVu2WJC6ejr3dRc/Xs4do/T0dLNnz57nbdO+fXszMjLSbNasmTl//vyg567OynvMHnnkEbN58+ZmVFSUGRsba6ampppvv/22NeGrqdKOl6QSPzt8nlUtFTlmgfg8c/zfmwMAACCE2faaOgAAgOqEUgcAAGADlDoAAAAboNQBAADYAKUOAADABih1AAAANkCpAwAAsAFKHQAAwCV45513NGDAADVu3FgOh0NLly4t92uYpqnHHntMrVq1kmEY+tnPfqa//e1vAclHqQMACxUUFMjhcGjr1q1WRwFwEadOnVJKSopmzpxZ4dcYO3as5syZo8cee0yffPKJli1bps6dOwckX42AvAoAAIDN9evXT/369SvzeZ/Pp/vvv1///ve/9fXXXys5OVmPPPKIUlNTJUkff/yx/vGPf2j79u1KSkqSJF122WUBy8dMHYBqrbi4WFOnTlWLFi1kGIaaNGniPxWybds2XXfddYqOjla9evU0atQonTx50r9tamqqMjIySrzewIEDNWLECP/jxMREPfzww7rttttUu3ZtNWnSRM8884z/+XO/0Dt06CCHw+H/5Q8g9IwZM0YbNmzQ888/r48++kg333yz+vbtq08//VSS9Nprr6lZs2Z6/fXXddlllykxMVG33367jh07FpD3p9QBqNYyMzM1ZcoUPfDAA9q5c6f+9a9/KS4uTqdOnVJaWprq1q2r999/X4sXL9Zbb72lMWPGlPs9pk2bpiuvvFJ5eXn605/+pNGjR2vXrl2SpE2bNkmS3nrrLe3fv18vv/xyQPcPQHDs27dP8+fP1+LFi9WjRw81b95cf/nLX9S9e3fNnz9fkvT5559r7969Wrx4sRYuXKicnBxt3rxZgwYNCkgGTr8CqLZOnDihJ554Qk899ZTS09MlSc2bN1f37t01e/Zsffvtt1q4cKFq1qwpSXrqqac0YMAAPfLII4qLi7vk97nhhhv0pz/9SZJ0zz33aPr06VqzZo2SkpLUoEEDSVK9evUUHx8f4D0EECzbtm1TUVGRWrVqVWK9z+dTvXr1JH1/ZsDn82nhwoX+cXPnzlWnTp20a9cu/ynZiqLUAai2Pv74Y/l8Pl1//fWlPpeSkuIvdJLUrVs3FRcXa9euXeUqde3atfP/t8PhUHx8vA4dOvTTwgOoUk6ePKnw8HBt3rxZ4eHhJZ6rVauWJKlRo0aqUaNGieLXunVrSd/P9FHqAKCCoqOjf9L2YWFhMk2zxLqzZ8+eNy4iIqLEY4fDoeLi4p/03gCqlg4dOqioqEiHDh1Sjx49Sh3TrVs3fffdd8rPz1fz5s0lSbt375YkNW3a9Cdn4Jo6ANVWy5YtFR0drdWrV5/3XOvWrfXhhx/q1KlT/nXr169XWFiY//+mGzRooP379/ufLyoq0vbt28uVITIy0r8tgKrt5MmT2rp1q/8riPbs2aOtW7dq3759atWqlX73u99p+PDhevnll7Vnzx5t2rRJbrdby5cvlyT16tVLHTt21G233aa8vDxt3rxZf/jDH9S7d+/zTttWBKUOQLUVFRWle+65RxMnTtTChQuVn5+v//73v5o7d65+97vfKSoqSunp6dq+fbvWrFmjP//5zxo2bJj/1Ot1112n5cuXa/ny5frkk080evRoff311+XK0LBhQ0VHR2vFihU6ePCgPB5PJewpgED44IMP1KFDB3Xo0EGSNG7cOHXo0EGTJk2SJM2fP1/Dhw/X+PHjlZSUpIEDB+r9999XkyZNJH0/u//aa6+pfv36uuaaa9S/f3+1bt1azz//fEDycfoVQLX2wAMPqEaNGpo0aZK++uorNWrUSH/84x8VExOjlStXauzYsbrqqqsUExOjm266SY8//rh/29tuu00ffvihhg8frho1aujuu+/WtddeW673r1Gjhp588klNnjxZkyZNUo8ePZSbmxvgvQQQCKmpqeddcvFDERERys7OVnZ2dpljGjdurCVLllRGPDnMC6UDAABASOD0KwAAgA1Q6gAAAGyAUgcAAGADlDoAAAAboNQBAADYAKUOAADABih1AAAANkCpAwAAsAFKHQAAgA1Q6gAAAGyAUgcAAGADlDoAAAAb+P9EYpjpexBJsAAAAABJRU5ErkJggg==", - "text/plain": [ - "
" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "name": "stdout", - "output_type": "stream", - "text": [ - "Characters in dataset: 15.0 MB\n" - ] - } - ], - "source": [ + "assert cp.sum(histogram) == len(values)\n", + "\n", "histogram_host = cp.asnumpy(histogram)\n", "pairs = sorted(((i, c) for i, c in enumerate(histogram_host) if c), key=lambda x: x[1], reverse=True)[:20]\n", "labels = [('SPACE' if i == 32 else chr(i)) if 32 <= i <= 126 else f'0x{i:02X}' for i, _ in pairs]\n", @@ -648,7 +429,7 @@ "plt.xlabel('count')\n", "plt.tight_layout()\n", "plt.title(\"Top 20 Bins\")\n", - "plt.show()\n" + "plt.show()" ] }, { @@ -663,7 +444,7 @@ }, { "cell_type": "code", - "execution_count": 11, + "execution_count": null, "id": "d637b6b1-fb0b-4807-b70b-c80227c0fd6f", "metadata": { "colab": { @@ -679,28 +460,16 @@ "id": "d637b6b1-fb0b-4807-b70b-c80227c0fd6f", "outputId": "e3790713-5a2e-4b51-dd2c-182b8bc27a5b" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "==PROF== Connected to process 486 (/usr/bin/python3.12)\n", - "==PROF== Profiling \"histogram_localized[abi:v1,cw51cXTLSUwv1sDUaKthoaNgqamjgOR3W3Cw6igA9duC0hkwnNGiHEkYkrqQBFBTDEwCyQe21eqwcFW3UoDGjzpQEsgzrtUEAA_3d_3d]\": 0%....50%....100% - 44 passes\n", - "==PROF== Disconnected from process 486\n", - "==PROF== Report: /accelerated-computing-hub/tutorials/accelerated-python/notebooks/kernels/solutions/histogram_localized.ncu-rep\n" - ] - } - ], + "outputs": [], "source": [ - "%%ncu -o histogram_localized.ncu-rep\n", + "%%ncu -o histogram_localized.ncu-rep --kernel-name regex:histogram_localized\n", "histogram[:] = 0\n", - "histogram_localized[localized_blocks, threads_per_block](values, histogram)\n", - "cp.cuda.runtime.deviceSynchronize()" + "histogram_localized[blocks, threads_per_block](values, histogram)" ] }, { "cell_type": "code", - "execution_count": 12, + "execution_count": null, "id": "114e8ff7-b6fb-42ad-abda-f6d53479c052", "metadata": { "colab": { @@ -752,87 +521,7 @@ "id": "114e8ff7-b6fb-42ad-abda-f6d53479c052", "outputId": "1936cb5f-6708-45bc-ee39-e01ddc857b89" }, - "outputs": [ - { - "data": { - "text/html": [ - "\n", - " \n", - " " - ], - "text/plain": [ - "" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "5f0096f2398a4d9a8acaec8a5d434344", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Dropdown(description='Kernel:', layout=Layout(width='400px'), options=('histogram_localized',), style=Descript…" - ] - }, - "metadata": {}, - "output_type": "display_data" - }, - { - "data": { - "application/vnd.jupyter.widget-view+json": { - "model_id": "5d54028815a74c27b703e31417de3b05", - "version_major": 2, - "version_minor": 0 - }, - "text/plain": [ - "Output()" - ] - }, - "metadata": {}, - "output_type": "display_data" - } - ], + "outputs": [], "source": [ "assert cp.sum(histogram) == len(values)" ] @@ -851,7 +540,7 @@ }, { "cell_type": "code", - "execution_count": 13, + "execution_count": null, "id": "2a3f9ca4-b61b-4536-9896-7a41498cc986", "metadata": { "colab": { @@ -867,19 +556,9 @@ "id": "2a3f9ca4-b61b-4536-9896-7a41498cc986", "outputId": "08bb2e07-1dda-4cf1-c2d3-7125d575cd8f" }, - "outputs": [ - { - "name": "stdout", - "output_type": "stream", - "text": [ - "histogram_global: 4.15 ms ± 0.60% (mean ± relative stdev of 15 runs)\n", - "histogram_localized: 0.163 ms ± 7.60% (mean ± relative stdev of 15 runs)\n", - "histogram_localized speedup over histogram_global: 25.46\n" - ] - } - ], + "outputs": [], "source": [ - "global_times = cpx.profiler.benchmark(launch_global, n_repeat=15, n_warmup=4).gpu_times[0]\n", + "global_times = cpx.profiler.benchmark(lambda: histogram_global[blocks, threads_per_block](values, histogram), n_repeat=15, n_warmup=4).gpu_times[0]\n", "localized_times = cpx.profiler.benchmark(launch_localized, n_repeat=15, n_warmup=4).gpu_times[0]\n", "histogram_global_duration = global_times.mean() * 1000\n", "histogram_localized_duration = localized_times.mean() * 1000\n", @@ -887,7 +566,7 @@ "\n", "print(f\"histogram_global: {histogram_global_duration:.3g} ms\")\n", "print(f\"histogram_localized: {histogram_localized_duration:.3g} ms\")\n", - "print(f\"histogram_localized speedup over histogram_global: {speedup:.2f}\")\n" + "print(f\"histogram_localized speedup over histogram_global: {speedup:.2f}\")" ] } ],