SZ神庙

从此开始,遁入幻想

2017年5月12日

秘术记闻>

Cuda学习笔记

已经四个月没有写博客了。这段时间真的是各种忙,忙着毕设,忙着找实习,时间怎么样都不够用,闲下来的时候真的是宁可躺床上发发呆,也懒得动笔写点什么。原本自己开这个博客的初心就是随便记录点学习上、生活上的事情,可是半年过去了,脑子里总还是空空的时候居多,所以这博客也是寂寥的像山间小居一样。经营blog真的也是需要很多精力的事情呢。

然而怠惰总不是件好事,所以还是动笔写一点什么吧。正好这段时间在用Cuda做并行计算,也踩了不少坑,就拿出来写一写吧。

一.Cuda环境配置

首先,你得有一台配备了支持Cuda的显卡的机器。然后,去Nvidia官网根据你的系统下载相应的开发包。我是Ubuntu14.04的系统,按照官网上的说明很容易就能顺利安装好。然后,新建工程,在CMakeLists.txt下加入下面的代码:

Cuda源文件以*.cu为后缀名。生成可执行文件的时候,只需要在add_executable或add_library的前面加上前缀cuda_,也就是像这样:

或者

最后记得链接上Cuda的库:

二. 一些心得体会

Cuda写起来真的是很方便,不用像写shader那样麻烦,你完全可以像写CPU程序那样写GPU程序。我这篇博文不想写成教程那样,所以一些基本知识就请阅读《CUDA by Example.An Introduction to General-Purpose GPU Programming》这本书吧,很赞,非常适合入门。我这里就记录一下一些书上没有提到的,我的一些心得。

2.1 数据传递

我们知道,GPU和CPU使用的是不同的存储空间。不过,Cuda6以后,Nvidia新增了Unified Memory, 也就是可以被CPU和GPU共享的、对双方都可见的存储空间。官网上给出来的例子是这个样子的:

看上去很方便,不过仍然要记得手动free,所以我就仿照shared_ptr,利用RAII自己封装了一个CudaPointer,这玩意内部带了一个引用计数,不用自己手动free了。在项目里我用上了,感觉还挺方便。比较遗憾的是没法直接在kernel函数里面使用,得转成裸指针。这也没办法,你不能让GPU去执行cudaMallocManaged()和cudaFree()呀。类似的我还自己写了个CudaVector。本来Cuda有个thrust库,这玩意把STL的容器底层用GPU重新实现了一遍,然而这玩意只能在CPU上用,GPU端用不了,我觉得很不爽,干脆自己写了个,在项目里用着还挺爽的。代码这里就不贴了。

2.2 线程同步

多线程程序里同步总是很关键的一环。在Cuda里,原生API只提供了thread级的同步,不提供block级的同步,所以你只能对单个block内部进行同步操作。

同步主要利用的是__syncthreads()这个函数。它的功能是设置一个barrier,当一个线程执行到这个函数以后,它会等待其他所用线程也执行到这个函数为止,然后才继续执行下面的代码。然而,使用这个函数同步的时候有一点非常重要,请务必保证所有thread执行__syncthreads()的次数是相等的。如果有的thread执行__syncthreads()的次数和其他thread不等,会引发未定义行为。

举个例子,在我的项目里,使用了大概类似这么一段代码:

这段代码的本意是什么呢,我是打算用这段代码来并行计算一些值,然后选出其中最小的那一个。我一共开了256个线程,也就是THREADS_PER_BLOCK是256,然而要计算的值不止这么多,那么我就想一部分一部分的来算,先算256个,求一个最小值,再算接下来的256个,再求一个最小值。。。看上去很美,但实际运行以后完全就不对!!!问题出在哪里呢?

问题就出在对所用线程来说,最外层那个大循环的执行次数可能是不一样的。比如如果我有400个值要计算,那么有144个线程执行了两次循环,而有112个线程只制定了一次循环。因此,各个线程对__syncthreads()的调用次数是不等的,也就发生了未定义行为,__syncthreads()失去了同步的作用了。这段代码正确的写法是这样的:

__syncthreads()只能对同一个block中的不同threads进行同步,如果我们想要实现不同block间的同步要怎么做呢?其实也是有办法的,Cuda支持原子操作,可以据此自己实现一个spinlock,这样就可以支持不同block间的同步,而且比起__syncthreads()这种barrier形式的同步能有更多的适用场景。具体做法就留到下一篇文章再写吧,今天就到这里了。晚安,世界。