Profile CUDA workloads with NVIDIA Nsight¶

Use NVIDIA Nsight Systems (nsys) to analyze application-wide CPU and GPU activity, and NVIDIA Nsight Compute (ncu) to inspect individual CUDA kernels.

A typical profiling workflow is:

1. Run a CUDA workload to confirm that the GPU environment works.
2. Use nsys to identify where the application spends time.
3. Use ncu to investigate individual CUDA kernels.

Verda GPU instances are configured by default to allow non-admin users to access NVIDIA GPU hardware performance counters, so standard Nsight profiling does not require additional system configuration.

Prerequisites¶

You need:

• A Verda GPU instance created from a current image (Ubuntu 24 or 26, with CUDA 12.9 or newer)

• SSH access to the instance

• SSH access to the instance

Current Verda images come with CUDA, nvcc, Nsight Systems, and Nsight Compute preinstalled. Confirm this on your instance:

nvidia-smi
nvcc --version
nsys --version
ncu --version

The CUDA tools are installed in /usr/local/cuda/bin, which is on the PATH for root but not for a newly created user. If you're signed in as a user other than root, add it first, before running the commands above:

export PATH=/usr/local/cuda/bin:$PATH

To keep it across sessions, add the same line to ~/.bashrc.

Create a test CUDA workload¶

Use this small CUDA program to verify the GPU and profiling tools before profiling your own application. The program allocates a million-plus floats (2^20 elements) on the GPU, then launches a kernel that adds 1.0 to every element, repeating the launch 100 times to give the profiling tools enough activity to measure.

Create the source file:

cat > verda_profile_test.cu <<'EOF'
#include <cuda_runtime.h>
#include <cstdio>

__global__ void profile_kernel(float* data, int count) {
    const int index = blockIdx.x * blockDim.x + threadIdx.x;

    if (index < count) {
        data[index] += 1.0f;
    }
}

int main() {
    const int count = 1 << 20;
    const size_t bytes = count * sizeof(float);
    float* data = nullptr;

    cudaError_t error = cudaMalloc(&data, bytes);
    if (error != cudaSuccess) {
        std::fprintf(stderr, "cudaMalloc failed: %s\n",
                     cudaGetErrorString(error));
        return 1;
    }

    error = cudaMemset(data, 0, bytes);
    if (error != cudaSuccess) {
        std::fprintf(stderr, "cudaMemset failed: %s\n",
                     cudaGetErrorString(error));
        cudaFree(data);
        return 1;
    }

    for (int launch = 0; launch < 100; ++launch) {
        profile_kernel<<<(count + 255) / 256, 256>>>(data, count);
    }

    error = cudaGetLastError();

    if (error == cudaSuccess) {
        error = cudaDeviceSynchronize();
    }

    cudaFree(data);

    if (error != cudaSuccess) {
        std::fprintf(stderr, "CUDA failure: %s\n",
                     cudaGetErrorString(error));
        return 1;
    }

    std::puts("CUDA sample completed successfully");
    return 0;
}
EOF

Compile the application for the GPU in the instance:

nvcc -O2 -arch=native verda_profile_test.cu -o verda_profile_test

Run the application:

./verda_profile_test

Profile with Nsight Systems¶

Nsight Systems provides an application-wide timeline of CPU and GPU activity. Use it first when the location of the performance bottleneck is unknown.

Profile the test workload:

nsys profile \
  --trace=cuda \
  --sample=none \
  --cpuctxsw=none \
  --force-overwrite=true \
  -o verda-timeline \
  ./verda_profile_test

Print the CUDA API and kernel summaries:

nsys stats \
  --force-export=true \
  --report cuda_api_sum \
  --report cuda_gpu_kern_sum \
  verda-timeline.nsys-rep

A successful profile lists profile_kernel, the kernel used by the test application, in the CUDA GPU Kernel Summary.

Example output:

 ** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):

 Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)              Name
 --------  ---------------  ---------  --------  --------  --------  --------  -----------  ----------------------------
    100.0           319137        100    3191.4    3200.0      3104      3296         48.9  profile_kernel(float *, int)

The two reports cover different halves of the picture:

• cuda_api_sum is time spent in CUDA API calls on the CPU side. The first CUDA call in a process also carries context initialization, which is why cudaMalloc dominates it in a short run.
• cuda_gpu_kern_sum is time spent executing kernels on the GPU. The top row is the kernel to investigate first.

In both, Time (%) is the share of that report's total and times are in nanoseconds by default. For the full column reference and the other reports nsys stats can produce, see the Nsight Systems Post-Collection Analysis Guide.

You can also open the .nsys-rep file produced above in the Nsight Systems GUI for deeper visual timeline analysis than the command-line summary provides. Since the GUI isn't installed on the instance, copy the file to your local machine first, with scp or any file transfer tool, then open it there.

To profile your own workload, replace ./verda_profile_test with your application command. Use the kernel names shown in the Nsight Systems report to select kernels for detailed analysis with Nsight Compute.

Profile with Nsight Compute¶

Nsight Compute collects detailed performance metrics for individual CUDA kernels.

Profile one launch of the test kernel:

ncu \
  --kernel-name 'regex:profile_kernel' \
  --launch-count 1 \
  -o verda-kernel \
  --force-overwrite \
  ./verda_profile_test

Example output:

==PROF== Connected to process 2360 (/home/user/verda_profile_test)
==PROF== Profiling "profile_kernel": 0%....50%....100% - 10 passes
CUDA sample completed successfully
==PROF== Disconnected from process 2360
==PROF== Report: /home/user/verda-kernel.ncu-rep

Nsight Compute may replay the selected kernel several times to collect hardware performance counters, which is the pass count shown above.

Inspect the saved report from the command line:

ncu \
  --import verda-kernel.ncu-rep \
  --page details

The details page prints one section per area of the kernel's behavior:

• GPU Speed Of Light Throughput is the high-level view: what share of the device's compute and memory peak the kernel reached.
• Launch Statistics is the launch configuration, including grid size, block size, registers per thread and shared memory.
• Occupancy compares active warps per multiprocessor against the theoretical maximum. The gap between theoretical and achieved occupancy is usually the first thing to look at.
• GPU and Memory Workload Distribution shows how evenly the work spreads across SMs, caches and DRAM.

Nsight Compute also prints OPT recommendations under most sections, each with an estimated speedup. For what every section measures, see Nsight Compute sections and rules.

You can also open the .ncu-rep file in the Nsight Compute GUI, copied across the same way as the Nsight Systems report.

To profile your own workload, replace profile_kernel with a kernel name identified using Nsight Systems and replace ./verda_profile_test with your application command. Using --launch-count 1 prevents Nsight Compute from profiling every matching kernel launch in a large workload.

Advanced: change profiling access¶

The setting behind this is the nvidia kernel module parameter NVreg_RestrictProfilingToAdminUsers, reported as RmProfilingAdminOnly in /proc/driver/nvidia/params. On Verda GPU instances it is set to 0, so tools such as Nsight Compute work without sudo.

If a non-admin user sees ERR_NVGPUCTRPERM when profiling, hardware-counter access is restricted on that instance. Follow Allow non-admin profiling to enable it.

Allow non-admin profiling¶

To enable non-admin hardware-counter access:

echo 'options nvidia NVreg_RestrictProfilingToAdminUsers=0' \
  | sudo tee /etc/modprobe.d/nvidia-profiling.conf

sudo reboot

After reconnecting, verify the active setting:

grep RmProfilingAdminOnly /proc/driver/nvidia/params

Restrict profiling to administrators¶

To restrict hardware performance-counter access:

echo 'options nvidia NVreg_RestrictProfilingToAdminUsers=1' \
  | sudo tee /etc/modprobe.d/nvidia-profiling.conf

sudo reboot

After reconnecting, verify the active setting:

grep RmProfilingAdminOnly /proc/driver/nvidia/params

Setting this value to 1 restricts hardware performance-counter access but does not prevent root or another sufficiently privileged user from profiling.

Related documentation¶

• NVIDIA Nsight Systems User Guide
• NVIDIA Nsight Compute CLI
• NVIDIA Nsight Compute system requirements
• NVIDIA Nsight Compute GPU support
• NVIDIA GPU performance-counter permissions
• CUDA C++ Best Practices Guide: Application profiling