CUDA编程经验技术总结 系列之《内存模块》

        共享内存、文理内存是cuda中经常用到的内存模式。

一、共享内存和同步：

、共享内存和同步：实现方式：__share__ float  array
;
N一般为threadsPerBlock。

N的确定方式之一：对于共享内存来说，由于编译器将为每个线程块生成共享变量的一个副本，因此，只需根据线程块中线程的数量来分配内存。

实用情形：

二、纹理内存（只读）：

1、一维纹理内存的使用（必须先把数据cudaMemcpy到设备上，然后把设备上的该数据绑定到纹理内存，见文档尾示例）：

定义：       texture <float>  texIn;

texture <float,1, cudaReadModeElementType>  texIn;

绑定：       cudaBindTexture(NULL,texIn,hostDataPtr,hostDataPtr.Size());

取纹理内存： float c=tex1Dfetch(texIn,x); x=threadIdx.x+blockIdx.x*blockDim.x;

解绑定：     cudaUnbindTexture(texIn);   用完后要解绑定。

 

2、二维纹理内存的使用:

定义：       texture <float,2, cudaReadModeElementType>  texIn;

绑定：cudaChannelFormatDescdesc=cudaCreateChannelDesc<float>();

cudaBindTexture2D(NULL,texIn,hostDataPtr, desc,W,H,sizeof(float)*Dim);

或：

cudaChannelFormatDescdesc=cudaCreateChannelDesc<float>();

cudaArray *devArray;

cudaMallocArray(&devArray, &desc, imgW, imgH);

cudaMemcpyToArray(devArray, 0, 0, imgPtr,imageSize,cudaMemcpyHostToDevice);

cudaBindTextureToArray(textImg, devArray, desc);

 

取纹理内存： float c=tex2D(texIn,x,y); x=threadIdx.x+blockIdx.x*blockDim.x;

解绑定：     cudaUnbindTexture(texIn);   用完后要解绑定。

linearmemory（我理解的线性内存）只能与一维纹理绑定，只有CUDA Array才能与二维纹理和三维纹理绑定。但《*实战》中的实例，linear memory与二维纹理绑定的例子确实也在用。但为什么运行代码就会出问题呢？

个人猜测：当使用linear memory与二维纹理绑定时，每一维的大小必须都是2的n次幂，且两个维度的大小需相同。至于为什么，我想是因为如果不是2的n次方，那么gpu读取global memory时就不能很直接得到这两个值（因为读取操作每次必然是按照2的n次幂大小的一个块来读的）

补充知识：从CUDA的内核函数访问纹理存储器的操作被称为纹理拾取(texture fetching)。纹理拾取使用的坐标与数据在显存中的地址可以不同，两者通过纹理参照系（texture reference）约定从数据的地址到纹理坐标的映射方式。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding)。显存中可以绑定到纹理的数据有两种，分别是普通的线性内存（Linear Memroy）和CUDA数组（CUDA Array）。CUDA数组则为纹理访问进行了优化，并且在Device端中只能通过纹理拾取访问。

 

Texture结构体

texture类型是从底层的textureReference中派生而来的。TextureReference是一个下面的代码描述的结构体。

structtextureReference {

     int normalized;

     enum cudaTextureFilterMode filterMode;

     enum cudaTextureAddressModeaddressMode[3];

     struct cudaChannelFormatDesc channelDesc;

}

normalized 设置是否对纹理坐标是否进行归一化。如果normalized是一个非零值，那么就会使用归一化到[0，1)的坐标进行寻址，否则对尺寸为width, height, depth的纹理使用坐标[0,width-1], [0,height-1], [0,depth-1]寻址。例如，一个尺寸为64×32的纹理可以通过x维度范围为[0，63]，y维度范围[0,31]的坐标寻址。如果采用归一化方式对尺寸为64×32的纹理进行寻址，在x和y维度上的坐标就都是[0.0,1.0)。这样就可以保证纹理的坐标与纹理的尺寸无关。

filterMode用于设置纹理的滤波模式，即如何根据坐标计算返回的纹理值。滤波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。滤波模式为CudaFilterModePoint时，返回值是与坐标最接近的像元的值。CudaFilterModeLinear模式只能对返回值为浮点型的纹理使用，启用这一种模式时将拾取纹理坐标周围的像元，然后根据坐标与这些像元之间的距离进行插值计算。对一维纹理可以使用线性滤波，对二维纹理可以使用双线性滤波。返回值会是对最接近纹理坐标的两个像元（对一维纹理），四个像元（对二维纹理）或者八个像元（对三维纹理）进行插值后得到的值。

addressmode说明了寻址模式，即如何处理超出寻址范围的纹理坐标；addressmode是一个大小为3的数组，三个元素分别说明对第一、二、三个纹理坐标的取址模式；取址模式可以是cudaAddressMo
4000
deClamp或cudaAddressModeWrap中的一种，前者将超出寻址范围的纹理坐标”钳位”到寻址范围内的最大或最小值，后者将超出寻址范围的纹理坐标“折叠”进合理范围。cudaAddressModeWrap只支持归一化的纹理坐标。

对非归一化的坐标，如果寻址的坐标超过了范围[0，N]，大于N的坐标将被钳位，设为N-1。

对归一化的坐标，有钳位和循环两种处理方式，在钳位方式下，超过[0.0，1.0)范围的坐标将被钳位到[0.0，1.0)；循环方式一般用于周期循环纹理，它只使用了纹理坐标中有用的小数部分，例如1.25会被当作0.25处理，而-1.25则会被当成0.75处理。

channelDesc描述纹理获取返回值类型。纹理参照系的返回值类型描述必须和与之绑定的CUDA array的数据类型描述相同，或者和与之绑定的线性内存中的元素类型相同。

normalized, addressMode和filterMode可以直接在主机端代码中修改。它们只适用于与CUDA数组绑定的纹理参照系。

 

N、归约算法：

最佳要求：array[threadsPerBlock]，设置的threadsPerBlock 必须是2的指数。