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

Question

默认nvprof输出很棒，但nvprof已被弃用，转而使用ncu. 我怎样才能ncu给我一个看起来更像的输出nvprof？

最小工作示例

我有 2 个range函数，其中一个函数以一种非常不理想的方式调用（仅使用 1 个线程）。它比其他函数需要更长的时间range。

简介.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;
};

汇编

#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;
};

nvprof剖析

nvcc profile.cu -o 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剖析

nvprof ./profile.exe

ncu输出

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

Answer 1

在“新”分析工具中，的功能nvprof已分为2 个独立的工具。Nsight Compute工具主要关注内核（即设备代码）分析的活动，虽然它可以报告内核持续时间，但当然，它对API调用活动和内存复制活动等事情不太感兴趣。

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

尝试：

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

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

在较新版本的 nsight 系统中，添加了一项新功能。某些nvprof命令选项（不是全部）可以通过 发出nsys。尝试：

nsys nvprof --print-gpu-trace ./profile.exe

nvprof --print-gpu-trace例如，这与该工具非常相似。