浅谈GPU编程.

1、CUDA略谈 引言 逻辑线程拓扑 CUDA存储器类型 设备端存储器 主机端存储器 参考书CUDA 统一计算设备架构 它是一种将GPU作为数据并行计算设备的软硬件体系。 开发GPU通用计算程序，需掌握并行算法和GPU架构方面的知识。 支持CUDA的GPU可有效利用原用于图形渲染的计算资源进行通用计算。 此处就基于GPU的CUDA编程谈一些认识。线程逻辑拓扑结构线程逻辑拓扑结构 线程拓扑分两个层次。 顶层是2维网格平面，以块为单位。 块为3维立方结构，以线程为单位，故称线程块。grid:44block:444threadblockCUDA程序结构内核函数：即设备端代码。形如：_global_ ke

2、rnel(形参)主机端代码void main()dim3 block(4,4,4);dim3 grid(4,4,1);kernel(实参);存储器层次结构存储器类型设备端存储器结构 寄存器 共享存储器 局部存储器 常数存储器 全局存储器 纹理存储器板载显存GPU片上寄存器 GPU片上高速存储器 每个寄存器32bit。 每个SM有大量寄存器，但由块内线程共享，故平均到每个线程的寄存器就很有限了。 在内核函数中声明的少量变量是寄存器变量，每个线程都维护各自的寄存器变量，它们是线程私有的。_global_ kernel()int bx=blockIdx.x;int by=blockIdx.y;int

3、 tx=threadIdx.x;int ty=threadIdx.y;int tz=threadIdx.z;局部存储器 板载显存 如果每个线程使用了过多的寄存器，或声明了大型结构体或数组，或者编译器无法确定数组的大小，线程的私有变量就有可能会被分配到局部存储器中。 例见P46共享存储器（SM） 线程块内共享 片上高速存储器 静态分配 动态分配_global_ kernel(形参) _shared_ int sm_static16;extern _shared_ int sm_dynamic;void main()int sm_dynamic_size=32;kernel(实参);共享存储器（S

4、M）SM的组织方式 在1.x计算能力的设备中，sm共16KB，被划分为16个bank，每个bank的宽度为32位。32bitbank001632bitbankii16+i32bitbank15153132bitbank1117共享存储器（SM）SM访问的bank冲突： 每个warp为32个线程，一个warp对sm的访问被分成两个half-warp的访问，只有同一个half-warp内的线程才可能发生bank冲突。 前half-warp的线程与后half-warp的线程之间则不会发生bank冲突。共享存储器（SM）bankhalf-warpbankhalf-warp顺序访问随机访问无bank冲突

5、bankhalf-warp3间隔访问共享存储器（SM）bankhalf-warpbankhalf-warp2间隔访问8间隔访问2路bank冲突bankhalf-warp广播机制8路bank冲突无bank冲突常数存储器 只读存储器，数据位于显存，但拥有缓存加速。 空间较小，只有64KB。 每个SM拥有8KB的常数存储器缓存。 在函数外定义，作用范围为文件域。主机端和设备端函数均可见。 使用方法： 方法1：定义时初始化，内核函数中直接使用。 方法2：先定义，后在主机端用函数赋值。常数存储器_constant_ char p_HelloCUDA11;/定义_constant_ int t_Hello

6、CUDA11=0,1,2,3,4,5,6,7,8,9,10;/定义并初始化_constant_ int num=11; /定义并初始化_global_ static void HelloCUDA(char* result)int i=0;for (i=0;inum;i+) resulti=p_HelloCUDAi+t_HelloCUDAi;int main()char helloCUDA=“Hello CUDA!”;cudaMemcpyToSymbol(p_HelloCUDA,helloCUDA,sizeof(char)*11);直接使用函数赋值全局存储器 位于板载显存，占据显存的绝大部分，没

7、有缓存。 也称线性内存。 可定义两种数据结构： 线性存储器：用于存放主机端传过来的数据或存放将要回传给主机的数据。 CUDA数组：用于纹理绑定。 为有效利用带宽，必须遵循合并访问要求，并避免分区冲突。全局存储器分区(partition)冲突：在中高端GPU中一般有多个存储器控制器。每个存储器控制器对应的地址空间称为一个分区，连续的256Byte数据存储于同一分区中，相邻的另一组256Byte数据则存储在相邻的分区中。GTX280中有8个分区，一个512个元素的浮点数组布局如下，在访问数据时应均匀分布在不同的分区中。256bytepartition006344851164127256byte25

8、6byte256bytepartition1partitionipartition7全局存储器合并访问： 按段长对齐可满足合并访问的要求 结构体数组的对齐定义P157 一维线性数组至少按256Byte对齐 二、三维线性数组采用填充（P49）进行对齐 满足对齐要求的数组分配函数及赋值函数见下表。分配函数赋值函数一维cudaMalloc()cudaMemcpy()二维cudaMallocPitch()cudaMemcpy2D()三维cudaMalloc3D()cudaMemcpy3D()纹理存储器GPU芯片全局存储器用于纹理的两种数据形式：普通线性存储器CUDA数组纹理存储器HOST显卡纹理参照系

9、纹理拾取绑定纹理存储器有关术语及描述： 只读存储器 容量比常数存储器大 纹理拾取：在内核函数中访问纹理存储器的操作。 纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参照系约定二者的映射方式。 将数据与纹理绑定：将显存中的数据与纹理参照系关联的操作。纹理存储器有关术语及描述： 显存中可绑定到纹理的数据形式： 普通线性存储器（非对齐） 只能与一维或者二维纹理绑定，采用整型纹理拾取坐标，坐标与数据在存储器中的位置相同。 CUDA数组 可以与一维、二维或者三维纹理绑定，纹理拾取坐标为归一化或者非归一化的浮点型。 像元：绑定到纹理的线性存储器或CUDA数组中的元素。纹理存储器有关术语及描述：

10、像元的数据类型：不支持三元组unsignedsignedfloat一维uchar1、ushort1、uint1、ulong1char1、short1、int1、long1float1二维uchar2、ushort2、uint2、ulong2char2、short2、int2、long2float2四维uchar4、ushort4、uint4、ulong4char4、short4、int4、long4float4纹理存储器有关术语及描述： 纹理存储器有缓存机制 纹理缓存一次预取拾取坐标对应位置附近的几个像元。 绑定到纹理的数据修改后，应重启内核函数且纹理缓存刷新后，才能拾取到被修改的数据。纹理存

11、储器 纹理存储器的特殊功能： 浮点型纹理拾取坐标 归一化：0.0，1.0) 非归一化：0.0，N)，N为该维度上的像元数。 寻址模式：输入坐标超范围时的处理方式。 钳位模式(clamp)：按上下限坐标拾取像元。 循环模式(wrap)：也称回绕，只用于归一化坐标。 类型转换：对8位或16位整型像元数据，其拾取的返回值可转换为归一化浮点型。 滤波：对CUDA数组绑定的纹理，拾取返回值为浮点型，则可对返回值进行滤波。 最近点取样：适用于查找表。 线性滤波：适用于图像处理。纹理存储器使用步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组元素的数据类型 用结构体cudaC

12、hannelFormatDesc来描述 struct cudaChannelFormatDescint x,y,z,w;/多元组数据中每个分量二进制位数enum cudaChannelFormatKind f; f取值cudaChannelFormatKindSigned，有符号整型cudaChannelFormatKindUnsigned，无符号整型cudaChannelFormatKindFloat，浮点型纹理存储器使用步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组元素的数据类型 如uchar2,则x、y、z、w分别是8，8，0，0，而f取值为cudaC

13、hannelFormatKindUnsigned 又如float4,则x,y,z,w分别是32,32,32,32;而f取值为cudaChannelFormatKindFloat纹理存储器使用步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组的维度 CUDA数组空间分配函数： cudaMalloc3DArray()：分配一维、二维或三维数组 cudaMallocArray()：一般用于二维数组 CUDA数组空间释放函数：cudaFreeArray() CUDA数组与其他CUDA数组或线性存储器的数据传输： cudaMemcpyToArray()或cudaMemcp

14、y3D()纹理存储器使用步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组尺寸 用结构体cudaExtent描述数组三个维度的大小cudaExtent extent=make_cudaextent(1,8192,0,0)cudaExtent extent=make_cudaextent(1,65535,1,32768,0)cudaExtent extent=make_cudaextent(1,2048,1,2048,1,2048)纹理存储器使用步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 例，声明一个数据类型为uchar2，643216的

15、CUDA 3D数组。cudaArray* cuArray=0;cudaExtent extent=make_cudaExtent(64,32,16);cudaChannelFormatDesc desc=cudaCreateChannelDesc(8,8,0,0, cudaChannelFormatKindUnsigned);cudaMalloc3DArray(&cuArray,&desc,extent);纹理存储器使用步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组声明完毕后，还需赋值。 可用以下函数完成CUDA数组赋值。 CUDA数组之间赋值cudaMem

16、cpyToArray() CUDA数组与线性存储器之间赋值cudaMemcpy3D()纹理存储器使用步骤2：声明纹理参照系 纹理参照系属性 编译时属性：编译前显式声明，编译时确定，一旦确定不能修改。 运行时属性：运行时设定，只适用于与CUDA数组绑定的纹理参照系。 在所有函数体外声明文件域的texture型变量，最好放在头文件中。形如下示： texture texRef;纹理存储器使用步骤2：声明纹理参照系texture texRef;编译时属性: Type：纹理拾取返回的数据类型，即CUDA像元 Dim：确定纹理参照系的维度，默认为1，可取值分别为1，2，3 ReadMode：确定返回值是否

17、进行类型转换 cudaReadModeNormalizedFloat：转换为浮点型 cudaReadModeElementType：不转换，此为默认值纹理存储器使用步骤2：声明纹理参照系运行时属性:通过结构体textureReference描述struct textureReferenceint normalized;/坐标是否归一化，非零值表示归一化enum cudaTextureFilterMode filterMode;/滤波模式enum cudaTextureAddressMode addressMode3;struct cudaChannelFormatDesc channelDes

18、c;/纹理拾取返回的数据类型，要与CUDA数组声明时的类型一致，前已述及。滤波模式有以下两种取值：cudaFilterModePoint表示最近点取样cudaFilterModeLinear表示线性滤波寻址模式，它是一个大小为3的数组，对应3个维度，可取以下两种值：cudaAddressModeClamp，表示钳位模式，cudaAddressModeWrap，表示循环模式。纹理存储器使用步骤3：纹理绑定 将纹理与数组按纹理参照系绑定，实现数组到纹理的映射。 绑定函数： cudaBindTexture()/用于纹理与线性存储器绑定 cudaBindTextureToArray ()/用于纹理与CUDA数组绑定 解除绑定：cudaUnbindTexture()纹理存储器使用步骤4：纹理拾取 函数见下表：绑定到纹理的数组拾取函数纹理坐标线性存储器tex1Dfetch()整型CUDA数组一维tex1D()浮点型二维