GPU编程之GUDA（四）——基础知识补充

 

主机函数：在CPU上调用，CPU上执行的函数

全局函数：在CPU上调用，GPU上执行的函数

设备函数：在GPU上调用，GPU上执行的函数——它的线程配置由调用关系中最近的全局函数决定

 

主机函数在声明时可以带有限定符 _host_

全局函数在声明时必须带有限定符 _global_

设备函数在声明时必须带有限定符 _device_

 

在调用全局函数时，除了函数名和形参表外，还有一个用三个小于号“<”和三个大于号“>”包含的部分，这一部分用来指定在并行计算式使用的线程组数和每个线程组白虎呢的线程数。

my_first_kernel<<<nBlocks,nThreads>>>(/*全局函数参数表*/)

一个grid包含完成一次函数调用的所有block，一个block由若干个线程组成

block和grid的尺寸都可以用三元向量来表示，这表明block和grid都是三位数组

 

grid的数组元素由block组成，block的数组元素由线程组成

 

表示线程数量时，可以用CUDA自带的dim3变量来声明三元向量，如果是二元向量，即和第三元置为一是等效的——可缺省，如果只是一维的，可以用int代替

dim3 gridsize(3,2,1);
dim3 blocksize(2,2,2);
my_first_kernel<<<gridsize,blocksize>>>();

dim3 gridsize(3,2);

 

内建向量是从基本体数据类型扩展而成的结构体，访问内建向量的元素可以用索引x,y.z.w

 

=======================================================================================================

运用GUDA运算时，与经典GPGPU方法一样，需要一个输入缓存来准备原始数据和一个输出缓存来接收GPU计算的结果

因此，需要为程序分配两个数组：一个设备上的数组pfDevice用来存放GPU生成的数据，一个在主机上的数组pfHost，用来接收从GU取回的数据

这两个数组应该拥有相同的大小和类型

=======================================================================================================

 

线程索引，是线程在每个block里的索引，由于block的尺寸是三维的，因此线程索引也是一个三元向量：threadIdx

一个grid中的block索引：block.x  block.y  block.z

一个block中的thread索引：thread.x  thread.y  thread.z 

一次函数调用只有一个grid，因此不存在grid索引

 

block尺寸保存在向量blockDim中，grid尺寸保存在gridDim中

线程号：每个线程唯一的标识符

当grid和block是一维的时候，线程号是这样直观定义的

int tid=threadIdx.x+blockDim.x*blockIdx.x;

 

在实际运用中，如果没有某些需要，一般都没有必要使用都为的block，因为一维的block下线程比较容易管理

什么是某些需要？

这样的“某些需要”可以和硬件有关（1、某些存储器为了二维元素而特别优化，2、纹理存储器），也可以是算法的需要，有些算法在使用二维block时较为直观，有助于用户实现算法

=======================================================================================================

内核：人们习惯将全局函数这样在GPU上执行的函数称为内核，简称核。它是函数执行时的基本单位

 

一个grid对应一个GPU或者多个多处理器

一个block中的所有线程在一个多处理器上并发执行

对于同一个block中的线程来说：一、同一个block中的线程可以同步

                                                    二、他们可以共同访问多处理器里的共享存储器

                                                    三：他们不能被拆分到多个多处理器上执行

对处理器将每个县城与一个标量处理器核心对应起来

 

warp：32个线程

如果block中的线程数为32的整数倍，会给调度带来方便。

 

CUDA存储器分为三层：1、私有本地存储器：可被单个线程访问

                                       2、共享存储器：可被block中所有线程访问

                                       3、全局存储器、本地存储器：可被一个应用程序中所有线程访问

除此之外还有程序中所有线程都可以访问的只读存储器：常量存储器和纹理存储器

 

自动变量：在全局函数和设备函数中不是用限定符说明的变量——系统决定他们存放在哪里——————一般放在寄存器中

寄存器：独立分布于贝格标量处理器上，为每个线程提供私有的存储空间

            特点：容量小，访问速度快。

 

共享存储器：位于么一个多处理器内——两种方式：静态分配方式，动态分配方式（动态分配时，存储器尺寸可以是变量）

                    在调用全局函数时不需要在执行配置表中给出共享存储器的尺寸，直接在内核中用限定符_shared_即可

 

对共享存储器访问是以半个warp为单位的

bank冲突：当半warp中的多个线程方位的数组元素位于同一个bank时，他们的读写操作就不能同时进行，也就是说，这几个对同一bank的访问会被串行化

 

无冲突访问模式：1、分在不同的bank中，2、广播

 

设备存储器分为：全局存储器、纹理存储器、常量存储器

设备存储器有两种分配方式：分配为线性存储器、也可以被分配为CUDA数组

 

CUDA数组是为读取纹理而特别优化的，因此一般将纹理存储器分配为CUDA数组

全局存储器最便捷的使用方式是线性存储的方式

常量存储器：容量小，只读

 

设备存储器不是安装在多处理器内的存储器，即片外存储器，因此对设备存储器的直接访问需要耗费大量的时间

缓存、寄存器、共享存储器访问速度快

 

使用全局存储器需要满足接合的要求：不仅要求访问连续的存储单元，还对访问的起始地址有一定的要求——数据的起始地址应为每个线程访问数据大小的16倍的整数倍

一维、二维、四维向量可以结合成较大的数据传输单元，二三维向量的结核性能稍逊，只能作为三个一维标量来处理

 

并发：

数据传输与内核的并发

内核的并发

数据传输的并发

 

流：不同的流在执行时是相互独立的，因此他们可能在时间轴上出现乱序的情况，即不是按照定义的顺序执行，故用户需要再设计时考虑到这一点可能带来的隐患

 

参考文献：

仇德元．《GPGPU编程技术——从GLSL、CDPU到OpenGL》[M]．河北省三河市：机械工业出版社，2012年：323．