使用ncu (NsightComputeCli)获取nvprof默认行为

Question

默认的nvprof输出是很好的，但是nvprof已经被废弃了，而倾向于ncu。如何使ncu给出一个更像nvprof的输出

最小工作示例

我有两个range函数，其中一个是以非常不理想的方式调用的(只使用一个线程)。它比其他range函数花费的时间要长得多。

profile.cu

#include <stdio.h>

//! makes sure both range functions executed correctly
bool check_range(int N, float *x_d) {
    float *x_h;
    cudaMallocHost(&x_h,N*sizeof(float));
    cudaMemcpy(x_h, x_d, N*sizeof(float), cudaMemcpyDeviceToHost);
    bool success=true;
    for( int i=0; i < N; i++)
        if( x_h[i] != i ) {
            printf("\33[31mERROR: x[%d]=%g\33[0m\n",i,x_h[i]);
            success=false;
            break;
        }
    cudaFreeHost(x_h);
    return success;
}

//! called with many threads
__global__ void range_fast(int N, float *x) {
    for( int i=threadIdx.x; i < N; i+=blockDim.x)
        x[i]=i;
}

//! only gets called with 1 thread. This is the bottleneck I want to detect
__global__ void range_slow(int N, float *x) {
    for( int i=threadIdx.x; i < N; i+=blockDim.x)
        x[i]=i;
}

int main(int argc, char *argv[]) {
    int N=(1<<20)*10;
    float *x_fast, *x_slow;
    cudaMalloc(&x_fast,N*sizeof(float));
    cudaMalloc(&x_slow,N*sizeof(float));
    range_fast<<<1,512>>>(N,x_fast);
    range_slow<<<1,1>>>(N,x_slow);
    check_range(N,x_fast);
    check_range(N,x_slow);
    cudaFree(x_fast);
    cudaFree(x_slow);
    return 0;
};

编译

nvcc profile.cu -o profile.exe

nvprof profiling

nvprof ./profile.exe

nvprof 输出

            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   99.17%  1.20266s         1  1.20266s  1.20266s  1.20266s  range_slow(int, float*)
                    0.53%  6.3921ms         2  3.1961ms  3.1860ms  3.2061ms  [CUDA memcpy DtoH]
                    0.31%  3.7273ms         1  3.7273ms  3.7273ms  3.7273ms  range_fast(int, float*)
      API calls:   88.79%  1.20524s         2  602.62ms  3.2087ms  1.20203s  cudaMemcpy
                    9.31%  126.39ms         2  63.196ms  100.62us  126.29ms  cudaMalloc
                    1.11%  15.121ms         2  7.5607ms  7.5460ms  7.5754ms  cudaHostAlloc
                    0.64%  8.6687ms         2  4.3344ms  4.2029ms  4.4658ms  cudaFreeHost
                    0.09%  1.2195ms         2  609.73us  103.80us  1.1157ms  cudaFree

这使我清楚地了解了哪些函数占用了运行时的大部分时间，而range_slow是瓶颈。

ncu profiling

ncu ./profile.exe

ncu 输出

ncu输出有更多的细节，其中大部分我并不真正关心。它也没有被很好的总结。

Answer 1

在“新”分析工具中，nvprof的功能已被分解为2个单独的工具。Nsight Compute工具主要关注内核的活动(即设备代码)分析，虽然它可以报告内核的持续时间，但是它对API调用活动和内存复制活动的兴趣不大。

具有此功能的工具是Nsight Systems。

尝试：

nsys profile --stats=true ./profile.exe

除其他外，您将得到GPU活动的pareto列表(分为单独的内核活动和内存复制活动的pareto列表)和API调用的pareto列表。