0

有没有办法让 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 个样本。

<code>both</code> 内核配置文件 - PC 采样

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

4

1 回答 1

2

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

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

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

于 2019-06-25T20:29:37.930 回答