Getting started with API tracing
This notebook shows how to get started with API tracing.
Tracing is the process of collecting and examining information about activities that happen during program execution.
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.
NVIDIA provides Nsight Systems for this purpose.
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.
As an example, we consider a program that launches four kernels with a diamond dependency structure. We programmatically trace its execution, examine the recorded CUDA API calls, and verify the concurrency structure.
About this page
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.
To run the example yourself, download the notebook and open it in JupyterLab.
Alternatively, copy-paste the code snippets successively into an interactive Python session.
The requirements are Python 3.10 or newer and ReProspect:
python -m pip install reprospect
JupyterLab may be installed as:
python -m pip install jupyterlab
The example also requires:
a CUDA Toolkit installation providing
nvcc;an Nsight Systems CLI installation providing the command-line tool
nsys;the NVTX header
nvtx3/nvtx3.hpp(included in recent CUDA Toolkits; otherwise available from NVTX installation);a C++20-capable toolchain.
It should be noted that ReProspect provides the script reprospect.utils.installers.nsight_systems for installing Nsight Systems through apt.
Because tracing executes the program, a GPU is required.
Source code
Let us consider a program that launches four kernels with a diamond dependency structure:
Four kernels with a diamond dependency structure
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. Two of the dependencies are thus implied by stream order: A -> B and B -> D. 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.
The source code features NVTX annotations, for two purposes.
On the one hand, the outer NVTX range start_end_range controls when Nsight Systems starts and stops collecting data.
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.
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.
CODE = """\
#include <cassert>
#include <source_location>
#include <sstream>
#include <stdexcept>
#include <cuda_runtime.h>
#include <nvtx3/nvtx3.hpp>
inline void check_cudart_call(
const cudaError_t status,
const char* const statement,
const std::source_location& loc = std::source_location::current()) {
if (status != cudaSuccess) {
std::ostringstream oss;
oss << statement << " failed: " << status << " (" << cudaGetErrorName(status)
<< "): " << cudaGetErrorString(status) << " (" << loc.file_name() << ":" << loc.line() << ")";
throw std::runtime_error(oss.str());
}
}
#define CHECK_CUDART_CALL(statement) check_cudart_call((statement), #statement)
__global__ void check_add_kernel(int* data, int prev, int val) {
assert(*data == prev);
*data += val;
}
__global__ void increment_kernel(int* data) {
atomicAdd(data, 1);
}
struct ExampleApiTracingDomain {
static constexpr char const * name{"example_api_tracing_domain"};
};
int main() {
cudaStream_t stream_0, stream_1;
CHECK_CUDART_CALL(cudaStreamCreate(&stream_0));
CHECK_CUDART_CALL(cudaStreamCreate(&stream_1));
cudaEvent_t event_fork, event_join;
CHECK_CUDART_CALL(cudaEventCreateWithFlags(&event_fork, cudaEventDisableTiming));
CHECK_CUDART_CALL(cudaEventCreateWithFlags(&event_join, cudaEventDisableTiming));
int* data;
CHECK_CUDART_CALL(cudaMallocAsync(&data, sizeof(int), stream_0));
CHECK_CUDART_CALL(cudaMemsetAsync(data, 0, sizeof(int), stream_0));
const auto start_end_range = nvtx3::start_range_in<ExampleApiTracingDomain>("start_end_range");
{
nvtx3::scoped_range_in<ExampleApiTracingDomain> diamond_range("diamond_range");
check_add_kernel<<<1, 1, 0, stream_0>>>(data, 0, 4); // kernel A
CHECK_CUDART_CALL(cudaEventRecord(event_fork, stream_0));
increment_kernel<<<1, 1, 0, stream_0>>>(data); // kernel B
CHECK_CUDART_CALL(cudaStreamWaitEvent(stream_1, event_fork));
increment_kernel<<<1, 1, 0, stream_1>>>(data); // kernel C
CHECK_CUDART_CALL(cudaEventRecord(event_join, stream_1));
CHECK_CUDART_CALL(cudaStreamWaitEvent(stream_0, event_join));
check_add_kernel<<<1, 1, 0, stream_0>>>(data, 6, 3); // kernel D
}
{
nvtx3::scoped_range_in<ExampleApiTracingDomain> check_range("check_range");
int data_h;
CHECK_CUDART_CALL(cudaMemcpyAsync(&data_h, data, sizeof(int), cudaMemcpyDeviceToHost, stream_0));
CHECK_CUDART_CALL(cudaStreamSynchronize(stream_0));
if (data_h != 9) {
throw std::runtime_error("wrong value");
}
}
nvtx3::end_range_in<ExampleApiTracingDomain>(start_end_range);
CHECK_CUDART_CALL(cudaFreeAsync(data, stream_0));
CHECK_CUDART_CALL(cudaEventDestroy(event_fork));
CHECK_CUDART_CALL(cudaEventDestroy(event_join));
CHECK_CUDART_CALL(cudaStreamDestroy(stream_0));
CHECK_CUDART_CALL(cudaStreamDestroy(stream_1));
}
"""
Compilation
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.
import pathlib
import subprocess
from reprospect.utils import rich_helpers
from reprospect.utils.detect import GPUDetector
print(subprocess.check_output(('nvcc', '--version')).decode().strip())
print(subprocess.check_output(('nsys', '--version')).decode().strip())
visible_gpus = GPUDetector().detect()
print(f'Visible GPUs:\n{rich_helpers.to_string(rich_helpers.df_to_table(visible_gpus))}')
workdir = pathlib.Path.cwd() / 'example_api_tracing'
workdir.mkdir(exist_ok=True)
print(f'Working directory: {workdir}')
source = workdir / 'concurrency.cu'
executable = workdir / 'concurrency'
source.write_text(CODE)
_ = subprocess.check_call(('nvcc', '-arch=native', '-std=c++20', '-O3', '-o', executable, source))
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Fri_Nov__7_07:23:37_PM_PST_2025
Cuda compilation tools, release 13.1, V13.1.80
Build cuda_13.1.r13.1/compiler.36836380_0
NVIDIA Nsight Systems version 2025.5.2.266-255236693005v0
Visible GPUs:
┏━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━┳━━━━━━━━━━━━━━┓
┃ uuid ┃ index ┃ name ┃ compute_cap ┃ architecture ┃
┡━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━╇━━━━━━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━╇━━━━━━━━━━━━━━┩
│ GPU-d27df728-7612-809e-697e-38da9785efa3 │ 0 │ NVIDIA GeForce RTX 5070 Ti │ 12.0 │ BLACKWELL120 │
│ GPU-f0b17181-5e70-f633-ed96-14cf7e77e8e9 │ 1 │ NVIDIA GeForce RTX 5070 Ti │ 12.0 │ BLACKWELL120 │
└──────────────────────────────────────────┴───────┴────────────────────────────┴─────────────┴──────────────┘
Working directory: /tmp/tmprqugcno6/example_api_tracing
Here, reprospect.utils.detect.GPUDetector is a utility class to detect the visible GPUs through an nvidia-smi query.
It returns the data collected about the visible GPUs as a pandas.DataFrame.
The module reprospect.utils.rich_helpers provides functionality for rich rendering of such data frames.
Running API tracing
Nsight Systems provides the command-line tool nsys for collecting data from the execution of a program.
Here, we invoke it through the ReProspect class reprospect.tools.nsys.session.Session.
The argument executable designates the executable on which to collect data.
The argument nvtx_capture is the NVTX range that controls when data collection starts and stops.
The value 'start_end_range@example_api_tracing_domain' follows the nsys convention <range>@<domain>.
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'.
from reprospect.tools.nsys import Command, Session
ns = Session(
command=Command(
executable=executable,
output=workdir / executable.name,
nvtx_capture='start_end_range@example_api_tracing_domain',
),
)
ns.run(cwd=workdir)
Capture range started in the application.
Capture range ended in the application.
Generating '/tmp/nsys-report-5c4c.qdstrm'
[1/1] [0% ] concurrency.nsys-repProcessing events...
[1/1] [========================100%] concurrency.nsys-rep
Generated:
/tmp/tmprqugcno6/example_api_tracing/concurrency.nsys-rep
CUDA API Trace default report
Nsight Systems provides its own functionality for post-processing the collected data into summary and trace reports, as described here. Nsight Systems generates these default reports through SQL queries on an SQLite export of the output file.
The ReProspect method reprospect.tools.nsys.session.Session.extract_statistical_report() allows such reports to be retrieved as a pandas.DataFrame.
Here, we retrieve the report named cuda_api_trace, which contains a trace of the CUDA API calls, with their start times and durations.
We first export the .nsys-rep report to an SQLite database.
The method extract_statistical_report() then invokes nsys stats on this database.
The same database will be queried directly in the sections below.
sqlite_database = ns.export_to_sqlite(cwd=workdir)
print(f'SQLite database: {sqlite_database}')
Processing 63 events:
SQLite database: /tmp/tmprqugcno6/example_api_tracing/concurrency.sqlite
[==========================================================================100%]
cuda_api_trace_report = ns.extract_statistical_report(report='cuda_api_trace')
print(f'Nsight Systems cuda_api_trace report:\n{rich_helpers.to_string(rich_helpers.df_to_table(cuda_api_trace_report))}')
Processing [/tmp/tmprqugcno6/example_api_tracing/concurrency.sqlite] with [/opt/nvidia/nsight-systems/2025.5.2/host-linux-x64/reports/cuda_api_trace.py] to [/tmp/tmprqugcno6/example_api_tracing/concurrency_cuda_api_trace.csv]... PROCESSED
Nsight Systems cuda_api_trace report:
┏━━━━━━━━━━━━┳━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━┳━━━━━━━━┳━━━━━┳━━━━━┳━━━━━━━┳━━━━━━━━━━━━━┓
┃ Start (us) ┃ Duration (us) ┃ Name ┃ Result ┃ CorrID ┃ Pid ┃ Tid ┃ T-Pri ┃ Thread Name ┃
┡━━━━━━━━━━━━╇━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━━╇━━━━━━━━╇━━━━━╇━━━━━╇━━━━━━━╇━━━━━━━━━━━━━┩
│ 25505.438 │ 96.14 │ cuLibraryLoadData │ 0 │ 1 │ 652 │ 652 │ 20 │ concurrency │
│ 25602.991 │ 0.4 │ cuLibraryGetKernel │ 0 │ 2 │ 652 │ 652 │ 20 │ concurrency │
│ 25607.239 │ 0.3 │ cuKernelGetName │ 0 │ 3 │ 652 │ 652 │ 20 │ concurrency │
│ 25608.892 │ 1989.212 │ cudaLaunchKernel │ 0 │ 4 │ 652 │ 652 │ 20 │ concurrency │
│ 27599.016 │ 5.851 │ cudaEventRecord │ 0 │ 5 │ 652 │ 652 │ 20 │ concurrency │
│ 27605.899 │ 0.34 │ cuLibraryGetKernel │ 0 │ 6 │ 652 │ 652 │ 20 │ concurrency │
│ 27606.84 │ 0.181 │ cuKernelGetName │ 0 │ 7 │ 652 │ 652 │ 20 │ concurrency │
│ 27607.141 │ 10.52 │ cudaLaunchKernel │ 0 │ 8 │ 652 │ 652 │ 20 │ concurrency │
│ 27618.041 │ 4.479 │ cudaStreamWaitEvent │ 0 │ 9 │ 652 │ 652 │ 20 │ concurrency │
│ 27622.74 │ 0.11 │ cuKernelGetName │ 0 │ 10 │ 652 │ 652 │ 20 │ concurrency │
│ 27622.961 │ 4.007 │ cudaLaunchKernel │ 0 │ 11 │ 652 │ 652 │ 20 │ concurrency │
│ 27627.279 │ 0.511 │ cudaEventRecord │ 0 │ 12 │ 652 │ 652 │ 20 │ concurrency │
│ 27627.96 │ 0.601 │ cudaStreamWaitEvent │ 0 │ 13 │ 652 │ 652 │ 20 │ concurrency │
│ 27628.741 │ 0.091 │ cuKernelGetName │ 0 │ 14 │ 652 │ 652 │ 20 │ concurrency │
│ 27628.942 │ 2.825 │ cudaLaunchKernel │ 0 │ 15 │ 652 │ 652 │ 20 │ concurrency │
│ 27634.131 │ 30.588 │ cudaMemcpyAsync │ 0 │ 16 │ 652 │ 652 │ 20 │ concurrency │
│ 27665.16 │ 3.476 │ cudaStreamSynchronize │ 0 │ 17 │ 652 │ 652 │ 20 │ concurrency │
└────────────┴───────────────┴───────────────────────┴────────┴────────┴─────┴─────┴───────┴─────────────┘
Querying CUDA API tracing data
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.
The Nsight Systems documentation provides the SQLite schema reference here.
The ReProspect class reprospect.tools.nsys.report.Report provides access to the database.
The class is used as a context manager, which opens and closes the connection to the database.
The property reprospect.tools.nsys.report.Report.tables provides a list of the names of the tables available in the database.
from reprospect.tools.nsys import Report
report = Report(db=ns.command.output.with_suffix('.sqlite'))
with report:
tables = report.tables
print(f'Tables:\n{tables}')
Tables:
['StringIds', 'ANALYSIS_FILE', 'ProcessStreams', 'TARGET_INFO_SYSTEM_ENV', 'TARGET_INFO_SESSION_START_TIME', 'ANALYSIS_DETAILS', 'DIAGNOSTIC_EVENT', 'NVTX_EVENTS', 'CUPTI_ACTIVITY_KIND_CUDA_EVENT', 'CUPTI_ACTIVITY_KIND_RUNTIME', 'CUPTI_ACTIVITY_KIND_OVERHEAD', 'CUPTI_ACTIVITY_KIND_KERNEL', 'CUPTI_ACTIVITY_KIND_SYNCHRONIZATION', 'CUPTI_ACTIVITY_KIND_MEMCPY', 'PROCESSES', 'ThreadNames', 'TARGET_INFO_GPU', 'TARGET_INFO_CUDA_DEVICE', 'TARGET_INFO_CUDA_CONTEXT_INFO', 'TARGET_INFO_CUDA_STREAM', 'ENUM_NSYS_EVENT_TYPE', 'ENUM_NSYS_EVENT_CLASS', 'ENUM_NSYS_GENERIC_EVENT_SOURCE', 'ENUM_NSYS_GENERIC_EVENT_GROUP', 'ENUM_NSYS_GENERIC_EVENT_FIELD_TYPE', 'ENUM_NSYS_GENERIC_EVENT_FIELD_ETW_PROPERTY', 'ENUM_NSYS_GENERIC_EVENT_FIELD_ETW_TYPE', 'ENUM_NSYS_GENERIC_EVENT_FIELD_ETW_FLAGS', 'ENUM_CUDA_MEMCPY_OPER', 'ENUM_CUDA_MEM_KIND', 'ENUM_CUDA_MEMPOOL_TYPE', 'ENUM_CUDA_MEMPOOL_OPER', 'ENUM_CUDA_DEV_MEM_EVENT_OPER', 'ENUM_CUDA_KERNEL_LAUNCH_TYPE', 'ENUM_CUDA_SHARED_MEM_LIMIT_CONFIG', 'ENUM_CUDA_UNIF_MEM_MIGRATION', 'ENUM_CUDA_UNIF_MEM_ACCESS_TYPE', 'ENUM_CUDA_FUNC_CACHE_CONFIG', 'ENUM_CUPTI_STREAM_TYPE', 'ENUM_CUPTI_SYNC_TYPE', 'ENUM_CUPTI_OVERHEAD_TYPE', 'ENUM_DIAGNOSTIC_SEVERITY_LEVEL', 'ENUM_DIAGNOSTIC_SOURCE_TYPE', 'ENUM_DIAGNOSTIC_TIMESTAMP_SOURCE', 'META_DATA_CAPTURE', 'META_DATA_EXPORT']
The method reprospect.tools.nsys.report.Report.table() retrieves a table as a pandas.DataFrame.
Let us retrieve the tables:
CUPTI_ACTIVITY_KIND_RUNTIMEwith the trace of the CUDA runtime API calls;CUPTI_ACTIVITY_KIND_CUDA_EVENTwith details relevant to the CUDA events;CUPTI_ACTIVITY_KIND_SYNCHRONIZATIONwith details relevant to the CUDA synchronization operations;CUPTI_ACTIVITY_KIND_KERNELwith details relevant to the kernel executions;ENUM_NSYS_EVENT_CLASS, which we will use below to interpret theeventClasscolumn of the runtime table;StringIds, which relates string identifiers to the strings themselves.
Here, CUPTI stands for the CUDA Profiling Tools Interface (CUPTI), which Nsight Systems uses to collect the associated data.
with report:
runtime = report.table(name='CUPTI_ACTIVITY_KIND_RUNTIME')
print(f'Table CUPTI_ACTIVITY_KIND_RUNTIME:\n{rich_helpers.to_string(rich_helpers.df_to_table(runtime))}')
cuda_event = report.table(name='CUPTI_ACTIVITY_KIND_CUDA_EVENT')
print(f'Table CUPTI_ACTIVITY_KIND_CUDA_EVENT:\n{rich_helpers.to_string(rich_helpers.df_to_table(cuda_event))}')
synchronization = report.table(name='CUPTI_ACTIVITY_KIND_SYNCHRONIZATION')
print(f'Table CUPTI_ACTIVITY_KIND_SYNCHRONIZATION:\n{rich_helpers.to_string(rich_helpers.df_to_table(synchronization))}')
kernel = report.table(name='CUPTI_ACTIVITY_KIND_KERNEL')
print(f'Table CUPTI_ACTIVITY_KIND_KERNEL, for brevity only its column names:\n{kernel.columns}')
enum_event_class = report.table(name='ENUM_NSYS_EVENT_CLASS')
print(f'Table ENUM_NSYS_EVENT_CLASS first two rows:\n{rich_helpers.to_string(rich_helpers.df_to_table(enum_event_class[:2]))}')
stringids = report.table(name='StringIds')
first_runtime_row = runtime.iloc[0]
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"]]))}')
Table CUPTI_ACTIVITY_KIND_RUNTIME:
┏━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━┳━━━━━━━━┳━━━━━━━━━━━━━┳━━━━━━━━━━━━━┓
┃ start ┃ end ┃ eventClass ┃ globalTid ┃ correlationId ┃ nameId ┃ returnValue ┃ callchainId ┃
┡━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━━━╇━━━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━━━╇━━━━━━━━╇━━━━━━━━━━━━━╇━━━━━━━━━━━━━┩
│ 25505438 │ 25601578 │ 1 │ 281485915456140 │ 1 │ 28 │ 0 │ None │
│ 25602991 │ 25603391 │ 1 │ 281485915456140 │ 2 │ 29 │ 0 │ None │
│ 25607239 │ 25607539 │ 1 │ 281485915456140 │ 3 │ 30 │ 0 │ None │
│ 25608892 │ 27598104 │ 0 │ 281485915456140 │ 4 │ 59 │ 0 │ None │
│ 27599016 │ 27604867 │ 0 │ 281485915456140 │ 5 │ 60 │ 0 │ None │
│ 27605899 │ 27606239 │ 1 │ 281485915456140 │ 6 │ 29 │ 0 │ None │
│ 27606840 │ 27607021 │ 1 │ 281485915456140 │ 7 │ 30 │ 0 │ None │
│ 27607141 │ 27617661 │ 0 │ 281485915456140 │ 8 │ 59 │ 0 │ None │
│ 27618041 │ 27622520 │ 0 │ 281485915456140 │ 9 │ 63 │ 0 │ None │
│ 27622740 │ 27622850 │ 1 │ 281485915456140 │ 10 │ 30 │ 0 │ None │
│ 27622961 │ 27626968 │ 0 │ 281485915456140 │ 11 │ 59 │ 0 │ None │
│ 27627279 │ 27627790 │ 0 │ 281485915456140 │ 12 │ 60 │ 0 │ None │
│ 27627960 │ 27628561 │ 0 │ 281485915456140 │ 13 │ 63 │ 0 │ None │
│ 27628741 │ 27628832 │ 1 │ 281485915456140 │ 14 │ 30 │ 0 │ None │
│ 27628942 │ 27631767 │ 0 │ 281485915456140 │ 15 │ 59 │ 0 │ None │
│ 27634131 │ 27664719 │ 0 │ 281485915456140 │ 16 │ 64 │ 0 │ None │
│ 27665160 │ 27668636 │ 0 │ 281485915456140 │ 17 │ 65 │ 0 │ None │
└──────────┴──────────┴────────────┴─────────────────┴───────────────┴────────┴─────────────┴─────────────┘
Table CUPTI_ACTIVITY_KIND_CUDA_EVENT:
┏━━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━━┳━━━━━━━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━┳━━━━━━━━━┳━━━━━━━━━━━━━┓
┃ timestamp ┃ deviceId ┃ contextId ┃ greenContextId ┃ streamId ┃ correlationId ┃ globalPid ┃ eventId ┃ eventSyncId ┃
┡━━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━━╇━━━━━━━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━━━━━╇━━━━━━━━━╇━━━━━━━━━━━━━┩
│ 0 │ 0 │ 1 │ 0 │ 14 │ 5 │ 281485915455488 │ 1 │ 1 │
│ 0 │ 0 │ 1 │ 0 │ 15 │ 12 │ 281485915455488 │ 2 │ 2 │
└───────────┴──────────┴───────────┴────────────────┴──────────┴───────────────┴─────────────────┴─────────┴─────────────┘
Table CUPTI_ACTIVITY_KIND_SYNCHRONIZATION:
┏━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━━┳━━━━━━━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━━━┳━━━━━━━━━━━━┳━━━━━━━━━━━━━┓
┃ start ┃ end ┃ deviceId ┃ contextId ┃ greenContextId ┃ streamId ┃ correlationId ┃ globalPid ┃ deprecatedSyncType ┃ syncType ┃ eventId ┃ eventSyncId ┃
┡━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━━╇━━━━━━━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━━━━━━━━╇━━━━━━━━━━╇━━━━━━━━━━━━╇━━━━━━━━━━━━━┩
│ 27618873 │ 27622259 │ 0 │ 1 │ None │ 15 │ 9 │ 281485915455488 │ None │ 2 │ 1 │ 1 │
│ 27628030 │ 27628481 │ 0 │ 1 │ None │ 14 │ 13 │ 281485915455488 │ None │ 2 │ 2 │ 2 │
│ 27665701 │ 27668486 │ 0 │ 1 │ None │ 14 │ 17 │ 281485915455488 │ None │ 2 │ 4294967295 │ 4294967295 │
└──────────┴──────────┴──────────┴───────────┴────────────────┴──────────┴───────────────┴─────────────────┴────────────────────┴──────────┴────────────┴─────────────┘
Table CUPTI_ACTIVITY_KIND_KERNEL, for brevity only its column names:
Index(['start', 'end', 'deviceId', 'contextId', 'greenContextId', 'streamId',
'correlationId', 'globalPid', 'demangledName', 'shortName',
'mangledName', 'launchType', 'cacheConfig', 'registersPerThread',
'gridX', 'gridY', 'gridZ', 'blockX', 'blockY', 'blockZ',
'staticSharedMemory', 'dynamicSharedMemory', 'localMemoryPerThread',
'localMemoryTotal', 'gridId', 'sharedMemoryExecuted', 'graphNodeId',
'sharedMemoryLimitConfig', 'qmdBulkReleaseDone', 'qmdPreexitDone',
'qmdLastCtaDone', 'graphId', 'clusterX', 'clusterY', 'clusterZ',
'clusterSchedulingPolicy', 'maxPotentialClusterSize',
'maxActiveClusters'],
dtype='str')
Table ENUM_NSYS_EVENT_CLASS first two rows:
┏━━━━┳━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━┓
┃ id ┃ name ┃ label ┃
┡━━━━╇━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━━━━━━━━┩
│ 0 │ TRACE_PROCESS_EVENT_CUDA_RUNTIME │ CUDA runtime │
│ 1 │ TRACE_PROCESS_EVENT_CUDA_DRIVER │ CUDA driver │
└────┴──────────────────────────────────┴──────────────┘
Table StringIds row related to first runtime table row:
┏━━━━┳━━━━━━━━━━━━━━━━━━━┓
┃ id ┃ value ┃
┡━━━━╇━━━━━━━━━━━━━━━━━━━┩
│ 28 │ cuLibraryLoadData │
└────┴───────────────────┘
As we can observe, analysing the data requires exploiting the relationships between the tables.
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.
Querying CUDA API tracing data by nested NVTX range
ReProspect provides functionalities to facilitate lookups in the SQLite database.
The method reprospect.tools.nsys.report.Report.get_events() provides focused lookups by nested NVTX range, with automatic correlation of string identifiers.
For tables featuring start and end columns, it selects the rows whose time span is within a given nested NVTX range.
The argument accessors designates the path of nested NVTX ranges, from the outermost inward.
with report:
print(report.nvtx_events)
NVTX events
├── example_api_tracing_domain (NvtxDomainCreate)
└── start_end_range (NvtxStartEndRange)
├── diamond_range (NvtxPushPopRange)
└── check_range (NvtxPushPopRange)
with report:
diamond_runtime = report.get_events(table='CUPTI_ACTIVITY_KIND_RUNTIME', accessors=['start_end_range', 'diamond_range'])
print(f'Runtime events in diamond range:\n{diamond_runtime}')
Runtime events in diamond range:
start end eventClass globalTid correlationId nameId \
0 25505438 25601578 1 281485915456140 1 28
1 25602991 25603391 1 281485915456140 2 29
2 25607239 25607539 1 281485915456140 3 30
3 25608892 27598104 0 281485915456140 4 59
4 27599016 27604867 0 281485915456140 5 60
5 27605899 27606239 1 281485915456140 6 29
6 27606840 27607021 1 281485915456140 7 30
7 27607141 27617661 0 281485915456140 8 59
8 27618041 27622520 0 281485915456140 9 63
9 27622740 27622850 1 281485915456140 10 30
10 27622961 27626968 0 281485915456140 11 59
11 27627279 27627790 0 281485915456140 12 60
12 27627960 27628561 0 281485915456140 13 63
13 27628741 27628832 1 281485915456140 14 30
14 27628942 27631767 0 281485915456140 15 59
returnValue callchainId name
0 0 None cuLibraryLoadData
1 0 None cuLibraryGetKernel
2 0 None cuKernelGetName
3 0 None cudaLaunchKernel_v7000
4 0 None cudaEventRecord_v3020
5 0 None cuLibraryGetKernel
6 0 None cuKernelGetName
7 0 None cudaLaunchKernel_v7000
8 0 None cudaStreamWaitEvent_v3020
9 0 None cuKernelGetName
10 0 None cudaLaunchKernel_v7000
11 0 None cudaEventRecord_v3020
12 0 None cudaStreamWaitEvent_v3020
13 0 None cuKernelGetName
14 0 None cudaLaunchKernel_v7000
Some runtime activity names carry a suffix such as _v7000 and _v3020.
The function reprospect.tools.nsys.report.strip_cuda_api_suffix() allows such suffixes to be stripped, so that the analysis does not depend on them.
The runtime table may also contain CUDA driver activity, such as cuLibraryLoadData, cuLibraryGetKernel, cuKernelGetName related to module loading and name resolution.
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.
from reprospect.tools.nsys import strip_cuda_api_suffix
RUNTIME_EVENT_CLASS_ID = Report.single_row(data=enum_event_class[enum_event_class['label'] == 'CUDA runtime'])['id']
diamond_cuda_runtime = (
diamond_runtime.loc[diamond_runtime['eventClass'] == RUNTIME_EVENT_CLASS_ID]
.assign(api=lambda df: df['name'].map(strip_cuda_api_suffix))
)
print(f'CUDA runtime api calls in diamond range:\n{diamond_cuda_runtime["api"].to_list()}')
CUDA runtime api calls in diamond range:
['cudaLaunchKernel', 'cudaEventRecord', 'cudaLaunchKernel', 'cudaStreamWaitEvent', 'cudaLaunchKernel', 'cudaEventRecord', 'cudaStreamWaitEvent', 'cudaLaunchKernel']
The method reprospect.tools.nsys.report.Report.get_correlated_row() facilitates correlated lookups across tables in the database.
By default, it uses the correlationId column to join both tables.
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:
kernel_a = Report.get_correlated_row(src=diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaLaunchKernel'].iloc[0], dst=kernel)
print('StreamId for kernel A:', kernel_a['streamId'])
kernel_a_demangled_name = Report.get_correlated_row(src=kernel_a, dst=stringids, correlation_src='demangledName', correlation_dst='id')
print(f'Demangled name for kernel A: {kernel_a_demangled_name["value"]}')
StreamId for kernel A: 14
Demangled name for kernel A: check_add_kernel(int *, int, int)
Assertions on CUDA API tracing data
With the collected data gathered in Python data structures, the analysis can now go all the way to test assertions. Here, we programmatically verify the concurrency structure.
First, we verify the sequence of the CUDA runtime API calls. 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.
assert diamond_cuda_runtime['api'].to_list() == [
'cudaLaunchKernel',
'cudaEventRecord',
'cudaLaunchKernel',
'cudaStreamWaitEvent',
'cudaLaunchKernel',
'cudaEventRecord',
'cudaStreamWaitEvent',
'cudaLaunchKernel',
]
Next, we verify the streams the kernels are launched on. 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.
launches = diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaLaunchKernel']
kernel_a, kernel_b, kernel_c, kernel_d = (Report.get_correlated_row(src=launches.iloc[idx], dst=kernel) for idx in range(4))
assert kernel_b['streamId'] != kernel_c['streamId']
assert kernel_a['streamId'] == kernel_b['streamId'] == kernel_d['streamId']
Finally, we verify the dependencies that form the diamond. 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.
event_records = diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaEventRecord']
stream_wait_events = diamond_cuda_runtime[diamond_cuda_runtime['api'] == 'cudaStreamWaitEvent']
event_record_fork = Report.get_correlated_row(src=event_records.iloc[0], dst=cuda_event)
stream_wait_event_fork = Report.get_correlated_row(src=stream_wait_events.iloc[0], dst=synchronization)
assert event_record_fork['streamId'] == kernel_a['streamId']
assert stream_wait_event_fork['streamId'] == kernel_c['streamId']
assert stream_wait_event_fork['eventId'] == event_record_fork['eventId']
event_record_join = Report.get_correlated_row(src=event_records.iloc[1], dst=cuda_event)
stream_wait_event_join = Report.get_correlated_row(src=stream_wait_events.iloc[1], dst=synchronization)
assert event_record_join['streamId'] == kernel_c['streamId']
assert stream_wait_event_join['streamId'] == kernel_d['streamId']
assert stream_wait_event_join['eventId'] == event_record_join['eventId']
Outlook
In the example of this notebook, the traced sequence can be read off the source code directly.
One context in which the proposed approach becomes most valuable is that of abstraction layers.
Portability libraries such as Kokkos let applications express computations at a higher level of abstraction and map them to CUDA API calls internally.
Likewise, the C++26 std::execution model, whose customization for CUDA is under development in NVIDIA’s CCCL library, lowers declarative descriptions of asynchronous work, including fork–join structures like the diamond of this notebook, to streams and events.
API tracing then allows verifying that these mappings produce the intended calls.
With ReProspect’s fully programmatic approach, such verifications can run as tests in CI/CD pipelines.
API tracing is used in our 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.