如何让 malloc 出现在 nvprof 的统计分析器中？

Posted 2023-03-23

【中文标题】如何让 malloc 出现在 nvprof 的统计分析器中？

【英文标题】：How to get malloc to show up in nvprof's statistical profiler?

【发布时间】：2019-11-01 15:54:15

【问题描述】：

有没有办法让 CUDA 的 nvprof 在其统计分析器中包含像 malloc 这样的函数调用？

我一直在努力提高我的应用程序的性能。当然，我一直在使用 nvprof 作为工具。

最近，为了减少我的应用程序的 GPU 内存占用，我编写了代码，使其运行时间增加了一倍。然而，导致速度变慢的新代码在分析器中只出现了少量（指令采样表明大约 10% 的时间都花在了新代码中，但天真的想法表明 50 % 的时间应该花在新代码上）。也许新代码导致更多的缓存抖动，也许将实现放在头文件中，这样它可能会被内联混淆探查器等。但是，没有充分的理由，我怀疑新代码对malloc的调用。

确实，在我减少了malloc 的调用次数后，我的性能提高了，几乎回到了合并新代码之前的水平。

这让我想到了一个类似的问题，为什么malloc 的调用没有出现在统计分析器中？malloc 调用是某种 GPU 系统调用，可以不被观察？

下面，我提供了一个示例程序及其展示此特定问题的输出。

#include <iostream>
#include <numeric>
#include <thread>
#include <stdlib.h>
#include <stdio.h>

static void CheckCudaErrorAux (const char *, unsigned, const char *, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

__global__ void countup()

    long sum = 0;
    for (long i = 0; i < (1 << 23); ++i) 
        sum += i;
    
    printf("sum is %li\n", sum);


__global__ void malloc_a_lot() 
    long sum = 0;
    for (int i = 0; i < (1 << 17) * 3; ++i) 
        int * v = (int *) malloc(sizeof(int));
        sum += (long) v;
        free(v);
    
    printf("sum is %li\n", sum);


__global__ void both() 
    long sum = 0;
    for (long i = 0; i < (1 << 23); ++i) 
        sum += i;
    
    printf("sum is %li\n", sum);

    sum = 0;
    for (int i = 0; i < (1 << 17) * 3; ++i) 
        int * v = (int *) malloc(sizeof(int));
        sum += (long) v;
        free(v);
    
    printf("sum is %li\n", sum);



int main(void)


    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    std::chrono::time_point<std::chrono::system_clock> t1 = std::chrono::system_clock::now();
    countup<<<8,1>>>();
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    std::chrono::time_point<std::chrono::system_clock> t2 = std::chrono::system_clock::now();
    malloc_a_lot<<<8,1>>>();
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    std::chrono::time_point<std::chrono::system_clock> t3 = std::chrono::system_clock::now();
    both<<<8,1>>>();
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    std::chrono::time_point<std::chrono::system_clock> t4 = std::chrono::system_clock::now();

    std::chrono::duration<double> duration_1_to_2 = t2 - t1;
    std::chrono::duration<double> duration_2_to_3 = t3 - t2;
    std::chrono::duration<double> duration_3_to_4 = t4 - t3;
    printf("timer for countup() took %.3lf\n", duration_1_to_2.count());
    printf("timer for malloc_a_lot() took %.3lf\n", duration_2_to_3.count());
    printf("timer for both() took %.3lf\n", duration_3_to_4.count());

    return 0;


static void CheckCudaErrorAux (const char *file, unsigned line, const char *statement, cudaError_t err)

    if (err == cudaSuccess)
        return;
    std::cerr << statement<<" returned " << cudaGetErrorString(err) << "("<<err<< ") at "<<file<<":"<<line << std::endl;
    exit (1);

结果的省略版本是：

sum is 35184367894528...
sum is -319453208467532096...
sum is 35184367894528...
sum is -319453208467332416...
timer for countup() took 4.034
timer for malloc_a_lot() took 4.306
timer for both() took 8.343

分析结果如下图所示。将鼠标悬停在浅蓝色条上时显示的数字与条的大小一致。具体来说，第 41 行有 16,515,077 个样本与之关联，但第 47 行只有 633,996 个样本。

顺便说一句，上面的程序是用调试信息编译的，可能没有优化——在 Nsight Eclipse 中编译的默认“调试”模式。如果我在“发布”模式下编译，会调用优化，countup() 调用的持续时间非常接近 0 秒。

【问题讨论】：

有一个想法，也许我需要在程序结束时加入对cudaProfilerStop() 的调用。所以，我试过了，但变化不大。

我目前正在运行带有标志 --track-memory-allocations on 的 nvprof 会话。将报告调查结果。

使用nvprof 和track-memory-allocations 似乎没有这样做。对于任何关注的人，这里是我发出的确切命令（然后将生成的两个文件导入 Nsight Eclispse/nvvp）。 nvprof -f -o ~/cuda-workspace/malloc/timeline.out ~/cuda-workspace/malloc/Debug/malloc; nvprof -f --track-memory-allocations on --kernels ::both:1 --analysis-metrics -o ~/cuda-workspace/malloc/analysis-metrics.out ~/cuda-workspace/malloc/Debug/malloc;

【参考方案1】：

当前的 NVIDIA GPU PC 采样器仅收集当前的 warp 程序计数器（不是调用堆栈）。 PC 采样器将正确采集 malloc 内部的样本；但是，该工具不显示内部系统调用的 SASS 或高级源代码。

该工具没有 UI 来显示系统调用模块内样本的聚合计数。 该工具不知道 malloc、free 或其他系统调用的 PC 范围，无法将样本正确归因于名为 syscall 的用户。

如果 (1) 或 (2) 是固定的，则数据将显示在单独的行中，仅标记为“syscall”或“malloc”。硬件不收集调用堆栈，因此无法将样本归因于 L48。

【讨论】：

我猜你的意思是写“硬件不收集调用堆栈...”？