cuda 学习笔记

uodate:有毒吧。kernel中出问题原来是不会报错的。。。。

请教了组里的hust学长orz..、

学到了cuda-memcheck命令和cudaGetLastError来查看问题。。可以参考What is the canonical way to check for errors using the CUDA runtime API?

 

先放一波资料。

 

 

cuda 提出的目的是能够让程序员透明地使用GPU来高效地进行并行运算。

kernel和c语言中的函数相似,函数名字前通常用global来标识。

 

下面考虑一个2个大小的1M的数组相加的例子。

总的思路是通过并行,来观察到计算速度的加快。

如果不考虑并行,2个数组相加的代码,如下:

 

如果用cuda的方式来搞,代码如下:

除了代码注释,还有几个地方要说明:

  • 函数名字前面的global 是 cuda kernel  标识符
  • cuda kernel的调用方式是 <<<,>>>  更具体地说,是add<<<numBlocks,blockSize>>>(N,x,y)
  • .cu是 cuda C++ 文件的后缀,类似.cpp
  • nvcc是cuda C++ 的编译器,其将source code分成host codedevice code两部分。前者通过c++编译器编译,后者通过nvidia编译器编译。

 

关于devide code 和 host code,参考下图。

 

现在我们单线程地跑了一个cuda kernel, 接下来是如何使它并行.关键在于<<<1,1>>>这部分。

这行代码告诉了cuda runtime 有多少个并行的线程要被执行。
这里有2个参数,不过我们可以先改变第二个,也就是一个线程block中线程的个数。
cuda GPu的kernel 使用的blocks中线程的个数应该是32的倍数(后面会解释32代表什么),所以256看起来很合理。

 

 

不过如果只是修改了<<<1,1>>> 到  <<<1,256>>>
那实际上是对于每个线程都算了整个array的相加,而没有将整个计算任务分给多个并行的线程。
为了解决这个问题,我们需要修改kernel的代码。
cuda C++ 提供了关键字,允许kernel得到正要执行的thread是哪一个
threadIdx.x  表示当前运行的thread是block中的哪一个
blockDim.x 表示block中的线程个数

关于threadIdx.x等 下标问题,参考下图。

 

我们需要观察到使用cuda的方法之后时间的变化。

可以使用nvprof命令

 

 

可以看出时间的变化,从167.48ms到3.5144ms,再到1.8084ms

 

 

我们注意到,对线程的管理实际上是三维的 :grid,block,thread.

为什么要这样设计?

一个这样做的目的是,在一个block中,thread可以通过share_memory 来共享数据。

通过在声明的变量前面添加 shared  来表示,这个变量是声明在share memory部分了。

share memory类似与缓存,容量小,但是速度快。不过这个cache是可以编程控制的。

在一个block中共享的data,对于其他block是不可见的。

我们不妨考虑一个例子,有2个数组a,b

进行如下运算:

 

为了加快运行速度,我们还是考虑多线程的办法。

让每一个线程处理一个输出。

然而我们发现,in中除了边界元素,每一个元素都被读了7次。

这显然是没有必要的。

问题的关键在于,不同的线程之间,不知道某个元素已经被读入了。

更进一步,不同线程之间可以共享数据吗?

答案是可以的。也就是上面提到的shared_memory

然而这样就可以了吗。。

由于线程的访问顺序是不固定的(?

会发生如下的问题:

 

解决办法很无脑。。。因为cuda并没有想象中得那么底层。。。

就是使用__syncthreads(); 来同步一个block中的所有进程。。。

完整代码如下:

 

 

需要特别强调的是,cuda代码的debug问题,很多错误,不用特定的工具查看是不会显示的。
以及,虽然cuda代码是在c++上添加了一些东西,但是device code部分用的是nvidia的编译器。
所以c/cpp中,对于访问非法内存不敏感的特点,在cuda代码中不存在(我猜是因为编译器…)
在访问之前,一定要check访问地址合法性。。。。

在访问之前,一定要check访问地址合法性。。。。

在访问之前,一定要check访问地址合法性。。。。

 

 

 

作者: CrazyKK

ex-ACMer@hust,stackoverflow-engineer@sensetime

说点什么

您将是第一位评论人!

提醒
wpDiscuz