Cuda里面的Texture Memory

1-dimension linear memory

How TO USE？

1. texture<float,1,cudaReadModeElementType> texreference;
建立一个全局的texture，<type，dimension , readtype>
   type 是数据类型，int ，uchar，float等等
   dimension是1，2 或 3
   readtype有cudaReadModeNormalizedFloat或者cudaReadModeElementType，默认是cudaReadModeElementType。如果是选择cudaReadModeNormalizedFloat，则会对输出数据进行转换，归一化为[0.0，1.0]（对无符号整型），或者[-1.0，1.0]（对有符号整型）

2. 在device端开辟一段空间，并在这段空间传入需要的数据
   如：cudaMemcpy(diarray, harray, sizeof(float)*size, cudaMemcpyHostToDevice);

3. 绑定texture

cudaBindTexture(size_t *offset, const struct textureReference *texref, const void *devPtr, size_t size )

offset是偏移空间，一般写NULL，即不偏移
texref是第一步的texture
devPtr是第二步开辟空间的指针
size是大小

最后记得解绑

cudaUnbindTexture(&text);

代码示例

#include <iostream>#include <cuda.h>#include <cuda_runtime.h>#include <device_functions.h>#include <device_launch_parameters.h>using namespace std;texture<float, 1, cudaReadModeElementType> texreference;__global__ void kernel(float* doarray, int size){    int index;    //calculate each thread global index    index = blockIdx.x*blockDim.x + threadIdx.x;    //fetch global memory through texture reference    doarray[index] = tex1Dfetch(texreference, index);    printf("%f", doarray[index]);    return;}int main(int argc, char** argv){    int size = 3200;    float* harray;    float* diarray;    float* doarray;    //allocate host and device memory    harray = (float*)malloc(sizeof(float)*size);    cudaMalloc((void**)&diarray, sizeof(float)*size);    cudaMalloc((void**)&doarray, sizeof(float)*size);    //initialize host array before usage    for (int loop = 0; loop<size; loop++)        harray[loop] = (float)rand() / (float)(RAND_MAX - 1);    //copy array from host to device memory    cudaMemcpy(diarray, harray, sizeof(float)*size, cudaMemcpyHostToDevice);    //bind texture reference with linear memory    cudaBindTexture(0, texreference, diarray, sizeof(float)*size);    //execute device kernel    kernel << <(int)ceil((float)size / 64), 64 >> >(doarray, size);    //unbind texture reference to free resource    cudaUnbindTexture(texreference);    //copy result array from device to host memory    cudaMemcpy(harray, doarray, sizeof(float)*size, cudaMemcpyDeviceToHost);    //free host and device memory    free(harray);    cudaUnbindTexture(&texreference);    cudaFree(diarray);    cudaFree(doarray);    return 0;}

2-dimension cuda array
1.声明CUDA数组之前，必须先以结构体channelDesc描述CUDA数组中的数据类型。

struct cudaChannelFormatDesc {     int x, y, z, w;     enum cudaChannelFormatKind f;};

cudaChannelFormatKind有三种类型
cudaChannelFormatKindSigned，如果这些成员是有符号整型；
cudaChannelFormatKindUnsigned，如果这些成员是无符号整型；
cudaChannelFormatKindFloat，如果这些成员是浮点型；
示例

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 0, 0,cudaChannelFormatKindunsigned);//一个char2类型

若觉得这麻烦，还有简单的使用方法，使用模板

cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar>();

2.为array开辟一段空间，把数据放进array中
示例

cudaArray* arr;cudaMallocArray(&arr, &desc, data.cols, data.rows);cudaMemcpyToArray(arr, 0, 0, a.data, sizeof(uchar)* data.cols*data.rows, cudaMemcpyHostToDevice);

3.绑定

cudaBindTextureToArray(&texture, arr, &desc);

4.最后别忘记解绑和释放空间

cudaUnbindTexture(&texture);cudaFreeArray(arr);

代码示例

#include <iostream>#include <cuda.h>#include <cuda_runtime.h>#include <device_functions.h>#include <device_launch_parameters.h>#include <opencv.hpp>using namespace std;using namespace cv;texture<uchar,2,cudaReadModeElementType> text;__global__ void test(){    //printf("df");    int x = threadIdx.x + blockDim.x*blockIdx.x;    int y = threadIdx.y + blockDim.y*blockIdx.y;    uchar a = tex2D(text, x, y);    printf("%d %d %d\n", a,x,y);}int main(){    Mat a = imread("t.jpg",0);    cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar>();    cudaArray* arr;    cudaMallocArray(&arr, &desc, a.cols, a.rows);    cudaMemcpyToArray(arr, 0, 0, a.data, sizeof(uchar)* a.cols*a.rows, cudaMemcpyHostToDevice);    text.filterMode = cudaFilterModePoint;    text.addressMode[0] = cudaAddressModeWrap;    text.addressMode[1] = cudaAddressModeWrap;    cudaBindTextureToArray(&text, arr, &desc);    dim3 th(16,16);    dim3 bl(ceil((float)a.cols / 16), ceil((float)a.rows / 16));    test << <bl, th >> >();    cudaUnbindTexture(&text);    cudaFreeArray(arr);    return 0;}

补充说明
texture的属性
1.normalized是设置纹理坐标是否进行归一化，如果是非0，则归一化到[0，1)的坐标进行寻址，否则对尺寸为width, height, depth的纹理使用坐标[0,width-1], [0,height-1], [0,depth-1]寻址。

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

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

参考 http://blog.csdn.net/darkstorm2111203/article/details/4294012 向这位作者致敬