浅谈GPU编程要点.ppt

纹理存储器使用 步骤1：在主机端声明显存中需要绑定到纹理的线性存储器或CUDA数组。 例，声明一个数据类型为uchar2，64×32×16的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数组之间赋值 cudaMemcpyToArray() CUDA数组与线性存储器之间赋值cudaMemcpy3D() 纹理存储器使用 步骤2：声明纹理参照系 纹理参照系属性 编译时属性：编译前显式声明，编译时确定，一旦确定不能修改。 运行时属性：运行时设定，只适用于与CUDA数组绑定的纹理参照系。 在所有函数体外声明文件域的texture型变量，最好放在头文件中。形如下示： textureType,Dim,ReadMode texRef; 纹理存储器使用 步骤2：声明纹理参照系 textureType,Dim,ReadMode texRef; 编译时属性: Type：纹理拾取返回的数据类型，即CUDA像元 Dim：确定纹理参照系的维度，默认为1，可取值分别为1，2，3 ReadMode：确定返回值是否进行类型转换 cudaReadModeNormalizedFloat：转换为浮点型 cudaReadModeElementType：不转换，此为默认值 纹理存储器使用 步骤2：声明纹理参照系 运行时属性:通过结构体textureReference描述 struct textureReference{ int normalized;//坐标是否归一化，非零值表示归一化 enum cudaTextureFilterMode filterMode;//滤波模式 enum cudaTextureAddressMode addressMode[3]; struct cudaChannelFormatDesc channelDesc;//纹理拾取返回的数据类型，要与CUDA数组声明时的类型一致，前已述及。 } 滤波模式有以下两种取值： cudaFilterModePoint表示最近点取样 cudaFilterModeLinear表示线性滤波 寻址模式，它是一个大小为3的数组，对应3个维度，可取以下两种值： cudaAddressModeClamp，表示钳位模式， cudaAddressModeWrap，表示循环模式。 纹理存储器使用 步骤3：纹理绑定 将纹理与数组按纹理参照系绑定，实现数组到纹理的映射。 绑定函数： cudaBindTexture()//用于纹理与线性存储器绑定 cudaBindTextureToArray ()//用于纹理与CUDA数组绑定 解除绑定：cudaUnbindTexture() 纹理存储器使用 步骤4：纹理拾取 函数见下表： 绑定到纹理的数组 拾取函数 纹理坐标 线性存储器 tex1Dfetch() 整型 CUDA数组 一维 tex1D() 浮点型 二维 tex2D() 三维 tex3D() 主机端内存 在CUDA编程模型中，可申请使用的内存有如下两类： 可分页内存(pageable memory)：通过操作系统API来分配，如malloc()。 页锁定内存(locked/pinned memory)：始终在物理内存中，并以DMA与设备通信。 分配函数：cudaHostAlloc() 释放函数：cudaFreeHost() 主机端内存 页锁定内存的标志： cudaHostAllocDefault：分配页锁定内存 cudaHostAllocPortable：可在线程间共享 cudaHostAllocWriteCombined： 此类内存在CPU访问时不缓存 仅适用于CPU写-GPU读的模式 cudaHostAllocMapped：将分配的内存映射到CUDA地址空间，以设备指针指向该类内存，可实现zero-copy。 主机端内存 mapped memory： 此类内存在内核函数中可直接访问，不必在内存与显存间进行数据拷贝，即zero-copy。 在内核函数只做少量读写时，这种内存就没有分配显存和数据拷贝的时间。 一块mapped memory有两个地址 主机端地址：由cudaHostAlloc(