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

Question

有没有办法让 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行关联了16515077个样本，而第47行只有633996个样本。

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

Answer 1

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

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

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

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

How to get malloc to show up in nvprof's statistical profiler?

profiling

cuda

nvprof