CUDA-Nsight Eclipse Edition

Nsight 是一个开发CUDA程序的IDE和debug工具。

  1. 使用Nsight打开samples程序。
    选择“new”,“CUDA C/C++ Project”,给project命名,Project type 选择”Import CUDA Sample“。接下来从你的机器的samples install location中选择想要打开的project。如果机器上有CUDA-enabled GPU,接下来的设置默认就好。此时可以看到,.cu文件存在于project下的src文件中。这表示,如果自己新建的project也应个先create一个src文件夹,来存放所有源文件。

  2. 使用Nsight创建自己的project。
    如上述,只需在project type选择“Empty Project”。然后在这个project中 新建一个“Source Folder”,取名为src。最后就可以把所有的 .cu, .cuh, .cpp, .h 等源码文件在src中创建。

Nsight 的强大之处在于debug。它可以告诉你你的程序使用了多少SM,多少warp,多少registers,以及每个register中所存放的内容,SM的利用率,硬件基本信息等等。除了debug,Nsight还集成了visual profiler的功能,即可视化程序每个部分的执行时间,以便找到程序可优化之处:
以如下简单code为例:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
#include <stdio.h>
__global__ void kernel(int N){
unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;
unsigned int stride = blockDim.x * gridDim.x;

for (int idd=tid; idd<N; idd+=stride){
printf("hello from thread: %d\n", idd);
}
}

int main(int argc, char** argv){

kernel<<<2, 12>>>(36);
cudaDeviceSynchronize(); //同步Device和Host,即,device 执行完后再执行Host

printf("hello from Host\n");
return 0;
}

即我有36个元素要处理,使用32个线程,并且32个线程分配到1个block中。当启用debug时,可以得到Device端的信息。
从硬件角度看:

可以看到我的GPU编号为0,共有5个SM,使用了2个SM,每个SM都有64分warp,只使用了1个warp。

具体看一个SM中的一个warp。一个warp有32个线程,此处只使用了12个。

从逻辑角度看:启用了两个block,分别在两个SM中。

每个block使用12个线程。

另外Nsight还会给出Host的信息,如下:

以下是GPU中registers中的信息:

以及dissambly信息:

当生成可执行文件后,便可以使用profiler测程序的性能。


CUDA