CUPTI: the profiling interface under Nsight, PyTorch and DCGM¶
Scope: CUPTI as the layer every GPU profiler is built on, and therefore as a thing you administer rather than call. Who is allowed to read the GPU's performance counters (NVreg_RestrictProfilingToAdminUsers, and the capabilities system replacing it), why only one tool can hold those counters at a time (which is why dcgm-exporter and an engineer's nsys run fight), what tracing costs, and how to turn an activity buffer into a number that is not a lie. For the tools built on it, see Nsight profiling workflow; for the management library that is not part of this stack, see NVML.
What it is¶
"The CUDA Profiling Tools Interface (CUPTI) provides C and Python APIs for building profiling and tracing tools for CUDA applications, providing detailed insight into how code executes on both the CPU and GPU."1 It ships with the CUDA Toolkit, not the driver: "The CUPTI SDK is part of the CUDA Toolkit, and will be installed along all the other CUDA libraries", under /usr/local/cuda-<version>/extras/CUPTI/, and it is not added to LD_LIBRARY_PATH for you.1
The docs draw a line that governs everything downstream, and it is worth internalizing before touching permissions or overhead:
"tracing means collecting timestamps and associated metadata for CUDA activities such as runtime and driver API calls, kernel launches, and memory copy operations during application execution."
"profiling refers to collecting GPU performance metrics for individual kernels or sets of kernels, often in isolation from the rest of the workload. Profiling may require replaying kernels or even the full application multiple times to gather all desired metrics under controlled conditions."1
Tracing is cheap, unprivileged, and mostly composable. Profiling is expensive, privileged, and exclusive. Nearly every operational surprise in this page follows from that one distinction.
The current API set:1
| API | What it does |
|---|---|
| Activity | "Asynchronously record CUDA activities, e.g. CUDA API, Kernel, memory copy" |
| Callback | "CUDA event callback mechanism to notify subscriber that a specific CUDA event executed e.g. 'Entering CUDA runtime memory copy'" |
| Host Profiling | "Host APIs for enumeration, configuration and evaluation of performance metrics" |
| Range Profiling | "Target APIs for collection of performance metrics for a range of execution" |
| PC Sampling | "Sampling of the warp program counter and warp scheduler state (stall reasons)" |
| SASS Metrics | "Collect kernel performance metrics at the source level using SASS patching" |
| PM Sampling | "Collect hardware metrics by sampling the GPU performance monitors (PM) periodically at fixed intervals" |
| Checkpoint | "Provides support for automatically saving and restoring the functional state of the CUDA device" |
| Profiling | Deprecated in CUDA 13.0; use Range Profiling |
The old Event API and Metric API are gone, not deprecated: in CUDA 13.0 "The CUPTI Event API from the header cupti_events.h and the CUPTI Metric API from the header cupti_metrics.h are dropped. Calling any Event or Metric API will return the error code CUPTI_ERROR_LEGACY_PROFILER_NOT_SUPPORTED."3 Any tool still on them stops working on CUDA 13, and the failure is a clean error code rather than a crash, which is at least a courtesy.
Who is built on it. Nsight Systems: "Nsight Systems uses CUPTI for CUDA profiling, including to collect the CUDA memory usage by the application processes."10 The frameworks: CUPTI's own troubleshooting text names "a framework with CUPTI integration (like PyTorch or TensorFlow)" as a CUPTI client,1 and PyTorch's kineto states "NVIDIA CUPTI: used to collect traces and metrics from NVIDIA GPUs."11 HPCToolkit wraps it too, and it is the standard integration point for third-party HPC profilers. Nsight Compute is the interesting exception: no NVIDIA source says it is CUPTI-based, and its docs treat "a client of CUPTI's Profiling API" as a separate process competing for the same reservation, so do not assume it.9
Why use it¶
You will most likely never write a cuptiSubscribe. You will still administer CUPTI, because every torch.profiler call, every nsys run, and every DCGM profiling metric on your cluster goes through its permission model and its exclusivity rules. Two consequences dominate the ticket queue:
Profiling is admin-only by default, and the error message does not say why. A user runs nsys profile or ncu and gets:
"ERR_NVGPUCTRPERM: The user running \<tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device."6
This is not a bug, it is a security mitigation, applied since "driver versions 419.17+ on Windows or 418.43+ on Linux."6
Only one client can hold the counters. DCGM says it plainly: "Due to current hardware limitations, collection of profiling metrics with DCGM will conflict with usage of other developer tools from NVIDIA such as Nsight Systems or Nsight Compute."8 If you run dcgm-exporter with profiling metrics enabled on your nodes, you have already taken the lock. Be careful about which error you go looking for, because the two sides report differently: the engineer's tool says the reservation failed (Nsight Compute: "Profiling failed because a driver resource was unavailable", and a CUPTI client gets CUPTI_ERROR_HARDWARE_BUSY, "The performance monitoring hardware is currently in use by another client"), while the Error setting watches ... the affected resource is in use string is what DCGM prints when a developer tool got there first.918
When to use it (and when not)¶
| Goal | Reach for | Not |
|---|---|---|
| Find the bottleneck in a training step | nsys / torch.profiler |
Writing a CUPTI tool; they are already CUPTI tools |
| Continuous, low-cost timeline data in your own harness | CUPTI Activity API | Range Profiling; counters are not free |
| Per-kernel counters (SM active, tensor active, memory throughput) | Range Profiling API (or ncu) |
The Event/Metric APIs; dropped in CUDA 13.0 |
| Fleet-wide utilization metrics | DCGM profiling metrics | Per-node CUPTI tools; DCGM already holds the counters |
| Device state, ECC, Xid, power | NVML | CUPTI; it does not do device management |
| Correctness bugs | Compute Sanitizer | CUPTI, which cannot even run at the same time |
Do not collect counters on a shared production GPU. Counter collection may serialize kernels and replay them, and it takes an exclusive reservation that blocks your own monitoring. Tracing is the mode that is safe to leave running; profiling is a scheduled, isolated activity.
Three hard incompatibilities to design around:
- Not with the debugger or the sanitizer. "CUPTI still does not support tracing or profiling along with other NVIDIA Developer Tools like cuda-gdb or Compute Sanitizer."1 This is the same wall from the other side of cuda-gdb's rule that core dumps are unsupported when another tool is attached.
- Not on a shared MIG compute instance. For a shared compute instance, "Due to this resource sharing, collecting profiling data from shared units is not permitted. Attempts to collect metrics from a shared unit will result in NaN values." Tracing still works, and profiling works on an isolated compute instance.5
- Not under MPS, at least for the (now deprecated) Profiling API, which is "not supported under MPS (Multi-Process Service), Confidential Compute, or SLI configured systems."1 Note the scope honestly: that sentence sits in the Profiling API section, and NVIDIA states no equivalent MPS restriction for the Range Profiling API that replaces it, so do not conclude counter collection under MPS is categorically impossible. Verify on your driver.
Architecture¶
flowchart TB
APP["CUDA application<br/>(PyTorch, vLLM, your kernel)"] --> CUPTI["CUPTI (CUDA Toolkit)"]
CUPTI --> TRACE["TRACE path<br/>Activity + Callback APIs<br/>cheap, unprivileged, composable"]
CUPTI --> PROF["PROFILE path<br/>Range Profiling / PM Sampling<br/>replay, serialize, exclusive"]
TRACE --> BUF["activity buffers<br/>(unordered, correlationId-linked)"]
PROF --> PERM{"NVreg_RestrictProfilingToAdminUsers<br/>or nvidia-capabilities (R610+)"}
PERM -->|"denied"| ERR["ERR_NVGPUCTRPERM"]
PERM -->|"allowed"| PM["GPU performance monitors"]
DCGM["DCGM (dcgm-exporter)"] --> PM
NCU["Nsight Compute"] --> PM
PM --> LOCK["one reservation holder at a time"]The bottom of the diagram is the whole operational problem: dcgm-exporter, Nsight Compute, and any CUPTI Profiling client all want the same single-holder resource. The trace path, above it, does not.
How to administer it: the permission model¶
This is the part that belongs to the platform team, not the user.
The classic control is a kernel module parameter. To open profiling to all users:6
# Legacy regkey method. Create /etc/modprobe.d/nvidia-profiling.conf:
options nvidia NVreg_RestrictProfilingToAdminUsers=0
A reboot "may be required" for the change to take effect, and on some systems the initramfs must be rebuilt or the file is ignored: dracut --regenerate-all -f on Red Hat derivatives, update-initramfs -u -k all on Debian derivatives.6 Verify what the running module actually has, rather than what your .conf says:
One caveat on newer drivers, straight from the same page: "On R610+ drivers, this flag only reflects the regkey state and does not account for capability-based access grants."6 So a 1 there does not prove your users cannot profile; they may hold a capability instead.
What is the default? NVIDIA's own header contradicts itself, so trust the code. In open-gpu-kernel-modules at driver 610.43.03, the comment above the parameter says "0: Do not restrict GPU counters (default)", while the actual definition, five hundred lines further down the same file, is NV_DEFINE_REG_ENTRY(__NV_RM_PROFILING_ADMIN_ONLY_PARAMETER, 1);, and that macro expands to a variable initialized to its second argument.7 The compiled-in default is 1, restricted, which is also what the ERR_NVGPUCTRPERM page implies when it says a recent driver installation "has disabled access to GPU Performance Counters for regular users". The comment is stale. If you grep the header and conclude profiling is open by default, you have been misled by NVIDIA's documentation, not by your cluster.
The alternatives to opening it globally:
- Capabilities. A user with
CAP_SYS_ADMINmay profile; "Starting in driver version R565, the CAP_PERFMON capability will also allow access."6CAP_PERFMONis the narrower grant and generally the better one, with one documented exception: "CAP_PERFMON will not work in secure execution mode unless profiling within a container as described above."6 - Containers. "When profiling within a container, access must be enabled on the host, or the container must be started with the appropriate permissions by passing
--cap-add=SYS_ADMINas an admin user."6 - The R610+ replacement. Newer drivers move to per-capability device nodes:
profiler-context("Grants profiling access to intra-context scope"),profiler-device(adds device-level scope), andtrace-device. NVIDIA states the regkey method is legacy and that "In a future release, the regkey-based method will be removed."6 Note the scoping limit: "Profiling capabilities are granted system-wide and cannot be scoped to individual GPUs."6 You cannot grant profiling on GPU 3 only.
The right posture for a multi-tenant cluster is usually: leave counters restricted on production nodes, and grant profiling (via CAP_PERFMON or a capability-enabled debug pod) on a dedicated pool where the serialization and exclusivity cost nobody else anything.
How to develop with it: a real client, then honest arithmetic¶
The Activity API is asynchronous and buffer-based. You register two callbacks with cuptiActivityRegisterCallbacks, one that hands CUPTI an empty buffer and one that receives a full one, and "For typical workloads, it's suggested to choose a size between 1 and 10 MB."1 Two properties of what arrives will break a naive tool:
- The buffer is unordered. "CUPTI doesn't guarantee any ordering of the activities in the activity buffer as activity records for few activity kinds are added lazily."1 Any logic that assumes records arrive in time order is wrong.
- The link back to the launch site is a correlation ID, and it is opt-in. "Every CUDA driver/runtime API invocation that CUPTI records is assigned a unique correlation ID", the resulting kernel and memcpy records "carry the same correlation ID", and critically: "For CUPTI to generate
correlationIdvalues, the CUDA API activity kinds i.e.CUPTI_ACTIVITY_KIND_RUNTIMEand/orCUPTI_ACTIVITY_KIND_DRIVERmust be enabled."1 Forget that, and you get kernel records that cannot be attributed to any line of code.
A minimal Activity API client, built against the real headers¶
Everything in this subsection is taken from the headers shipped in Ubuntu 24.04's libcupti-dev 12.0.146~12.0.1-4build4 (/usr/include/cupti.h, /usr/include/cupti_activity.h, /usr/include/cupti_result.h, /usr/include/cupti_version.h) and verified by compiling the client below against them, not from memory. The buffer contract, quoted from cupti_activity.h:13
- Records are packed and 8-byte aligned, so the buffers you hand CUPTI should be too. The header defines
ACTIVITY_RECORD_ALIGNMENT 8and declares every recordPACKED_ALIGNMENT, which on GCC expands to__attribute__ ((__packed__)) __attribute__ ((aligned (ACTIVITY_RECORD_ALIGNMENT))); "all activity record types are padded and aligned to ensure that each member of the record is naturally aligned." NVIDIA's packaged sample (activity_trace_async.cpp, installed bylibcupti-doc) allocatesBUF_SIZE + ALIGN_SIZEwith#define ALIGN_SIZE (8)and rounds the pointer up by hand; C11aligned_alloc(ACTIVITY_RECORD_ALIGNMENT, BUF_SIZE)achieves the same alignment without the pointer arithmetic. - You may decline a buffer request, at a price. "The callback function can decline the request by setting
*bufferto NULL. In this case CUPTI may drop activity records." Returning*maxNumRecords = 0means "the buffer is filled with as many records as possible." - A completed buffer is yours again, and it is global. "After this call CUPTI relinquished ownership of the buffer and will not use it anymore." The
contextandstreamIdparameters of the completion callback are "deprecated as of CUDA 6.0 and will always be NULL"; there is no per-stream buffer, you parse per-stream records out of a global one. - Iteration terminates by error code, not by count.
cuptiActivityGetNextRecordreturnsCUPTI_ERROR_MAX_LIMIT_REACHED"if no more records in the buffer"; that return is the loop exit, not a failure, and it is the one CUPTI error a client must not treat as fatal. - Dropped records are counted, never signaled. Nothing interrupts you when CUPTI drops records (because your request callback declined, or your buffers were too small or returned too slowly). The completion callback is where "The number of dropped records can be read using
cuptiActivityGetNumDroppedRecords", and a client that skips that call reports clean traces with holes in them. - Record struct names are versioned; read your own header. The kind-to-struct mapping lives in the
CUpti_ActivityKindenum comments: in this header,CUPTI_ACTIVITY_KIND_CONCURRENT_KERNELand_KERNELmap toCUpti_ActivityKernel9,_MEMCPYtoCUpti_ActivityMemcpy5, and_RUNTIME/_DRIVERtoCUpti_ActivityAPI.CUpti_ActivityKernel2throughKernel9all coexist in the same file; copying a struct name from a blog post written against a different CUPTI is how a client silently reads garbage fields.
The complete client. Registration, four activity kinds, aligned buffers, record iteration, dropped-record accounting, and the flush/disable/finalize shutdown, with every CUPTI call behind an error macro:
/* Minimal CUPTI Activity API client: asynchronous buffer callbacks,
* record iteration, dropped-record accounting, flush/disable/finalize
* shutdown. Written against libcupti-dev 12.0.146~12.0.1-4build4
* (headers in /usr/include, CUPTI_API_VERSION 18). */
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <inttypes.h>
#include <cuda_runtime_api.h>
#include <cupti.h>
/* Every CUPTI call goes through this. Fail fast on the first error and
* print the descriptive string from cuptiGetResultString. */
#define CUPTI_CALL(call) \
do { \
CUptiResult _status = (call); \
if (_status != CUPTI_SUCCESS) { \
const char *_msg = "unknown"; \
cuptiGetResultString(_status, &_msg); \
fprintf(stderr, "%s:%d: %s failed: %s (%d)\n", \
__FILE__, __LINE__, #call, _msg, (int)_status); \
exit(EXIT_FAILURE); \
} \
} while (0)
/* 1 MB per buffer; the CUPTI docs suggest 1 to 10 MB for typical
* workloads. ACTIVITY_RECORD_ALIGNMENT is 8, from cupti_activity.h. */
#define BUF_SIZE (1024 * 1024)
static uint64_t g_records = 0;
static uint64_t g_dropped = 0;
/* CUPTI asks for an empty buffer. Records are packed and 8-byte
* aligned (PACKED_ALIGNMENT in cupti_activity.h), so hand CUPTI an
* 8-byte aligned allocation, as NVIDIA's activity_trace_async sample
* does with ALIGN_SIZE (8). Setting *buffer to NULL declines the
* request and CUPTI may then drop records. */
static void CUPTIAPI buffer_requested(uint8_t **buffer, size_t *size,
size_t *maxNumRecords) {
*buffer = (uint8_t *)aligned_alloc(ACTIVITY_RECORD_ALIGNMENT, BUF_SIZE);
*size = (*buffer != NULL) ? BUF_SIZE : 0;
*maxNumRecords = 0; /* 0 = fill with as many records as fit */
}
static void print_record(const CUpti_Activity *record) {
switch (record->kind) {
case CUPTI_ACTIVITY_KIND_KERNEL:
case CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL: {
const CUpti_ActivityKernel9 *k = (const CUpti_ActivityKernel9 *)record;
printf("KERNEL corr=%u dev=%u stream=%u %" PRIu64 "..%" PRIu64
" ns grid=(%d,%d,%d) %s\n",
k->correlationId, k->deviceId, k->streamId, k->start, k->end,
k->gridX, k->gridY, k->gridZ, k->name);
break;
}
case CUPTI_ACTIVITY_KIND_MEMCPY: {
const CUpti_ActivityMemcpy5 *m = (const CUpti_ActivityMemcpy5 *)record;
printf("MEMCPY corr=%u dev=%u stream=%u %" PRIu64 "..%" PRIu64
" ns copyKind=%u bytes=%" PRIu64 "\n",
m->correlationId, m->deviceId, m->streamId, m->start, m->end,
(unsigned)m->copyKind, m->bytes);
break;
}
case CUPTI_ACTIVITY_KIND_RUNTIME:
case CUPTI_ACTIVITY_KIND_DRIVER: {
const CUpti_ActivityAPI *a = (const CUpti_ActivityAPI *)record;
printf("%s cbid=%u corr=%u %" PRIu64 "..%" PRIu64 " ns\n",
record->kind == CUPTI_ACTIVITY_KIND_RUNTIME ? "RUNTIME" : "DRIVER ",
a->cbid, a->correlationId, a->start, a->end);
break;
}
default:
printf("OTHER kind=%d\n", (int)record->kind);
break;
}
}
/* CUPTI returns a full (or force-flushed) buffer. Iterate it with
* cuptiActivityGetNextRecord until CUPTI_ERROR_MAX_LIMIT_REACHED, then
* account for drops with cuptiActivityGetNumDroppedRecords. Ownership
* of the buffer is back with the client; free it here. */
static void CUPTIAPI buffer_completed(CUcontext ctx, uint32_t streamId,
uint8_t *buffer, size_t size,
size_t validSize) {
(void)size; /* total capacity; validSize is what matters */
CUpti_Activity *record = NULL;
if (validSize > 0) {
for (;;) {
CUptiResult status =
cuptiActivityGetNextRecord(buffer, validSize, &record);
if (status == CUPTI_ERROR_MAX_LIMIT_REACHED)
break; /* no more records in this buffer */
CUPTI_CALL(status);
print_record(record);
g_records++;
}
}
size_t dropped = 0;
CUPTI_CALL(cuptiActivityGetNumDroppedRecords(ctx, streamId, &dropped));
if (dropped > 0) {
fprintf(stderr, "WARNING: %zu activity records dropped\n", dropped);
g_dropped += dropped;
}
free(buffer);
}
/* On a GPU host this produces MEMCPY and RUNTIME records (a kernel
* launch needs device code compiled by nvcc; the client is unchanged).
* The first CUDA runtime error is returned, not swallowed, so a
* driverless host reports its real boundary. */
static cudaError_t run_workload(void) {
void *dev = NULL;
static uint8_t host[4096];
cudaError_t err = cudaMalloc(&dev, sizeof(host));
if (err != cudaSuccess)
return err;
err = cudaMemcpy(dev, host, sizeof(host), cudaMemcpyHostToDevice);
if (err == cudaSuccess)
err = cudaMemcpy(host, dev, sizeof(host), cudaMemcpyDeviceToHost);
if (err == cudaSuccess)
err = cudaDeviceSynchronize();
cudaFree(dev);
return err;
}
int main(void) {
uint32_t version = 0;
CUPTI_CALL(cuptiGetVersion(&version));
printf("CUPTI API version: %u (built against CUPTI_API_VERSION %d)\n",
version, CUPTI_API_VERSION);
/* Register buffer callbacks and enable activity kinds BEFORE any
* CUDA call, so initialization itself is traced. RUNTIME/DRIVER
* must be on or kernel/memcpy records carry no correlationId. */
CUPTI_CALL(cuptiActivityRegisterCallbacks(buffer_requested,
buffer_completed));
CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_MEMCPY));
CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME));
CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_DRIVER));
cudaError_t err = run_workload();
if (err != cudaSuccess)
fprintf(stderr, "CUDA workload failed: %s (%s)\n",
cudaGetErrorName(err), cudaGetErrorString(err));
/* Shutdown lifecycle: force-flush so buffers with incomplete
* records are still delivered, disable every kind that was
* enabled, then detach CUPTI from the process. */
CUPTI_CALL(cuptiActivityFlushAll(CUPTI_ACTIVITY_FLAG_FLUSH_FORCED));
CUPTI_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
CUPTI_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_MEMCPY));
CUPTI_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_RUNTIME));
CUPTI_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_DRIVER));
CUPTI_CALL(cuptiFinalize());
printf("done: %" PRIu64 " records, %" PRIu64 " dropped\n",
g_records, g_dropped);
return (err == cudaSuccess) ? EXIT_SUCCESS : EXIT_FAILURE;
}
This compiles for real, warnings as errors, against CUPTI from libcupti-dev 12.0.146~12.0.1-4build4 on Ubuntu 24.04 (CUDA 12.0 toolkit). The exact command that succeeded:
gcc -std=c11 -Wall -Wextra -Werror -o cupti_activity_client cupti_activity_client.c \
-I/usr/include -L/usr/lib/x86_64-linux-gnu -lcupti -lcudart
On Ubuntu the distro packages put the headers in /usr/include and the library in /usr/lib/x86_64-linux-gnu, so the -I/-L are technically redundant there; on a system installed with NVIDIA's own installer the equivalents are -I/usr/local/cuda/extras/CUPTI/include -L/usr/local/cuda/extras/CUPTI/lib64, and you must also put that lib directory on LD_LIBRARY_PATH yourself, because the toolkit does not.1 ldd on the produced binary resolves libcupti.so.12 and libcudart.so.12. Link -lcuda as well only if you call the driver API directly; this client does not.
The shutdown lifecycle is three separate steps, and each covers a different hole. First, cuptiActivityFlushAll(CUPTI_ACTIVITY_FLAG_FLUSH_FORCED): a default flush (flag 0) "doesn't return buffers which have one or more incomplete records", so the forced flush is what delivers the tail of the trace, and the header suggests exactly this, "before the termination of the profiling session to allow remaining buffers to be delivered. In general, it can be done in the at-exit handler." Note the trap in the same doc comment: the flush "doesn't issue any CUDA synchronization calls implicitly", so records for still-running GPU work are not guaranteed complete unless you synchronize first, which is why the workload above ends in cudaDeviceSynchronize(). Second, cuptiActivityDisable per enabled kind, the symmetric close of every cuptiActivityEnable. Third, cuptiFinalize(), which "detaches the CUPTI from the running process" and frees its resources; the header recommends invoking it "from the exit callsite of any of the CUDA Driver or Runtime API" callback when detaching mid-run, and otherwise the client owns the synchronize-and-flush ordering before calling it.13
What actually happens with no GPU. The host this page was built on has the CUDA 12.0 toolkit but no NVIDIA kernel module and no /dev/nvidia* nodes, so the GPU path of this client was compiled but not executed here. The run itself is real; this is the complete output, verbatim (exit status 1):
CUPTI API version: 18 (built against CUPTI_API_VERSION 18)
cupti_activity_client.c:137: cuptiActivityRegisterCallbacks(buffer_requested, buffer_completed) failed: CUPTI_ERROR_NOT_INITIALIZED (15)
Two findings from that boundary are worth keeping. First, cuptiGetVersion is a pure host-side call and succeeds with no driver at all, and then the very first stateful call, cuptiActivityRegisterCallbacks (line 137 is the registration in main), fails with the same CUPTI_ERROR_NOT_INITIALIZED the maintenance section describes for a too-old driver; cupti_result.h defines code 15 as "CUPTI is unable to initialize its connection to the CUDA driver."13 A CUPTI client on a broken node therefore dies at registration, before any CUDA call, not at the first kernel launch. Second, a version quirk to not trip over: the comment table in cupti_version.h ends "v19 : CUDA Toolkit 12.0", yet the define this CUDA 12.0 package actually ships is CUPTI_API_VERSION 18, and cuptiGetVersion agrees at runtime. Gate on the value cuptiGetVersion returns, not on that comment table.13
Turning the records into a busy number¶
What follows validates the consumer-side arithmetic only; it is a Python model of buffer contents and calls no CUPTI. The C client above is what exercises CUPTI itself, at compile time against the real headers and at runtime up to the documented no-driver boundary.
Now the arithmetic. The question every trace consumer asks first is "how busy was the GPU", and the obvious implementation, summing kernel durations, is wrong the moment two streams overlap. It can report more than 100% busy, which should be the tell. The correct quantity is the union of the kernel intervals, and it must be computed without assuming buffer order.
"""Turning a CUPTI activity buffer into GPU busy time and launch-gap analysis.
Two documented properties of the buffer drive every assertion here:
1. "CUPTI doesn't guarantee any ordering of the activities in the activity buffer"
2. a kernel record is tied to its launching API call only by correlationId, and only
if CUPTI_ACTIVITY_KIND_RUNTIME / _DRIVER were enabled.
Timestamps are in nanoseconds.
"""
import numpy as np
# (correlationId, stream, start_ns, end_ns) for CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL
KERNELS = np.array([
(11, 7, 1_000, 5_000), # gemm on stream 7
(12, 7, 5_200, 9_000), # gemm on stream 7, after a 200 ns launch gap
(13, 13, 3_000, 4_000), # allreduce on stream 13, OVERLAPS the first gemm
(14, 13, 4_000, 6_000), # allreduce on stream 13, touches 13's end, overlaps 12
(15, 7, 9_000, 9_000), # a kernel that recorded a zero-length duration
], dtype=np.int64)
# (correlationId, start_ns, end_ns) for CUPTI_ACTIVITY_KIND_RUNTIME (the launch calls)
RUNTIME = np.array([
(11, 900, 950),
(12, 5_050, 5_100),
(13, 2_900, 2_950),
(14, 3_950, 3_990),
(15, 8_900, 8_950),
], dtype=np.int64)
def union_busy(iv):
"""Total wall time during which at least one kernel was resident. Order-independent."""
if len(iv) == 0:
return 0
iv = iv[np.argsort(iv[:, 0], kind="stable")] # sort by start; the buffer is unordered
total, cur_s, cur_e = 0, iv[0, 0], iv[0, 1]
for s, e in iv[1:]:
if s > cur_e: # a real gap: close the run
total += cur_e - cur_s
cur_s, cur_e = s, e
else:
cur_e = max(cur_e, e) # overlap or touch: extend
return int(total + cur_e - cur_s)
def naive_busy(iv):
"""What a per-kernel sum reports. Double counts every overlap."""
return int((iv[:, 1] - iv[:, 0]).sum())
iv = KERNELS[:, 2:4]
span = int(iv[:, 1].max() - iv[:, 0].min()) # wall clock across the trace
# --- Happy path: one stream, no overlap, sum and union agree ---------------
serial = KERNELS[KERNELS[:, 1] == 7][:, 2:4]
serial = serial[:2] # the two gemms, disjoint
assert union_busy(serial) == naive_busy(serial) == 7_800
print(f"single stream, disjoint kernels: union == sum == {union_busy(serial)} ns")
# --- Adversarial 1: two streams overlap; the naive sum exceeds wall clock --
u, n = union_busy(iv), naive_busy(iv)
assert u == 8_000, u # 1000..9000 is continuously busy
assert n == 10_800, n # 4000 + 3800 + 1000 + 2000 + 0
assert n > span == 8_000 # a "GPU busy" of 135% of the window
assert u <= span # the union never can
print(f"2 streams: union={u} ns, naive sum={n} ns, wall span={span} ns "
f"-> naive reports {100 * n / span:.0f}% busy, which is impossible")
# --- Adversarial 2: the buffer is unordered, so the result must not depend on order
rng = np.random.default_rng(7)
for _ in range(200):
shuffled = iv[rng.permutation(len(iv))]
assert union_busy(shuffled) == u
print(f"200 random buffer orderings: union invariant at {u} ns")
# --- Adversarial 3: touching intervals are not a gap, and zero-length is not busy
touching = np.array([[0, 100], [100, 200]], dtype=np.int64)
assert union_busy(touching) == 200 # 3000-4000 and 4000-6000 above
zero = np.array([[9_000, 9_000]], dtype=np.int64)
assert union_busy(zero) == 0 # kernel 15 adds nothing
assert union_busy(np.vstack([iv, zero])) == u # and cannot inflate the total
print("touching intervals merge (no phantom gap); a zero-length kernel adds 0 ns")
# --- Adversarial 4: idle gaps are the launch overhead you came to find -----
merged, order = [], iv[np.argsort(iv[:, 0], kind="stable")]
cs, ce = order[0]
for s, e in order[1:]:
if s > ce:
merged.append((cs, ce)); cs, ce = s, e
else:
ce = max(ce, e)
merged.append((cs, ce))
gaps = [int(merged[i + 1][0] - merged[i][1]) for i in range(len(merged) - 1)]
assert merged == [(1_000, 9_000)], merged # one continuous busy run
assert gaps == [] # nothing idle inside the window
assert span - u == 0
# Now drop the allreduce stream: the gemm-only timeline exposes the 200 ns launch gap.
gemm = KERNELS[(KERNELS[:, 1] == 7) & (KERNELS[:, 3] > KERNELS[:, 2])][:, 2:4]
gemm_span = int(gemm[:, 1].max() - gemm[:, 0].min())
gemm_idle = gemm_span - union_busy(gemm)
assert gemm_idle == 200, gemm_idle # 5000 -> 5200
print(f"stream 7 alone: {gemm_idle} ns idle inside an {gemm_span} ns window "
f"({100 * gemm_idle / gemm_span:.1f}% of it), the launch gap the union exposes")
# --- Adversarial 5: correlationId is the only link back to the launch site -
def attribute(kernels, api_records):
"""Join kernel records to their launching API call. Must tolerate missing records."""
by_corr = {int(c): (int(s), int(e)) for c, s, e in api_records}
launched, orphans = {}, []
for corr, _, k_start, _ in kernels:
rec = by_corr.get(int(corr)) # .get, not [], is the whole point
if rec is None:
orphans.append(int(corr))
continue
api_end = rec[1]
assert api_end <= k_start # the launch precedes the execution
launched[int(corr)] = int(k_start) - api_end # launch-to-execute queue latency
return launched, orphans
latency, orphans = attribute(KERNELS, RUNTIME)
assert latency == {11: 50, 12: 100, 13: 50, 14: 10, 15: 50} and orphans == []
# A buffer can be delivered with only some API records present: the client may have
# enabled only CUPTI_ACTIVITY_KIND_DRIVER, or dropped records by returning buffers late.
partial = RUNTIME[[0, 2, 4]] # correlations 11, 13, 15 survive
latency_p, orphans_p = attribute(KERNELS, partial)
assert sorted(orphans_p) == [12, 14] # these kernels have no launch site
assert set(latency_p) == {11, 13, 15}
missing = None
try: # the naive join a tool actually ships
{int(c): dict((int(a), (int(s), int(e))) for a, s, e in partial)[int(c)] for c in KERNELS[:, 0]}
raise AssertionError("expected a KeyError from the unguarded join")
except KeyError as exc:
missing = int(str(exc))
assert missing == 12
print(f"correlationId join: {len(latency)} kernels attributed (queue latencies "
f"{sorted(latency.values())} ns); with API records partly missing, {len(orphans_p)} "
f"kernels orphan and an unguarded dict lookup raises KeyError({missing})")
Executed output:
single stream, disjoint kernels: union == sum == 7800 ns
2 streams: union=8000 ns, naive sum=10800 ns, wall span=8000 ns -> naive reports 135% busy, which is impossible
200 random buffer orderings: union invariant at 8000 ns
touching intervals merge (no phantom gap); a zero-length kernel adds 0 ns
stream 7 alone: 200 ns idle inside an 8000 ns window (2.5% of it), the launch gap the union exposes
correlationId join: 5 kernels attributed (queue latencies [10, 50, 50, 50, 100] ns); with API records partly missing, 2 kernels orphan and an unguarded dict lookup raises KeyError(12)
The 135% is the point. A tool that sums kernel durations across streams reports a GPU busier than wall clock, which is impossible, and nobody notices because the number merely looks high rather than absurd. The same union, inverted, is the useful measurement: the gaps between merged intervals are the launch bubbles, and on stream 7 alone they are 2.5% of the window. That is the quantity comms-compute overlap work is trying to drive to zero.
How to maintain it: versions and the driver pairing¶
CUPTI is versioned with the CUDA Toolkit, and it is one of the few components you may need to update between toolkit releases: "Normally packaged with the CUDA Toolkit, NVIDIA occasionally uses this page to provide CUPTI improvements and bug fixes between toolkit releases."2 cuptiGetVersion reports the loaded version at runtime.
The compatibility direction that is documented is old tool, new driver: CUPTI "adheres to CUDA Toolkit compatibility requirements with CUDA drivers, which includes support for Backward, Forward and Enhanced compatibilities", and "A profiling tool based on an older version of CUPTI can still operate with a more recent CUDA driver."1 The reverse direction, a newer CUPTI against an older driver, is not promised in the same sentence; the docs send you to the toolkit's minimum-driver table, and "Attempting to use CUPTI calls with an incompatible CUDA driver version will result in a CUPTI_ERROR_NOT_INITIALIZED error code."1 The CUDA 13.3 multiple-subscriber feature is a concrete instance: a new CUPTI capability that requires driver r610+ regardless of toolkit.
Maintenance work that is worth scheduling rather than discovering during an incident:
- Audit for the dropped APIs before a CUDA 13 rollout. Anything still calling the Event or Metric API returns
CUPTI_ERROR_LEGACY_PROFILER_NOT_SUPPORTEDand collects nothing.3 This includes vendored or unmaintained internal profilers, which is where it usually hides. - Pin the driver alongside the toolkit the same way you pin everything else on the node (driver install and lifecycle). A CUPTI feature that silently no-ops on an older driver is worse than one that fails loudly.
- Re-check the permission mechanism when you move to R610 or later. The regkey still works, but NVIDIA has said it goes away, and the capabilities system that replaces it is configured differently.6
How to run it in production¶
Pause your fleet monitoring before anyone profiles. DCGM's own guidance, with the error you will see first:8
# The symptom, from dcgmi or the API:
# Error setting watches. Result: The requested operation could not be completed
# because the affected resource is in use.
dcgmi profile --pause # release the counters
# ... run nsys / ncu / your CUPTI tool ...
dcgmi profile --resume
The programmatic equivalents are dcgmProfPause() and dcgmProfResume(). Note the consequence for your dashboards: "When paused, DCGM will publish BLANK values for profiling metrics", testable with DCGM_FP64_IS_BLANK(value).8 A profiling session therefore punches a hole in your utilization graphs, and an alert rule that treats a missing value as zero will fire. Nsight Compute suggests exactly this workflow from its side: "If you expect the problem to be caused by DCGM, consider using dcgmi profile --pause to stop its monitoring while profiling with NVIDIA Nsight Compute."9
Understand what serializes. Not all tracing is equal:
CUPTI_ACTIVITY_KIND_KERNEL(serial kernel trace) "can significantly change the overall performance characteristics of the application because all kernel executions are serialized on the GPU."4CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL"doesn't affect the concurrency of the kernels in the application", but CUPTI instruments the kernel code to do it, so it costs more per launch.4
For a serving or training workload, CONCURRENT_KERNEL is the correct choice: serializing kernels does not just slow the app, it changes the thing you are trying to measure. Nsight Compute, separately, "serializes kernel launches within the profiled application, potentially across multiple processes profiled by one or more instances of the tool at the same time", because "some GPU and driver objects can only be acquired by a single process at a time... meaning only one process can profile a given device at a time."9
Expect no single headline overhead number; expect per-feature ones. CUPTI's own statement is qualitative: overhead "largely depends on the density of the CUDA activities in the application", and "In general overhead of tracing i.e. activity APIs is much lesser than the metrics profiling APIs."4 Individual features do carry published figures, though: Nsight Systems warns that "Collecting Unified Memory CPU page faults can cause overhead of up to 70% in testing. Use this functionality only when needed."10 And HPCToolkit, measuring CUPTI directly, reports two: "we have seen the execution time double when profiling and tracing a GPU-accelerated application that launches kernels very frequently", and separately that for PC sampling "the overhead is less than 5x".12 So the honest budget is 2x for launch-heavy tracing, and up to an order of magnitude once you turn on sampling or counters. Measure your own; do not extrapolate one figure across features.
Multiple tracers: check your CUDA version before you promise it. This changed recently and the documentation has not fully caught up. As of CUDA 13.3, "CUPTI supports multiple concurrent subscribers to trace CUDA activities simultaneously", requiring "CUDA 13.3+ and driver r610+".1 Before that, a second subscriber got CUPTI_ERROR_MULTIPLE_SUBSCRIBERS_NOT_SUPPORTED: with Nsight Systems and PyTorch both documented as CUPTI clients, that is the mechanism behind the familiar failure of nesting one inside the other.1 The limits still bite: it "applies only to the CUPTI Activity API. It is not applicable to CUPTI Profiling APIs", and V1 and V2 API styles cannot be mixed in one process.1 Be aware that CUPTI's own troubleshooting section still describes the single-subscriber rule as current, contradicting its multiple-subscribers section in the same document; the multiple-subscribers section is the newer text.1
Failure modes¶
| Symptom | Cause | Fix |
|---|---|---|
ERR_NVGPUCTRPERM on nsys / ncu / torch.profiler counters |
Counter access is restricted to admins by default since driver 418.436 | NVreg_RestrictProfilingToAdminUsers=0 plus initramfs rebuild, or grant CAP_PERFMON, or use the R610+ capabilities |
The .conf file is right but nothing changed |
The initramfs still carries the old parameters, or no reboot6 | dracut --regenerate-all -f or update-initramfs -u -k all, then confirm RmProfilingAdminOnly in /proc/driver/nvidia/params |
| You read the driver header and concluded profiling is open by default | The comment says default 0; the code compiles in 17 | Trust /proc/driver/nvidia/params, not the comment |
Your profiler reports "a driver resource was unavailable", or CUPTI returns CUPTI_ERROR_HARDWARE_BUSY |
Something else holds the counter reservation, commonly dcgm-exporter91 |
dcgmi profile --pause, profile, dcgmi profile --resume |
DCGM logs Error setting watches ... the affected resource is in use |
The reverse case: a developer tool holds the counters and DCGM cannot watch them8 | Finish the profiling session; DCGM resumes on its own watches |
| Utilization panels go blank during a profiling session | DCGM publishes BLANK values while paused8 |
Expect it; do not let an alert rule read blank as zero |
Metrics come back as NaN under MIG |
The compute instance is shared, and profiling shared units is not permitted5 | Profile on an isolated compute instance; tracing works either way |
Profiler and cuda-gdb or Compute Sanitizer refuse to coexist |
"CUPTI still does not support tracing or profiling along with other NVIDIA Developer Tools like cuda-gdb or Compute Sanitizer."1 | Run one tool at a time |
| A tool built on the Event/Metric API dies on CUDA 13 | Both APIs were dropped; calls return CUPTI_ERROR_LEGACY_PROFILER_NOT_SUPPORTED3 |
Port to the Range Profiling API |
| The app got much slower under tracing and the timeline looks serial | CUPTI_ACTIVITY_KIND_KERNEL serializes all kernel execution4 |
Use CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL |
| "GPU busy" exceeds 100% | Kernel durations were summed across overlapping streams (executed model above) | Take the union of intervals, not the sum |
| Kernel records cannot be traced back to any code | CUPTI_ACTIVITY_KIND_RUNTIME / _DRIVER were not enabled, so no correlation IDs1 |
Enable the CUDA API activity kinds |
Every stateful CUPTI call fails with CUPTI_ERROR_NOT_INITIALIZED, starting at cuptiActivityRegisterCallbacks |
"CUPTI is unable to initialize its connection to the CUDA driver": no kernel module or device nodes, or a driver too old for this CUPTI131 | Fix the node's driver install (driver install and lifecycle); reproduced verbatim above on a driverless host |
| The trace has silent holes | Buffer requests were declined or buffers came back too slowly, and nothing checked the drop counter | Call cuptiActivityGetNumDroppedRecords in every completion callback (client above) |
References¶
- NVIDIA, CUPTI documentation (v2026.2.1, last updated 2026-06-12): https://docs.nvidia.com/cupti/
- NVIDIA, CUPTI product page (packaging, out-of-band updates, feature list): https://developer.nvidia.com/cupti
- NVIDIA, CUPTI Usage (tracing vs profiling, activity buffers, correlation IDs, multiple subscribers, overhead): https://docs.nvidia.com/cupti/main/main.html
- NVIDIA, CUPTI release notes (Event/Metric API dropped in CUDA 13.0): https://docs.nvidia.com/cupti/release-notes/release-notes.html
- NVIDIA, CUPTI special configurations (MIG, vGPU, WSL): https://docs.nvidia.com/cupti/special-configurations/special-configurations.html
- NVIDIA, ERR_NVGPUCTRPERM: permission issue with performance counters: https://developer.nvidia.com/nvidia-development-tools-solutions-ERR_NVGPUCTRPERM-permission-issue-performance-counters
- NVIDIA, DCGM feature overview: concurrent usage of NVIDIA profiling tools: https://docs.nvidia.com/datacenter/dcgm/latest/user-guide/feature-overview.html
- NVIDIA, Nsight Compute profiling guide (serialization, driver-resource reservation, DCGM conflict): https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html
- NVIDIA, Nsight Systems user guide (uses CUPTI for CUDA profiling): https://docs.nvidia.com/nsight-systems/UserGuide/index.html
- NVIDIA, open-gpu-kernel-modules (
kernel-open/nvidia/nv-reg.h, driver 610.43.03): https://github.com/NVIDIA/open-gpu-kernel-modules/blob/main/kernel-open/nvidia/nv-reg.h - NVIDIA CUPTI headers and samples as packaged by Ubuntu 24.04 (noble),
libcupti-dev/libcupti-doc12.0.146~12.0.1-4build4, built from thenvidia-cuda-toolkit12.0.1-4build4 source package (/usr/include/cupti_activity.h,cupti_result.h,cupti_version.h; sampleactivity_trace_async.cpp), read locally 2026-07-17: https://launchpad.net/ubuntu/+source/nvidia-cuda-toolkit - PyTorch, kineto (uses CUPTI to collect traces and metrics): https://github.com/pytorch/kineto/blob/main/libkineto/README.md
- HPCToolkit, GPU measurement documentation (observed 2x execution time on launch-heavy code): https://hpctoolkit.gitlab.io/hpctoolkit/users/gpu/gpu.html
Related: Profiling GPUs: Nsight Systems and Nsight Compute · NVML · CUDA-GDB · Compute Sanitizer · GPU Diagnostics and Validation · Observability and Monitoring · Comms-Compute Overlap · PyTorch Performance Regression CI
-
NVIDIA, CUPTI documentation (v2026.2.1). Note which page carries which text: the CUPTI definition, the API table, the tracing-versus-profiling definitions, "The CUPTI SDK is part of the CUDA Toolkit, and will be installed along all the other CUDA libraries", and the
/usr/local/cuda-<version>/extras/CUPTI/path are on the documentation landing page (https://docs.nvidia.com/cupti/), not on the Usage page. Verified on the Usage page: the activity-buffer callbacks and "For typical workloads, it's suggested to choose a size between 1 and 10 MB"; "CUPTI doesn't guarantee any ordering of the activities in the activity buffer as activity records for few activity kinds are added lazily"; the Correlation ID section, including "For CUPTI to generatecorrelationIdvalues, the CUDA API activity kinds i.e.CUPTI_ACTIVITY_KIND_RUNTIMEand/orCUPTI_ACTIVITY_KIND_DRIVERmust be enabled"; the Profiling API being "not supported under MPS (Multi-Process Service), Confidential Compute, or SLI configured systems"; the framework-integration reference to "a framework with CUPTI integration (like PyTorch or TensorFlow)"; and the Multiple Subscribers section ("Starting with CUDA 13.3, CUPTI supports multiple concurrent subscribers to trace CUDA activities simultaneously", requiring "CUDA 13.3+ and driver r610+", "applies only to the CUPTI Activity API", and "CUPTI still does not support tracing or profiling along with other NVIDIA Developer Tools like cuda-gdb or Compute Sanitizer"); and the Compatibility and Requirements section ("CUPTI adheres to CUDA Toolkit compatibility requirements with CUDA drivers, which includes support for Backward, Forward and Enhanced compatibilities", "A profiling tool based on an older version of CUPTI can still operate with a more recent CUDA driver", the pointer to the CUDA Toolkit and Corresponding Driver Versions table, and "Attempting to use CUPTI calls with an incompatible CUDA driver version will result in aCUPTI_ERROR_NOT_INITIALIZEDerror code"). https://docs.nvidia.com/cupti/main/main.html · Internal inconsistency, flagged rather than resolved: the Troubleshooting section of the same document still describesCUPTI_ERROR_MULTIPLE_SUBSCRIBERS_NOT_SUPPORTEDas meaning CUPTI "only allows one callback subscriber... to be active at a time within a process", which contradicts the Multiple Subscribers section. The Multiple Subscribers text is the newer of the two. ↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩ -
NVIDIA, CUPTI product page: "Normally packaged with the CUDA Toolkit, NVIDIA occasionally uses this page to provide CUPTI improvements and bug fixes between toolkit releases." The feature list includes "Normalized timestamps for CPU and GPU trace". https://developer.nvidia.com/cupti ↩
-
NVIDIA, CUPTI release notes: CUDA 12.8 deprecated the Event API (
cupti_events.h) and Metric API (cupti_metrics.h); CUDA 13.0 dropped them, and "Calling any Event or Metric API will return the error codeCUPTI_ERROR_LEGACY_PROFILER_NOT_SUPPORTED. It is recommended to use the CUPTI Range Profiling API as an alternative." The PC Sampling Activity API fromcupti_activity.hwas also dropped in 13.0. https://docs.nvidia.com/cupti/release-notes/release-notes.html ↩↩↩ -
NVIDIA, CUPTI Usage, overhead: "Overhead can vary significantly from one application to another. It largely depends on the density of the CUDA activities in the application"; "In general overhead of tracing i.e. activity APIs is much lesser than the metrics profiling APIs"; serial kernel trace via
CUPTI_ACTIVITY_KIND_KERNEL"can significantly change the overall performance characteristics of the application because all kernel executions are serialized on the GPU", whileCUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL"doesn't affect the concurrency of the kernels in the application" at the cost of code instrumentation. https://docs.nvidia.com/cupti/main/main.html ↩↩↩↩ -
NVIDIA, CUPTI special configurations, MIG: an isolated compute instance supports tracing and profiling; for a shared compute instance, "Due to this resource sharing, collecting profiling data from shared units is not permitted. Attempts to collect metrics from a shared unit will result in NaN values." "Tracing works for shared Compute Instances." https://docs.nvidia.com/cupti/special-configurations/special-configurations.html ↩↩
-
NVIDIA, ERR_NVGPUCTRPERM solution page: the error text; the restriction applying to "driver versions 419.17+ on Windows or 418.43+ on Linux";
options nvidia NVreg_RestrictProfilingToAdminUsers=0in a.conffile under/etc/modprobe.d; thedracut --regenerate-all -fandupdate-initramfs -u -k allinitramfs rebuilds; theRmProfilingAdminOnlyflag in/proc/driver/nvidia/params;CAP_SYS_ADMIN, and "Starting in driver version R565, the CAP_PERFMON capability will also allow access"; "When profiling within a container, access must be enabled on the host, or the container must be started with the appropriate permissions by passing--cap-add=SYS_ADMINas an admin user"; and the R610+nvidia-capabilitiessystem (profiler-context,profiler-device,trace-device), including "Profiling capabilities are granted system-wide and cannot be scoped to individual GPUs" and "In a future release, the regkey-based method will be removed." https://developer.nvidia.com/nvidia-development-tools-solutions-ERR_NVGPUCTRPERM-permission-issue-performance-counters ↩↩↩↩↩↩↩↩↩↩↩↩↩ -
NVIDIA,
open-gpu-kernel-modules,kernel-open/nvidia/nv-reg.h, read atNVIDIA_VERSION = 610.43.03. The doc comment above the parameter reads "0: Do not restrict GPU counters (default)", while the definition isNV_DEFINE_REG_ENTRY(__NV_RM_PROFILING_ADMIN_ONLY_PARAMETER, 1);and the macro expands tostatic NvU32 __NV_REG_VAR(regkey) = (default_value);. The compiled-in default is therefore 1 (restricted) and the comment is stale. https://github.com/NVIDIA/open-gpu-kernel-modules/blob/main/kernel-open/nvidia/nv-reg.h ↩↩ -
NVIDIA, DCGM feature overview, Concurrent Usage of NVIDIA Profiling Tools: "Due to current hardware limitations, collection of profiling metrics with DCGM will conflict with usage of other developer tools from NVIDIA such as Nsight Systems or Nsight Compute."; the error "The requested operation could not be completed because the affected resource is in use.";
dcgmi profile --pause/--resumeand thedcgmProfPause()/dcgmProfResume()APIs; "When paused, DCGM will publish BLANK values for profiling metrics." https://docs.nvidia.com/datacenter/dcgm/latest/user-guide/feature-overview.html ↩↩↩↩↩↩ -
NVIDIA, Nsight Compute profiling guide: "NVIDIA Nsight Compute serializes kernel launches within the profiled application, potentially across multiple processes profiled by one or more instances of the tool at the same time."; "Serialization across processes is necessary since for the collection of HW performance metrics, some GPU and driver objects can only be acquired by a single process at a time. This is done on a per-CUDA device or MIG instance basis, meaning only one process can profile a given device at a time."; the driver-resource error names "DCGM, a client of CUPTI's Profiling API, Nsight Graphics, or another instance of NVIDIA Nsight Compute" as possible reservation holders; "If you expect the problem to be caused by DCGM, consider using
dcgmi profile --pause..." Note that no NVIDIA source states Nsight Compute is itself built on CUPTI, and these passages treat a CUPTI Profiling API client as a distinct competing process; this page therefore does not claim it is. https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html ↩↩↩↩↩ -
NVIDIA, Nsight Systems user guide: "Nsight Systems uses CUPTI for CUDA profiling, including to collect the CUDA memory usage by the application processes." The same guide publishes a per-feature overhead figure: "Collecting Unified Memory CPU page faults can cause overhead of up to 70% in testing. Use this functionality only when needed." https://docs.nvidia.com/nsight-systems/UserGuide/index.html ↩↩
-
PyTorch, kineto
libkineto/README.md: "NVIDIA CUPTI: used to collect traces and metrics from NVIDIA GPUs." https://github.com/pytorch/kineto/blob/main/libkineto/README.md ↩ -
HPCToolkit GPU documentation: "we have seen the execution time double when profiling and tracing a GPU-accelerated application that launches kernels very frequently", and, on PC sampling, "Our experience with CUPTI's serialization-based API for PC sampling is that the overhead is less than 5x." These are HPCToolkit's own measurements, not NVIDIA figures. https://hpctoolkit.gitlab.io/hpctoolkit/users/gpu/gpu.html ↩
-
NVIDIA CUPTI headers as packaged by Ubuntu 24.04,
libcupti-dev12.0.146~12.0.1-4build4, read locally on 2026-07-17 (Doxygen\p/\refmarkers elided from quotes). From/usr/include/cupti_activity.h:#define ACTIVITY_RECORD_ALIGNMENT 8withPACKED_ALIGNMENTexpanding on GCC to__attribute__ ((__packed__)) __attribute__ ((aligned (ACTIVITY_RECORD_ALIGNMENT))), and "Note that all activity record types are padded and aligned to ensure that each member of the record is naturally aligned."; the buffer-request contract ("The callback function can decline the request by setting *buffer to NULL. In this case CUPTI may drop activity records."; formaxNumRecords, "If 0 then the buffer is filled with as many records as possible"); the buffer-completion contract ("The buffer contains validSize bytes of activity records which should be read using cuptiActivityGetNextRecord. The number of dropped records can be read using cuptiActivityGetNumDroppedRecords. After this call CUPTI relinquished ownership of the buffer and will not use it anymore.", with thecontext/streamIdparameters marked "deprecated as of CUDA 6.0 and will always be NULL");cuptiActivityGetNextRecordreturning "CUPTI_ERROR_MAX_LIMIT_REACHED if no more records in the buffer"; thecuptiActivityFlushAllsemantics (default flush "doesn't return buffers which have one or more incomplete records", forced flush "suggested for clients to do the force flush before the termination of the profiling session to allow remaining buffers to be delivered. In general, it can be done in the at-exit handler.", and "This is a blocking call but it doesn't issue any CUDA synchronization calls implicitly");cuptiFinalize("This API detaches the CUPTI from the running process", "it is recommended this API is invoked from the exit callsite of any of the CUDA Driver or Runtime API"); and the kind-to-struct mapping in theCUpti_ActivityKindenum comments (CUPTI_ACTIVITY_KIND_KERNELand_CONCURRENT_KERNEL: "The corresponding activity record structure is CUpti_ActivityKernel9";_MEMCPY:CUpti_ActivityMemcpy5;_DRIVERand_RUNTIME:CUpti_ActivityAPI). From/usr/include/cupti_result.h:CUPTI_ERROR_NOT_INITIALIZED = 15, documented as "CUPTI is unable to initialize its connection to the CUDA driver." From/usr/include/cupti_version.h: the version comment table ending "v18 : CUDA Toolkit 11.8 / v19 : CUDA Toolkit 12.0" sitting directly above#define CUPTI_API_VERSION 18in this CUDA 12.0 package. The packaged sample/usr/share/doc/libcupti-dev/examples/activity_trace_async/activity_trace_async.cpp(shipped in these Ubuntu packages) aligns client buffers with#define ALIGN_SIZE (8)over amalloc(BUF_SIZE + ALIGN_SIZE)allocation. Compile and run evidence, both performed 2026-07-17 on this page's build host (Ubuntu 24.04, CUDA 12.0 toolkit, no NVIDIA kernel module, no/dev/nvidia*): the client compiled cleanly with the exactgcc -std=c11 -Wall -Wextra -Werror ... -lcupti -lcudartcommand shown, and the run produced exactly the two output lines quoted, exiting 1. The GPU path was not executed; no GPU run is claimed. ↩↩↩↩↩