CUDA入门

CUDA是什么

CUDA,Compute Unified Device Architecture的简称,是由NVIDIA公司创立的基于他们公司生产的图形处理器GPUs(Graphics Processing Units,可以通俗的理解为显卡)的一个并行计算平台和编程模型。

通过CUDA,GPUs可以很方便地被用来进行通用计算(有点像在CPU中进行的数值计算等等)。在没有CUDA之前,GPUs一般只用来进行图形渲染(如通过OpenGL,DirectX)。

开发人员可以通过调用CUDA的API,来进行并行编程,达到高性能计算目的。NVIDIA公司为了吸引更多的开发人员,对CUDA进行了编程语言扩展,如CUDA C/C++,CUDA Fortran语言。注意CUDA C/C++可以看作一个新的编程语言,因为NVIDIA配置了相应的编译器nvcc,CUDA Fortran一样。更多信息可以参考文献。

主要概念与名称

主机

将CPU及系统的内存(内存条)称为主机。

设备

将GPU及GPU本身的显示内存称为设备。

线程(Thread)

一般通过GPU的一个核进行处理。(可以表示成一维,二维,三维,具体下面再细说)。

线程块(Block)

1. 由多个线程组成(可以表示成一维,二维,三维,具体下面再细说)。

2. 各block是并行执行的,block间无法通信,也没有执行顺序。

3. 注意线程块的数量限制为不超过65535(硬件限制)。

线程格(Grid)

由多个线程块组成(可以表示成一维,二维,三维,具体下面再细说)。


线程束

在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行,线程束中的每个线程都将在不同数据上执行相同的命令。

核函数(Kernel)

1. 在GPU上执行的函数通常称为核函数。

2. 一般通过标识符__global__修饰,调用通过<<<参数1,参数2>>>,用于说明内核函数中的线程数量,以及线程是如何组织的。

3. 以线程格(Grid)的形式组织,每个线程格由若干个线程块(block)组成,而每个线程块又由若干个线程(thread)组成。

4. 是以block为单位执行的。

5. 叧能在主机端代码中调用。

6. 调用时必须声明内核函数的执行参数。

7. 在编程时,必须先为kernel函数中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误,例如越界或报错,甚至导致蓝屏和死机。


dim3结构类型

1. dim3是基亍uint3定义的矢量类型,相当亍由3个unsigned int型组成的结构体。uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z;

2. 可使用亍一维、二维或三维的索引来标识线程,构成一维、二维或三维线程块。

3. dim3结构类型变量用在核函数调用的<<<,>>>中。

4. 相关的几个内置变量

4.1. threadIdx,顾名思义获取线程thread的ID索引;如果线程是一维的那么就取threadIdx.x,二维的还可以多取到一个值threadIdx.y,以此类推到三维threadIdx.z。

4.2. blockIdx,线程块的ID索引;同样有blockIdx.x,blockIdx.y,blockIdx.z。

4.3. blockDim,线程块的维度,同样有blockDim.x,blockDim.y,blockDim.z。

4.4. gridDim,线程格的维度,同样有gridDim.x,gridDim.y,gridDim.z。

5. 对于一维的block,线程的threadID=threadIdx.x。

6. 对于大小为(blockDim.x, blockDim.y)的 二维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。

7. 对于大小为(blockDim.x, blockDim.y, blockDim.z)的 三维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。

8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride。

函数修饰符

1. __global__,表明被修饰的函数在设备上执行,但在主机上调用。

2. __device__,表明被修饰的函数在设备上执行,但只能在其他__device__函数或者__global__函数中调用。

常用的GPU内存函数

cudaMalloc()

1. 函数原型: cudaError_t cudaMalloc (void **devPtr, size_t size)。

2. 函数用处:与C语言中的malloc函数一样,只是此函数在GPU的内存你分配内存。

3. 注意事项:

3.1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数;

3.2. 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作;

3.3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数;

3.4. 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作(即不能进行解引用)。

cudaMemcpy()

1. 函数原型:cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。

2. 函数作用:与c语言中的memcpy函数一样,只是此函数可以在主机内存和GPU内存之间互相拷贝数据。

3. 函数参数:cudaMemcpyKind kind表示数据拷贝方向,如果kind赋值为cudaMemcpyDeviceToHost表示数据从设备内存拷贝到主机内存。

4. 与C中的memcpy()一样,以同步方式执行,即当函数返回时,复制操作就已经完成了,并且在输出缓冲区中包含了复制进去的内容。

5. 相应的有个异步方式执行的函数cudaMemcpyAsync(),这个函数详解请看下面的流一节有关内容。

cudaFree()

1. 函数原型:cudaError_t cudaFree ( void* devPtr )。

2. 函数作用:与c语言中的free()函数一样,只是此函数释放的是cudaMalloc()分配的内存。

下面实例用于解释上面三个函数


GPU内存分类

全局内存

通俗意义上的设备内存。

共享内存

1. 位置:设备内存。

2. 形式:关键字__shared__添加到变量声明中。如__shared__ float cache[10]。

3. 目的:对于GPU上启动的每个线程块,CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存,但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

常量内存

1. 位置:设备内存

2. 形式:关键字__constant__添加到变量声明中。如__constant__ float s[10];。

3. 目的:为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下,用常量内存替换全局内存能有效地减少内存带宽。

4. 特点:常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时,静态地分配空间。

5. 要求:当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol(),而cudaMemcpy()会复制到全局内存。

6. 性能提升的原因:

6.1. 对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。(为什么是15,因为“邻近”指半个线程束,一个线程束包含32个线程的集合。)

6.2. 常量内存的数据将缓存起来,因此对相同地址的连续读操作将不会产生额外的内存通信量。

纹理内存

1. 位置:设备内存

2. 目的:能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计,意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图:


3. 纹理变量(引用)必须声明为文件作用域内的全局变量。

4. 形式:分为一维纹理内存 和 二维纹理内存。

4.1. 一维纹理内存

4.1.1. 用texture<类型>类型声明,如texture texIn。

4.1.2. 通过cudaBindTexture()绑定到纹理内存中。

4.1.3. 通过tex1Dfetch()来读取纹理内存中的数据。

4.1.4. 通过cudaUnbindTexture()取消绑定纹理内存。

4.2. 二维纹理内存

4.2.1. 用texture<类型,数字>类型声明,如texture texIn。

4.2.2. 通过cudaBindTexture2D()绑定到纹理内存中。

4.2.3. 通过tex2D()来读取纹理内存中的数据。

4.2.4. 通过cudaUnbindTexture()取消绑定纹理内存。

固定内存

1. 位置:主机内存。

2. 概念:也称为页锁定内存或者不可分页内存,操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址,因为这块内存将不会破坏或者重新定位。

3. 目的:提高访问速度。由于GPU知道主机内存的物理地址,因此可以通过“直接内存访问DMA(Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。

4. 缺点:使用固定内存,将失去虚拟内存的所有功能;系统将更快的耗尽内存。

5. 建议:对cudaMemcpy()函数调用中的源内存或者目标内存,才使用固定内存,并且在不再需要使用它们时立即释放。

6. 形式:通过cudaHostAlloc()函数来分配;通过cudaFreeHost()释放。

7. 只能以异步方式对固定内存进行复制操作。

原子性

1. 概念:如果操作的执行过程不能分解为更小的部分,我们将满足这种条件限制的操作称为原子操作。

2. 形式:函数调用,如atomicAdd(addr,y)将生成一个原子的操作序列,这个操作序列包括读取地址addr处的值,将y增加到这个值,以及将结果保存回地址addr。

常用线程操作函数

1. 同步方法__syncthreads(),这个函数的调用,将确保线程块中的每个线程都执行完__syscthreads()前面的语句后,才会执行下一条语句。

使用事件来测量性能

1. 用途:为了测量GPU在某个任务上花费的时间。CUDA中的事件本质上是一个GPU时间戳。由于事件是直接在GPU上实现的。因此不适用于对同时包含设备代码和主机代码的混合代码设计。

2. 形式:首先创建一个事件,然后记录事件,再计算两个事件之差,最后销毁事件。如:



1. 扯一扯:并发重点在于一个极短时间段内运行多个不同的任务;并行重点在于同时运行一个任务。

2. 任务并行性:是指并行执行两个或多个不同的任务,而不是在大量数据上执行同一个任务。

3. 概念:CUDA流表示一个GPU操作队列,并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作,如核函数启动,内存复制以及事件的启动和结束等。这些操作的添加到流的顺序也是它们的执行顺序。可以将每个流视为GPU上的一个任务,并且这些任务可以并行执行。

4. 硬件前提:必须是支持设备重叠功能的GPU。支持设备重叠功能,即在执行一个核函数的同时,还能在设备与主机之间执行复制操作。

5. 声明与创建:声明cudaStream_t stream;,创建cudaSteamCreate(&stream);。

6. cudaMemcpyAsync():前面在cudaMemcpy()中提到过,这是一个以异步方式执行的函数。在调用cudaMemcpyAsync()时,只是放置一个请求,表示在流中执行一次内存复制操作,这个流是通过参数stream来指定的。当函数返回时,我们无法确保复制操作是否已经启动,更无法保证它是否已经结束。我们能够得到的保证是,复制操作肯定会当下一个被放入流中的操作之前执行。传递给此函数的主机内存指针必须是通过cudaHostAlloc()分配好的内存。(流中要求固定内存)

7. 流同步:通过cudaStreamSynchronize()来协调。

8. 流销毁:在退出应用程序之前,需要销毁对GPU操作进行排队的流,调用cudaStreamDestroy()。

9. 针对多个流:

9.1. 记得对流进行同步操作。

9.2. 将操作放入流的队列时,应采用宽度优先方式,而非深度优先的方式,换句话说,不是首先添加第0个流的所有操作,再依次添加后面的第1,2,…个流。而是交替进行添加,比如将a的复制操作添加到第0个流中,接着把a的复制操作添加到第1个流中,再继续其他的类似交替添加的行为。

9.3. 要牢牢记住操作放入流中的队列中的顺序影响到CUDA驱动程序调度这些操作和流以及执行的方式。

技巧

1. 当线程块的数量为GPU中处理数量的2倍时,将达到最优性能。

2. 核函数执行的第一个计算就是计算输入数据的偏移。每个线程的起始偏移都是0到线程数量减1之间的某个值。然后,对偏移的增量为已启动线程的总数。

最后编辑于
©著作权归作者所有,转载或内容合作请联系作者
  • 序言:七十年代末,一起剥皮案震惊了整个滨河市,随后出现的几起案子,更是在滨河造成了极大的恐慌,老刑警刘岩,带你破解...
    沈念sama阅读 206,602评论 6 481
  • 序言:滨河连续发生了三起死亡事件,死亡现场离奇诡异,居然都是意外死亡,警方通过查阅死者的电脑和手机,发现死者居然都...
    沈念sama阅读 88,442评论 2 382
  • 文/潘晓璐 我一进店门,熙熙楼的掌柜王于贵愁眉苦脸地迎上来,“玉大人,你说我怎么就摊上这事。” “怎么了?”我有些...
    开封第一讲书人阅读 152,878评论 0 344
  • 文/不坏的土叔 我叫张陵,是天一观的道长。 经常有香客问我,道长,这世上最难降的妖魔是什么? 我笑而不...
    开封第一讲书人阅读 55,306评论 1 279
  • 正文 为了忘掉前任,我火速办了婚礼,结果婚礼上,老公的妹妹穿的比我还像新娘。我一直安慰自己,他们只是感情好,可当我...
    茶点故事阅读 64,330评论 5 373
  • 文/花漫 我一把揭开白布。 她就那样静静地躺着,像睡着了一般。 火红的嫁衣衬着肌肤如雪。 梳的纹丝不乱的头发上,一...
    开封第一讲书人阅读 49,071评论 1 285
  • 那天,我揣着相机与录音,去河边找鬼。 笑死,一个胖子当着我的面吹牛,可吹牛的内容都是我干的。 我是一名探鬼主播,决...
    沈念sama阅读 38,382评论 3 400
  • 文/苍兰香墨 我猛地睁开眼,长吁一口气:“原来是场噩梦啊……” “哼!你这毒妇竟也来了?” 一声冷哼从身侧响起,我...
    开封第一讲书人阅读 37,006评论 0 259
  • 序言:老挝万荣一对情侣失踪,失踪者是张志新(化名)和其女友刘颖,没想到半个月后,有当地人在树林里发现了一具尸体,经...
    沈念sama阅读 43,512评论 1 300
  • 正文 独居荒郊野岭守林人离奇死亡,尸身上长有42处带血的脓包…… 初始之章·张勋 以下内容为张勋视角 年9月15日...
    茶点故事阅读 35,965评论 2 325
  • 正文 我和宋清朗相恋三年,在试婚纱的时候发现自己被绿了。 大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
    茶点故事阅读 38,094评论 1 333
  • 序言:一个原本活蹦乱跳的男人离奇死亡,死状恐怖,灵堂内的尸体忽然破棺而出,到底是诈尸还是另有隐情,我是刑警宁泽,带...
    沈念sama阅读 33,732评论 4 323
  • 正文 年R本政府宣布,位于F岛的核电站,受9级特大地震影响,放射性物质发生泄漏。R本人自食恶果不足惜,却给世界环境...
    茶点故事阅读 39,283评论 3 307
  • 文/蒙蒙 一、第九天 我趴在偏房一处隐蔽的房顶上张望。 院中可真热闹,春花似锦、人声如沸。这庄子的主人今日做“春日...
    开封第一讲书人阅读 30,286评论 0 19
  • 文/苍兰香墨 我抬头看了看天上的太阳。三九已至,却和暖如春,着一层夹袄步出监牢的瞬间,已是汗流浃背。 一阵脚步声响...
    开封第一讲书人阅读 31,512评论 1 262
  • 我被黑心中介骗来泰国打工, 没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留,地道东北人。 一个月前我还...
    沈念sama阅读 45,536评论 2 354
  • 正文 我出身青楼,却偏偏与公主长得像,于是被迫代替她去往敌国和亲。 传闻我的和亲对象是个残疾皇子,可洞房花烛夜当晚...
    茶点故事阅读 42,828评论 2 345

推荐阅读更多精彩内容

  • CUDA从入门到精通(零):写在前面 本文原版链接: 在老板的要求下,本博主从2012年上高性能计算课程开始接触C...
    Pitfalls阅读 3,600评论 1 3
  • 1. CPU vs. GPU 1.1 四种计算机模型 GPU设计的初衷就是为了减轻CPU计算的负载,将一部分图形计...
    王侦阅读 20,787评论 3 20
  • 一直很好奇GPU做矩阵运算是怎么并行加速的,今天看了一些粗浅的东西,并总结整理出来。version:cuda 8 ...
    bidai541阅读 10,293评论 0 3
  • CUDA是一种新的操作GPU计算的硬件和软件架构,它将GPU视作一个数据并行计算设备,而且无需把这些计算映射到图形...
    ai领域阅读 9,037评论 0 8
  • CUDA 简介 CUDA 是由英伟达公司开发的并行计算平台和编程模型的统称。通过这个平台我们可以利用 GPU 并行...
    ClarenceYK阅读 5,120评论 0 7