Nsight Systems & Compute: Profiling Your CUDA Kernels

# Basic profile - generates .nsys-rep file
nsys profile --output=my_app_profile ./my_cuda_app

# Profile with specific options
nsys profile \
  --trace=cuda,nvtx,osrt \          # What to trace
  --sample=cpu \                     # CPU sampling
  --cudabacktrace=all \              # CUDA API backtraces
  --output=detailed_profile \
  ./my_cuda_app

# Profile for specific duration
nsys profile --duration=10 ./my_cuda_app

# Profile with NVTX ranges (see below)
nsys profile --trace=cuda,nvtx ./my_cuda_app

# Generate stats summary
nsys stats my_app_profile.nsys-rep

#include <nvtx3/nvToolsExt.h>

void train_epoch() {
    // Create a named range for this epoch
    nvtxRangePush("Epoch");
    
    for (int batch = 0; batch < num_batches; batch++) {
        // Nested range for each phase
        nvtxRangePush("Forward Pass");
        forward<<>>(input, output);
        nvtxRangePop();
        
        nvtxRangePush("Backward Pass");
        backward<<>>(grad_output, grad_input);
        nvtxRangePop();
        
        nvtxRangePush("Weight Update");
        update_weights<<>>(weights, gradients, lr);
        nvtxRangePop();
    }
    
    nvtxRangePop();  // End Epoch range
}

// Python with PyTorch
import torch.cuda.nvtx as nvtx

with nvtx.range("training_loop"):
    for batch in dataloader:
        with nvtx.range("forward"):
            output = model(batch)
        with nvtx.range("backward"):
            loss.backward()

# Profile all kernels (basic)
ncu --set full -o kernel_report ./my_cuda_app

# Profile specific kernel by name
ncu --kernel-name "matmul_kernel" --set full ./my_cuda_app

# Profile specific kernel launch (by index)
ncu --launch-skip 10 --launch-count 5 ./my_cuda_app

# Collect specific metrics
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
dram__throughput.avg.pct_of_peak_sustained_elapsed \
./my_cuda_app

# Generate CSV for analysis
ncu --csv --set full ./my_cuda_app > metrics.csv

# Compare two implementations
ncu --set full -o baseline ./my_app_v1
ncu --set full -o optimized ./my_app_v2
ncu --page raw --diff baseline.ncu-rep optimized.ncu-rep

// Naive matrix multiply - each thread computes one element
__global__ void matmul_naive(float* A, float* B, float* C, 
                              int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += A[row * K + k] * B[k * N + col];  // Global memory!
        }
        C[row * N + col] = sum;
    }
}

#define TILE_SIZE 32

// Tiled matrix multiply using shared memory
__global__ void matmul_tiled(float* A, float* B, float* C,
                             int M, int N, int K) {
    // Shared memory tiles
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    // Loop over tiles
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // Cooperative load into shared memory
        if (row < M && t * TILE_SIZE + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && t * TILE_SIZE + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
        
        __syncthreads();  // Wait for tile to load
        
        // Compute on tile (from shared memory - fast!)
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        
        __syncthreads();  // Wait before loading next tile
    }
    
    if (row < M && col < N)
        C[row * N + col] = sum;
}