1、CUDA略谈,引言 逻辑线程拓扑 CUDA存储器类型 设备端存储器 主机端存储器 参考书,CUDA,统一计算设备架构 它是一种将GPU作为数据并行计算设备的软硬件体系。 开发GPU通用计算程序，需掌握并行算法和GPU架构方面的知识。 支持CUDA的GPU可有效利用原用于图形渲染的计算资源进行通用计算。 此处就基于GPU的CUDA编程谈一些认识。,线程逻辑拓扑结构,线程逻辑拓扑结构,线程拓扑分两个层次。 顶层是2维网格平面，以块为单位。 块为3维立方结构，以线程为单位，故称线程块。,grid:44,block:444,thread,block,CUDA程序结构,内核函数：即设备端代码。 形如：
2、_global_ kernel(形参) ,主机端代码 void main() dim3 block(4,4,4); dim3 grid(4,4,1); kernel(实参); ,存储器层次结构,存储器类型,设备端存储器结构,寄存器 共享存储器 局部存储器 常数存储器 全局存储器 纹理存储器,板载显存,GPU片上,寄存器,GPU片上高速存储器 每个寄存器32bit。 每个SM有大量寄存器，但由块内线程共享，故平均到每个线程的寄存器就很有限了。 在内核函数中声明的少量变量是寄存器变量，每个线程都维护各自的寄存器变量，它们是线程私有的。,_global_ kernel() int bx=blockI
3、dx.x; int by=blockIdx.y;int 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
4、sm_dynamic_size=32; kernel(实参); ,共享存储器（SM）,SM的组织方式 在1.x计算能力的设备中，sm共16KB，被划分为16个bank，每个bank的宽度为32位。,32bit,bank0,0,16,32bit,banki,i,16+i,32bit,bank15,15,31,32bit,bank1,1,17,共享存储器（SM）,SM访问的bank冲突： 每个warp为32个线程，一个warp对sm的访问被分成两个half-warp的访问，只有同一个half-warp内的线程才可能发生bank冲突。 前half-warp的线程与后half-warp的线程之间则不会
5、发生bank冲突。,共享存储器（SM）,bank,half-warp,bank,half-warp,顺序访问,随机访问,无bank冲突,bank,half-warp,3间隔访问,共享存储器（SM）,bank,half-warp,bank,half-warp,2间隔访问,8间隔访问,2路bank冲突,bank,half-warp,广播机制,8路bank冲突,无bank冲突,常数存储器,只读存储器，数据位于显存，但拥有缓存加速。 空间较小，只有64KB。 每个SM拥有8KB的常数存储器缓存。 在函数外定义，作用范围为文件域。主机端和设备端函数均可见。 使用方法： 方法1：定义时初始化，内核函数中直
6、接使用。 方法2：先定义，后在主机端用函数赋值。,常数存储器,_constant_ char p_HelloCUDA11;/定义 _constant_ int t_HelloCUDA11=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
7、!”; cudaMemcpyToSymbol(p_HelloCUDA,helloCUDA,sizeof(char)*11); ,直接使用,函数赋值,全局存储器,位于板载显存，占据显存的绝大部分，没有缓存。 也称线性内存。 可定义两种数据结构： 线性存储器：用于存放主机端传过来的数据或存放将要回传给主机的数据。 CUDA数组：用于纹理绑定。 为有效利用带宽，必须遵循合并访问要求，并避免分区冲突。,全局存储器,分区(partition)冲突： 在中高端GPU中一般有多个存储器控制器。 每个存储器控制器对应的地址空间称为一个分区，连续的256Byte数据存储于同一分区中，相邻的另一组256Byte数
8、据则存储在相邻的分区中。 GTX280中有8个分区，一个512个元素的浮点数组布局如下，在访问数据时应均匀分布在不同的分区中。,256byte,partition0,0,63,448,511,64,127,256byte,256byte,256byte,partition1,partitioni,partition7,全局存储器,合并访问： 按段长对齐可满足合并访问的要求 结构体数组的对齐定义P157 一维线性数组至少按256Byte对齐 二、三维线性数组采用填充（P49）进行对齐 满足对齐要求的数组分配函数及赋值函数见下表。,纹理存储器,GPU芯片,全局存储器用于纹理的 两种数据形式：普通线
9、性存储器CUDA数组,纹理存储器,HOST,显卡,纹理参照系,纹理拾取,绑定,纹理存储器,有关术语及描述： 只读存储器 容量比常数存储器大 纹理拾取：在内核函数中访问纹理存储器的操作。 纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参照系约定二者的映射方式。 将数据与纹理绑定：将显存中的数据与纹理参照系关联的操作。,纹理存储器,有关术语及描述： 显存中可绑定到纹理的数据形式： 普通线性存储器（非对齐） 只能与一维或者二维纹理绑定，采用整型纹理拾取坐标，坐标与数据在存储器中的位置相同。 CUDA数组 可以与一维、二维或者三维纹理绑定，纹理拾取坐标为归一化或者非归一化的浮点型。 像元：绑
10、定到纹理的线性存储器或CUDA数组中的元素。,纹理存储器,有关术语及描述： 像元的数据类型：不支持三元组,纹理存储器,有关术语及描述： 纹理存储器有缓存机制 纹理缓存一次预取拾取坐标对应位置附近的几个像元。 绑定到纹理的数据修改后，应重启内核函数且纹理缓存刷新后，才能拾取到被修改的数据。,纹理存储器,纹理存储器的特殊功能： 浮点型纹理拾取坐标 归一化：0.0，1.0) 非归一化：0.0，N)，N为该维度上的像元数。 寻址模式：输入坐标超范围时的处理方式。 钳位模式(clamp)：按上下限坐标拾取像元。 循环模式(wrap)：也称回绕，只用于归一化坐标。 类型转换：对8位或16位整型像元数据，其
11、拾取的返回值可转换为归一化浮点型。 滤波：对CUDA数组绑定的纹理，拾取返回值为浮点型，则可对返回值进行滤波。 最近点取样：适用于查找表。 线性滤波：适用于图像处理。,纹理存储器使用,步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组元素的数据类型 用结构体cudaChannelFormatDesc来描述 struct cudaChannelFormatDesc int x,y,z,w;/多元组数据中每个分量二进制位数 enum cudaChannelFormatKind f; f取值,cudaChannelFormatKindSigned，有符号整型,cuda
12、ChannelFormatKindUnsigned，无符号整型,cudaChannelFormatKindFloat，浮点型,纹理存储器使用,步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组元素的数据类型 如uchar2,则x、y、z、w分别是8，8，0，0，而f取值为cudaChannelFormatKindUnsigned 又如float4,则x,y,z,w分别是32,32,32,32;而f取值为cudaChannelFormatKindFloat,纹理存储器使用,步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组的维度 C
13、UDA数组空间分配函数： cudaMalloc3DArray()：分配一维、二维或三维数组 cudaMallocArray()：一般用于二维数组 CUDA数组空间释放函数：cudaFreeArray() CUDA数组与其他CUDA数组或线性存储器的数据传输： cudaMemcpyToArray()或cudaMemcpy3D(),纹理存储器使用,步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组尺寸 用结构体cudaExtent描述数组三个维度的大小 cudaExtent extent=make_cudaextent(1,8192,0,0) cudaExtent
14、extent=make_cudaextent(1,65535,1,32768,0) cudaExtent extent=make_cudaextent(1,2048,1,2048,1,2048),纹理存储器使用,步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 例，声明一个数据类型为uchar2，643216的CUDA 3D数组。 cudaArray* cuArray=0; cudaExtent extent=make_cudaExtent(64,32,16); cudaChannelFormatDesc desc= cudaCreateChannelDesc(8,8,0,
15、0, cudaChannelFormatKindUnsigned); cudaMalloc3DArray(,纹理存储器使用,步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 CUDA数组声明完毕后，还需赋值。 可用以下函数完成CUDA数组赋值。 CUDA数组之间赋值 cudaMemcpyToArray() CUDA数组与线性存储器之间赋值cudaMemcpy3D(),纹理存储器使用,步骤2：声明纹理参照系 纹理参照系属性 编译时属性：编译前显式声明，编译时确定，一旦确定不能修改。 运行时属性：运行时设定，只适用于与CUDA数组绑定的纹理参照系。 在所有函数体外声明文件域的t
16、exture型变量，最好放在头文件中。形如下示：texture texRef;,纹理存储器使用,步骤2：声明纹理参照系 texture texRef; 编译时属性: Type：纹理拾取返回的数据类型，即CUDA像元 Dim：确定纹理参照系的维度，默认为1，可取值分别为1，2，3 ReadMode：确定返回值是否进行类型转换 cudaReadModeNormalizedFloat：转换为浮点型 cudaReadModeElementType：不转换，此为默认值,纹理存储器使用,步骤2：声明纹理参照系 运行时属性:通过结构体textureReference描述 struct textureRefe
17、rence int normalized;/坐标是否归一化，非零值表示归一化 enum cudaTextureFilterMode filterMode;/滤波模式 enum cudaTextureAddressMode addressMode3; struct cudaChannelFormatDesc channelDesc;/纹理拾取返回的数据类型，要与CUDA数组声明时的类型一致，前已述及。 ,滤波模式有以下两种取值： cudaFilterModePoint表示最近点取样 cudaFilterModeLinear表示线性滤波,寻址模式，它是一个大小为3的数组，对应3个维度，可取以下两种
18、值： cudaAddressModeClamp，表示钳位模式， cudaAddressModeWrap，表示循环模式。,纹理存储器使用,步骤3：纹理绑定 将纹理与数组按纹理参照系绑定，实现数组到纹理的映射。 绑定函数： cudaBindTexture()/用于纹理与线性存储器绑定 cudaBindTextureToArray ()/用于纹理与CUDA数组绑定 解除绑定：cudaUnbindTexture(),纹理存储器使用,步骤4：纹理拾取 函数见下表：,主机端内存,在CUDA编程模型中，可申请使用的内存有如下两类： 可分页内存(pageable memory)：通过操作系统API来分配，如m
19、alloc()。 页锁定内存(page-locked/pinned memory)：始终在物理内存中，并以DMA与设备通信。 分配函数：cudaHostAlloc() 释放函数：cudaFreeHost(),主机端内存,页锁定内存的标志： cudaHostAllocDefault：分配页锁定内存 cudaHostAllocPortable：可在线程间共享 cudaHostAllocWriteCombined： 此类内存在CPU访问时不缓存 仅适用于CPU写-GPU读的模式 cudaHostAllocMapped：将分配的内存映射到CUDA地址空间，以设备指针指向该类内存，可实现zero-copy。,主机端内存,mapped memory： 此类内存在内核函数中可直接访问，不必在内存与显存间进行数据拷贝，即zero-copy。 在内核函数只做少量读写时，这种内存就没有分配显存和数据拷贝的时间。 一块mapped memory有两个地址 主机端地址：由cudaHostAlloc()获得 设备端地址：由cudaHostGetDevicePointer获得 在内核函数中访问页锁定内存时，需要将设备端指针作为参数传入。,参考书,GPU高性能运算之CUDA张舒等,
