如何对CUDA代码的内存分配与释放情况进行性能分析？

一、使用NVIDIA官方性能分析工具 #

1. Nsight Systems（系统级追踪） #

• 直接追踪所有CUDA内存API调用（cudaMalloc/cudaFree/cudaMallocManaged等），在timeline视图中可查看每个操作的内存大小、耗时、调用栈，还能自动统计全局内存的总分配/释放量。
• 操作流程：启动Nsight Systems → 选择你的CUDA程序 → 运行分析 → 在"CUDA API"栏筛选内存相关调用即可。

2. Nsight Compute（内核级分析） #

• 针对共享内存：可查看内核中静态分配（__shared__变量）和动态分配（启动时指定的大小）的共享内存使用量，还能检测bank conflict等内存访问问题。
• 编译时加-lineinfo参数，分析时能关联到代码行，定位共享内存的具体使用位置。

二、代码内轻量统计（无需额外工具） #

1. 全局内存使用情况查询 #

封装一个工具函数，在关键节点（如内存操作前后）调用，避免手动重复写打印：

void checkGPUmem(const char* tag) {
    size_t free, total;
    cudaMemGetInfo(&free, &total);
    size_t used = total - free;
    printf("[%s] GPU Mem: Used %zu MB | Free %zu MB | Total %zu MB\n",
           tag, used/(1024*1024), free/(1024*1024), total/(1024*1024));
}

用法示例：

checkGPUmem("Before cudaMalloc");
float* d_data;
cudaMalloc(&d_data, 1024*1024*4); // 4MB
checkGPUmem("After cudaMalloc");
cudaFree(d_data);
checkGPUmem("After cudaFree");

2. 共享内存编译时检查 #

编译CUDA代码时添加nvcc -Xptxas -v参数，编译器会输出每个内核的共享内存使用量，例如：

ptxas info : Used 64 bytes smem, 48 bytes cmem[0]

其中smem就是共享内存的使用大小。

三、自动追踪所有内存操作（CUDA Profiler API） #

通过注册回调函数，自动捕获所有CUDA API调用，无需手动插桩：

#include <cuda_profiler_api.h>
#include <stdio.h>
#include <string.h>

void CUDAAPI memTraceCallback(void* userData, cudaProfilerCallbackDomain domain, 
                              cudaProfilerCallbackId cbid, const char* apiName, void* callbackData) {
    if (domain != CUDA_PROFILER_DOMAIN_API) return;

    if (strcmp(apiName, "cudaMalloc") == 0) {
        auto mallocData = (cudaMallocCallbackData*)callbackData;
        printf("[TRACE] cudaMalloc: Size = %zu bytes | Ptr = %p\n", mallocData->size, mallocData->devPtr);
    } else if (strcmp(apiName, "cudaFree") == 0) {
        auto freeData = (cudaFreeCallbackData*)callbackData;
        printf("[TRACE] cudaFree: Ptr = %p\n", freeData->devPtr);
    } else if (strcmp(apiName, "cudaMallocManaged") == 0) {
        auto managedData = (cudaMallocManagedCallbackData*)callbackData;
        printf("[TRACE] cudaMallocManaged: Size = %zu bytes | Ptr = %p\n", managedData->size, managedData->ptr);
    }
}

int main() {
    // 注册回调
    cudaProfilerRegisterCallback(memTraceCallback, NULL);
    cudaProfilerStart();

    // 你的CUDA内存操作代码
    float* d_a;
    cudaMalloc(&d_a, 2*1024*1024);
    cudaFree(d_a);

    cudaProfilerStop();
    return 0;
}

编译时直接用nvcc -o mem_trace mem_trace.cu即可，运行时会自动打印所有内存分配/释放的细节。

内容的提问来源于stack exchange，提问作者emely_pi