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

Question

默认的 nvprof 输出很好，但是 nvprof 已被弃用，取而代之的是 ncu。我怎样才能让 ncu 给我一个看起来更像 nvprof 的输出？

最小工作示例

我有 2 个 range 函数，其中一个以非常不理想的方式调用（仅使用 1 个线程）。它比其他 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("[31mERROR: x[%d]=%g[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 分析

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 分析

ncu ./profile.exe

ncu输出

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

Answer 1

nvprof 的功能已在“新”分析工具中分解为 2 separate tools。 Nsight Compute 工具主要专注于 activity 内核（即设备代码）分析，虽然它可以报告内核持续时间，但它对 API 调用 [=21] 之类的东西不太感兴趣=] 和内存复制 activity.

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

尝试：

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

除其他外，您将获得 GPU 活动的帕累托列表（分为内核活动和内存复制活动的单独帕累托列表）和 API 次调用的帕累托列表。

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

Get nvprof default behavior with ncu (NsightComputeCli)

cuda