{ "cells": [ { "cell_type": "markdown", "id": "0", "metadata": {}, "source": [ "Getting started with API tracing\n", "================================\n", "\n", "This notebook shows how to get started with API tracing.\n", "\n", "Tracing is the process of collecting and examining information about activities that happen during program execution.\n", "In particular, CUDA API tracing records the CUDA API calls issued by a program: their timing and their correlation with the GPU activities they launch, such as kernel executions.\n", "NVIDIA provides [Nsight Systems](https://developer.nvidia.com/nsight-systems) for this purpose.\n", "`ReProspect` enables a fully programmatic use of this tool: it launches it, collects its output into Python data structures, and supports analyses of these data that can go all the way to test assertions.\n", "\n", "As an example, we consider a program that launches four kernels with a diamond dependency structure.\n", "We programmatically trace its execution, examine the recorded CUDA API calls, and verify the concurrency structure.\n" ] }, { "cell_type": "markdown", "id": "1", "metadata": {}, "source": [ "About this page\n", "---------------\n", "\n", "This page is rendered from a Jupyter notebook, executed at documentation build time, located at `docs/source/getting-started/example_api_tracing.ipynb` in the repository.\n", "\n", "To run the example yourself, {download}`download the notebook ` and open it in JupyterLab.\n", "Alternatively, copy-paste the code snippets successively into an interactive Python session.\n", "\n", "The requirements are Python 3.10 or newer and `ReProspect`:\n", "\n", "```bash\n", "python -m pip install reprospect\n", "```\n", "\n", "JupyterLab may be installed as:\n", "\n", "```bash\n", "python -m pip install jupyterlab\n", "```\n", "\n", "The example also requires:\n", "- a CUDA Toolkit installation providing `nvcc`;\n", "- an [Nsight Systems CLI installation](https://docs.nvidia.com/nsight-systems/InstallationGuide/index.html#installation-guide) providing the command-line tool `nsys`;\n", "- the NVTX header `nvtx3/nvtx3.hpp` (included in recent CUDA Toolkits; otherwise available from [NVTX installation](https://github.com/NVIDIA/NVTX#how-do-i-get-nvtx));\n", "- a C++20-capable toolchain.\n", "\n", "It should be noted that `ReProspect` provides the script {py:mod}`reprospect.utils.installers.nsight_systems` for installing Nsight Systems through `apt`.\n", "\n", "Because tracing executes the program, a GPU is required." ] }, { "cell_type": "markdown", "id": "2", "metadata": {}, "source": [ "Source code\n", "-----------\n", "\n", "Let us consider a program that launches four kernels with a diamond dependency structure:\n", "\n", "```{tikz} Four kernels with a diamond dependency structure\n", ":align: center\n", "\n", "\\begin{tikzpicture}[\n", " kernel/.style={\n", " draw,\n", " rounded corners,\n", " minimum width=1.2cm,\n", " minimum height=0.7cm\n", " },\n", " >=stealth\n", "]\n", "\\node[kernel] (A) at (0, 2) {kernel A};\n", "\\node[kernel] (B) at (-2, 0) {kernel B};\n", "\\node[kernel] (C) at (2, 0) {kernel C};\n", "\\node[kernel] (D) at (0, -2) {kernel D};\n", "\n", "\\draw[->] (A) -- (B);\n", "\\draw[->, dashed] (A) -- node[right=2pt] {\\scriptsize event\\_fork} (C);\n", "\\draw[->] (B) -- (D);\n", "\\draw[->, dashed] (C) -- node[right=2pt] {\\scriptsize event\\_join} (D);\n", "\\end{tikzpicture}\n", "```\n", "\n", "The four kernels are launched across two CUDA streams: kernels A, B, and D on one stream, and kernel C on the other, so that kernels B and C may potentially execute concurrently.\n", "Two of the dependencies are thus implied by stream order: A -> B and B -> D.\n", "The two cross-stream dependencies, A -> C and C -> D, are expressed with CUDA events, recorded on one stream and waited for on the other.\n", "\n", "The source code features NVTX annotations, for two purposes.\n", "On the one hand, the outer NVTX range `start_end_range` controls when Nsight Systems starts and stops collecting data.\n", "On the other hand, the inner ranges `diamond_range` and `check_range` allow the analyses of the recorded activity to be scoped to these specific regions, as we will see below.\n", "The ranges belong to a dedicated NVTX domain, which keeps them separate from other NVTX annotations, such as annotations that libraries used by an application may emit themselves." ] }, { "cell_type": "code", "execution_count": null, "id": "3", "metadata": {}, "outputs": [], "source": [ "CODE = \"\"\"\\\n", "#include \n", "#include \n", "#include \n", "#include \n", "\n", "#include \n", "#include \n", "\n", "inline void check_cudart_call(\n", " const cudaError_t status,\n", " const char* const statement,\n", " const std::source_location& loc = std::source_location::current()) {\n", " if (status != cudaSuccess) {\n", " std::ostringstream oss;\n", " oss << statement << \" failed: \" << status << \" (\" << cudaGetErrorName(status)\n", " << \"): \" << cudaGetErrorString(status) << \" (\" << loc.file_name() << \":\" << loc.line() << \")\";\n", "\n", " throw std::runtime_error(oss.str());\n", " }\n", "}\n", "\n", "#define CHECK_CUDART_CALL(statement) check_cudart_call((statement), #statement)\n", "\n", "__global__ void check_add_kernel(int* data, int prev, int val) {\n", " assert(*data == prev);\n", " *data += val;\n", "}\n", "\n", "__global__ void increment_kernel(int* data) {\n", " atomicAdd(data, 1);\n", "}\n", "\n", "struct ExampleApiTracingDomain {\n", " static constexpr char const * name{\"example_api_tracing_domain\"};\n", "};\n", "\n", "int main() {\n", " cudaStream_t stream_0, stream_1;\n", " CHECK_CUDART_CALL(cudaStreamCreate(&stream_0));\n", " CHECK_CUDART_CALL(cudaStreamCreate(&stream_1));\n", "\n", " cudaEvent_t event_fork, event_join;\n", " CHECK_CUDART_CALL(cudaEventCreateWithFlags(&event_fork, cudaEventDisableTiming));\n", " CHECK_CUDART_CALL(cudaEventCreateWithFlags(&event_join, cudaEventDisableTiming));\n", "\n", " int* data;\n", " CHECK_CUDART_CALL(cudaMallocAsync(&data, sizeof(int), stream_0));\n", " CHECK_CUDART_CALL(cudaMemsetAsync(data, 0, sizeof(int), stream_0));\n", "\n", " const auto start_end_range = nvtx3::start_range_in(\"start_end_range\");\n", "\n", " {\n", " nvtx3::scoped_range_in diamond_range(\"diamond_range\");\n", " check_add_kernel<<<1, 1, 0, stream_0>>>(data, 0, 4); // kernel A\n", "\n", " CHECK_CUDART_CALL(cudaEventRecord(event_fork, stream_0));\n", "\n", " increment_kernel<<<1, 1, 0, stream_0>>>(data); // kernel B\n", "\n", " CHECK_CUDART_CALL(cudaStreamWaitEvent(stream_1, event_fork));\n", " increment_kernel<<<1, 1, 0, stream_1>>>(data); // kernel C\n", "\n", " CHECK_CUDART_CALL(cudaEventRecord(event_join, stream_1));\n", " CHECK_CUDART_CALL(cudaStreamWaitEvent(stream_0, event_join));\n", "\n", " check_add_kernel<<<1, 1, 0, stream_0>>>(data, 6, 3); // kernel D\n", " }\n", "\n", " {\n", " nvtx3::scoped_range_in check_range(\"check_range\");\n", " int data_h;\n", "\n", " CHECK_CUDART_CALL(cudaMemcpyAsync(&data_h, data, sizeof(int), cudaMemcpyDeviceToHost, stream_0));\n", " CHECK_CUDART_CALL(cudaStreamSynchronize(stream_0));\n", "\n", " if (data_h != 9) {\n", " throw std::runtime_error(\"wrong value\");\n", " }\n", " }\n", "\n", " nvtx3::end_range_in(start_end_range);\n", "\n", " CHECK_CUDART_CALL(cudaFreeAsync(data, stream_0));\n", "\n", " CHECK_CUDART_CALL(cudaEventDestroy(event_fork));\n", " CHECK_CUDART_CALL(cudaEventDestroy(event_join));\n", "\n", " CHECK_CUDART_CALL(cudaStreamDestroy(stream_0));\n", " CHECK_CUDART_CALL(cudaStreamDestroy(stream_1));\n", "}\n", "\"\"\"" ] }, { "cell_type": "markdown", "id": "4", "metadata": {}, "source": [ "Compilation\n", "-----------\n", "\n", "Because this example will execute the program, we compile the source code for the native architecture, i.e., the architecture of the GPU present on the machine." ] }, { "cell_type": "code", "execution_count": null, "id": "5", "metadata": {}, "outputs": [], "source": [ "import pathlib\n", "import subprocess\n", "\n", "from reprospect.utils import rich_helpers\n", "from reprospect.utils.detect import GPUDetector\n", "\n", "print(subprocess.check_output(('nvcc', '--version')).decode().strip())\n", "\n", "print(subprocess.check_output(('nsys', '--version')).decode().strip())\n", "\n", "visible_gpus = GPUDetector().detect()\n", "print(f'Visible GPUs:\\n{rich_helpers.to_string(rich_helpers.df_to_table(visible_gpus))}')\n", "\n", "workdir = pathlib.Path.cwd() / 'example_api_tracing'\n", "workdir.mkdir(exist_ok=True)\n", "print(f'Working directory: {workdir}')\n", "\n", "source = workdir / 'concurrency.cu'\n", "executable = workdir / 'concurrency'\n", "\n", "source.write_text(CODE)\n", "_ = subprocess.check_call(('nvcc', '-arch=native', '-std=c++20', '-O3', '-o', executable, source))" ] }, { "cell_type": "markdown", "id": "6", "metadata": {}, "source": [ "Here, {py:class}`reprospect.utils.detect.GPUDetector` is a utility class to detect the visible GPUs through an `nvidia-smi` query.\n", "It returns the data collected about the visible GPUs as a {py:class}`pandas.DataFrame`.\n", "The module {py:mod}`reprospect.utils.rich_helpers` provides functionality for rich rendering of such data frames." ] }, { "cell_type": "markdown", "id": "7", "metadata": {}, "source": [ "Running API tracing\n", "-------------------\n", "\n", "Nsight Systems provides the command-line tool `nsys` for collecting data from the execution of a program. \n", "Here, we invoke it through the `ReProspect` class {py:class}`reprospect.tools.nsys.session.Session`.\n", "The argument `executable` designates the executable on which to collect data.\n", "The argument `nvtx_capture` is the NVTX range that controls when data collection starts and stops.\n", "The value `'start_end_range@example_api_tracing_domain'` follows the `nsys` [convention](https://docs.nvidia.com/nsight-systems/UserGuide/index.html#cli-profile-command-switch-options) `@`.\n", "The argument `output` determines the output file; for the passed value `workdir / executable.name`, the report is written to `workdir / f'{executable.name}.nsys-rep'`." ] }, { "cell_type": "code", "execution_count": null, "id": "8", "metadata": {}, "outputs": [], "source": [ "from reprospect.tools.nsys import Command, Session\n", "\n", "ns = Session(\n", " command=Command(\n", " executable=executable,\n", " output=workdir / executable.name,\n", " nvtx_capture='start_end_range@example_api_tracing_domain',\n", " ),\n", ")\n", "\n", "ns.run(cwd=workdir)" ] }, { "cell_type": "markdown", "id": "9", "metadata": {}, "source": [ "CUDA API Trace default report\n", "-----------------------------\n", "\n", "Nsight Systems provides its own functionality for post-processing the collected data into summary and trace reports, as described [here](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html#statistical-reports-shipped-with-product-name).\n", "Nsight Systems generates these default reports through SQL queries on an SQLite export of the output file.\n", "\n", "The `ReProspect` method {py:meth}`reprospect.tools.nsys.session.Session.extract_statistical_report` allows such reports to be retrieved as a {py:class}`pandas.DataFrame`.\n", "Here, we retrieve the report named `cuda_api_trace`, which contains a trace of the CUDA API calls, with their start times and durations.\n", "\n", "We first export the `.nsys-rep` report to an SQLite database.\n", "The method {py:meth}`~reprospect.tools.nsys.session.Session.extract_statistical_report` then invokes `nsys stats` on this database.\n", "The same database will be queried directly in the sections below." ] }, { "cell_type": "code", "execution_count": null, "id": "10", "metadata": {}, "outputs": [], "source": [ "sqlite_database = ns.export_to_sqlite(cwd=workdir)\n", "print(f'SQLite database: {sqlite_database}')" ] }, { "cell_type": "code", "execution_count": null, "id": "11", "metadata": {}, "outputs": [], "source": [ "cuda_api_trace_report = ns.extract_statistical_report(report='cuda_api_trace')\n", "print(f'Nsight Systems cuda_api_trace report:\\n{rich_helpers.to_string(rich_helpers.df_to_table(cuda_api_trace_report))}')" ] }, { "cell_type": "markdown", "id": "12", "metadata": {}, "source": [ "Querying CUDA API tracing data\n", "------------------------------\n", "\n", "Beyond the default reports of the previous section, the Nsight Systems documentation recommends the SQLite export of the `.nsys-rep` file as the primary means for custom analyses of the collected data, with full access to all recorded details and their correlations.\n", "The Nsight Systems documentation provides the SQLite schema reference [here](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html#sqlite-schema-reference).\n", "\n", "The `ReProspect` class {py:class}`reprospect.tools.nsys.report.Report` provides access to the database.\n", "The class is used as a context manager, which opens and closes the connection to the database.\n", "\n", "The property {py:attr}`reprospect.tools.nsys.report.Report.tables` provides a list of the names of the tables available in the database." ] }, { "cell_type": "code", "execution_count": null, "id": "13", "metadata": {}, "outputs": [], "source": [ "from reprospect.tools.nsys import Report\n", "\n", "report = Report(db=ns.command.output.with_suffix('.sqlite'))\n", "\n", "with report:\n", " tables = report.tables\n", " print(f'Tables:\\n{tables}')" ] }, { "cell_type": "markdown", "id": "14", "metadata": {}, "source": [ "The method {py:meth}`reprospect.tools.nsys.report.Report.table` retrieves a table as a {py:class}`pandas.DataFrame`.\n", "Let us retrieve the tables:\n", "- `CUPTI_ACTIVITY_KIND_RUNTIME` with the trace of the CUDA runtime API calls;\n", "- `CUPTI_ACTIVITY_KIND_CUDA_EVENT` with details relevant to the CUDA events;\n", "- `CUPTI_ACTIVITY_KIND_SYNCHRONIZATION` with details relevant to the CUDA synchronization operations;\n", "- `CUPTI_ACTIVITY_KIND_KERNEL` with details relevant to the kernel executions;\n", "- `ENUM_NSYS_EVENT_CLASS`, which we will use below to interpret the `eventClass` column of the runtime table;\n", "- `StringIds`, which relates string identifiers to the strings themselves.\n", "\n", "Here, `CUPTI` stands for the [CUDA Profiling Tools Interface (CUPTI)](https://developer.nvidia.com/cupti), which Nsight Systems uses to collect the associated data." ] }, { "cell_type": "code", "execution_count": null, "id": "15", "metadata": {}, "outputs": [], "source": [ "with report:\n", " runtime = report.table(name='CUPTI_ACTIVITY_KIND_RUNTIME')\n", " print(f'Table CUPTI_ACTIVITY_KIND_RUNTIME:\\n{rich_helpers.to_string(rich_helpers.df_to_table(runtime))}')\n", "\n", " cuda_event = report.table(name='CUPTI_ACTIVITY_KIND_CUDA_EVENT')\n", " print(f'Table CUPTI_ACTIVITY_KIND_CUDA_EVENT:\\n{rich_helpers.to_string(rich_helpers.df_to_table(cuda_event))}')\n", "\n", " synchronization = report.table(name='CUPTI_ACTIVITY_KIND_SYNCHRONIZATION')\n", " print(f'Table CUPTI_ACTIVITY_KIND_SYNCHRONIZATION:\\n{rich_helpers.to_string(rich_helpers.df_to_table(synchronization))}')\n", "\n", " kernel = report.table(name='CUPTI_ACTIVITY_KIND_KERNEL')\n", " print(f'Table CUPTI_ACTIVITY_KIND_KERNEL, for brevity only its column names:\\n{kernel.columns}')\n", "\n", " enum_event_class = report.table(name='ENUM_NSYS_EVENT_CLASS')\n", " print(f'Table ENUM_NSYS_EVENT_CLASS first two rows:\\n{rich_helpers.to_string(rich_helpers.df_to_table(enum_event_class[:2]))}')\n", "\n", " stringids = report.table(name='StringIds')\n", " first_runtime_row = runtime.iloc[0]\n", " print(f'Table StringIds row related to first runtime table row:\\n{rich_helpers.to_string(rich_helpers.df_to_table(stringids[stringids[\"id\"] == first_runtime_row[\"nameId\"]]))}')" ] }, { "cell_type": "markdown", "id": "16", "metadata": {}, "source": [ "As we can observe, analysing the data requires exploiting the relationships between the tables.\n", "Indeed, to retrieve the name of the activity described in the first row of the runtime table, we must read the `nameId` from the runtime table and correlate it via the `id` in the `StringIds` table with its corresponding `value`." ] }, { "cell_type": "markdown", "id": "17", "metadata": {}, "source": [ "Querying CUDA API tracing data by nested NVTX range\n", "---------------------------------------------------\n", "\n", "`ReProspect` provides functionalities to facilitate lookups in the SQLite database.\n", "The method {py:meth}`reprospect.tools.nsys.report.Report.get_events` provides focused lookups by nested NVTX range, with automatic correlation of string identifiers.\n", "For tables featuring `start` and `end` columns, it selects the rows whose time span is within a given nested NVTX range.\n", "The argument `accessors` designates the path of nested NVTX ranges, from the outermost inward." ] }, { "cell_type": "code", "execution_count": null, "id": "18", "metadata": {}, "outputs": [], "source": [ "with report:\n", " print(report.nvtx_events)" ] }, { "cell_type": "code", "execution_count": null, "id": "19", "metadata": {}, "outputs": [], "source": [ "with report:\n", " diamond_runtime = report.get_events(table='CUPTI_ACTIVITY_KIND_RUNTIME', accessors=['start_end_range', 'diamond_range'])\n", " print(f'Runtime events in diamond range:\\n{diamond_runtime}')" ] }, { "cell_type": "markdown", "id": "20", "metadata": {}, "source": [ "Some runtime activity names carry a suffix such as `_v7000` and `_v3020`.\n", "The function {py:func}`reprospect.tools.nsys.report.strip_cuda_api_suffix` allows such suffixes to be stripped, so that the analysis does not depend on them.\n", "The runtime table may also contain CUDA driver activity, such as `cuLibraryLoadData`, `cuLibraryGetKernel`, `cuKernelGetName` related to module loading and name resolution.\n", "Such entries may depend on the execution environment rather than on the program; we exclude them by selecting only the rows whose `eventClass` corresponds to the label `'CUDA runtime'` in the `ENUM_NSYS_EVENT_CLASS` table." ] }, { "cell_type": "code", "execution_count": null, "id": "21", "metadata": {}, "outputs": [], "source": [ "from reprospect.tools.nsys import strip_cuda_api_suffix\n", "\n", "RUNTIME_EVENT_CLASS_ID = Report.single_row(data=enum_event_class[enum_event_class['label'] == 'CUDA runtime'])['id']\n", "\n", "diamond_cuda_runtime = (\n", " diamond_runtime.loc[diamond_runtime['eventClass'] == RUNTIME_EVENT_CLASS_ID]\n", " .assign(api=lambda df: df['name'].map(strip_cuda_api_suffix))\n", ")\n", "print(f'CUDA runtime api calls in diamond range:\\n{diamond_cuda_runtime[\"api\"].to_list()}')" ] }, { "cell_type": "markdown", "id": "22", "metadata": {}, "source": [ "The method {py:meth}`reprospect.tools.nsys.report.Report.get_correlated_row` facilitates correlated lookups across tables in the database.\n", "By default, it uses the `correlationId` column to join both tables.\n", "The columns to use for the correlation can also be designated explicitly, as needed for instance when retrieving the name of a kernel by correlating its `demangledName` identifier in the `CUPTI_ACTIVITY_KIND_KERNEL` table with the `id` in the `StringIds` table:" ] }, { "cell_type": "code", "execution_count": null, "id": "23", "metadata": {}, "outputs": [], "source": [ "kernel_a = Report.get_correlated_row(src=diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaLaunchKernel'].iloc[0], dst=kernel)\n", "print('StreamId for kernel A:', kernel_a['streamId'])\n", "\n", "kernel_a_demangled_name = Report.get_correlated_row(src=kernel_a, dst=stringids, correlation_src='demangledName', correlation_dst='id')\n", "print(f'Demangled name for kernel A: {kernel_a_demangled_name[\"value\"]}')" ] }, { "cell_type": "markdown", "id": "24", "metadata": {}, "source": [ "Assertions on CUDA API tracing data\n", "-----------------------------------\n", "\n", "With the collected data gathered in Python data structures, the analysis can now go all the way to test assertions.\n", "Here, we programmatically verify the concurrency structure." ] }, { "cell_type": "markdown", "id": "25", "metadata": {}, "source": [ "First, we verify the sequence of the CUDA runtime API calls.\n", "The sequence assertion checks for exact equality: after stripping the name suffixes and excluding the driver activity, the remaining sequence is fully determined by the order of the host-side calls in the source code." ] }, { "cell_type": "code", "execution_count": null, "id": "26", "metadata": {}, "outputs": [], "source": [ "assert diamond_cuda_runtime['api'].to_list() == [\n", " 'cudaLaunchKernel',\n", " 'cudaEventRecord',\n", " 'cudaLaunchKernel',\n", " 'cudaStreamWaitEvent',\n", " 'cudaLaunchKernel',\n", " 'cudaEventRecord',\n", " 'cudaStreamWaitEvent',\n", " 'cudaLaunchKernel',\n", "]" ] }, { "cell_type": "markdown", "id": "27", "metadata": {}, "source": [ "Next, we verify the streams the kernels are launched on.\n", "In particular, we check that the potentially concurrent kernels B and C are launched on different streams and that kernels A, B, and D are launched on the same stream." ] }, { "cell_type": "code", "execution_count": null, "id": "28", "metadata": {}, "outputs": [], "source": [ "launches = diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaLaunchKernel']\n", "kernel_a, kernel_b, kernel_c, kernel_d = (Report.get_correlated_row(src=launches.iloc[idx], dst=kernel) for idx in range(4))\n", "assert kernel_b['streamId'] != kernel_c['streamId']\n", "\n", "assert kernel_a['streamId'] == kernel_b['streamId'] == kernel_d['streamId']" ] }, { "cell_type": "markdown", "id": "29", "metadata": {}, "source": [ "Finally, we verify the dependencies that form the diamond.\n", "In particular, for each dependency, we check that the event is recorded on the predecessor's stream, that the wait executes on the successor's stream, and that record and wait reference the same event." ] }, { "cell_type": "code", "execution_count": null, "id": "30", "metadata": {}, "outputs": [], "source": [ "event_records = diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaEventRecord']\n", "stream_wait_events = diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaStreamWaitEvent']\n", "\n", "event_record_fork = Report.get_correlated_row(src=event_records.iloc[0], dst=cuda_event)\n", "stream_wait_event_fork = Report.get_correlated_row(src=stream_wait_events.iloc[0], dst=synchronization)\n", "assert event_record_fork['streamId'] == kernel_a['streamId']\n", "assert stream_wait_event_fork['streamId'] == kernel_c['streamId']\n", "assert stream_wait_event_fork['eventId'] == event_record_fork['eventId']\n", "\n", "event_record_join = Report.get_correlated_row(src=event_records.iloc[1], dst=cuda_event)\n", "stream_wait_event_join = Report.get_correlated_row(src=stream_wait_events.iloc[1], dst=synchronization)\n", "assert event_record_join['streamId'] == kernel_c['streamId']\n", "assert stream_wait_event_join['streamId'] == kernel_d['streamId']\n", "assert stream_wait_event_join['eventId'] == event_record_join['eventId']" ] }, { "cell_type": "markdown", "id": "31", "metadata": {}, "source": [ "Outlook\n", "-------\n", "\n", "In the example of this notebook, the traced sequence can be read off the source code directly.\n", "One context in which the proposed approach becomes most valuable is that of abstraction layers.\n", "Portability libraries such as `Kokkos` let applications express computations at a higher level of abstraction and map them to CUDA API calls internally.\n", "Likewise, the C++26 `std::execution` model, whose customization for CUDA is under development in [NVIDIA's `CCCL` library](https://github.com/NVIDIA/cccl), lowers declarative descriptions of asynchronous work, including fork–join structures like the diamond of this notebook, to streams and events.\n", "API tracing then allows verifying that these mappings produce the intended calls.\n", "With `ReProspect`'s fully programmatic approach, such verifications can run as tests in CI/CD pipelines.\n", "API tracing is used in our {ref}`Kokkos View allocation case study `, where it elucidates benchmarking results by identifying the CUDA API calls that `Kokkos` issues when allocating a `Kokkos::View` under different scenarios." ] } ], "metadata": { "language_info": { "name": "python" } }, "nbformat": 4, "nbformat_minor": 5 }