CUDA纹理内存相关参数解释

纹理内存的读取模式：
cudaReadModeNormalizedFloat：如果纹理元素是一个16位或8位的整型，那么返回值将作为浮点数返回，
如果是无符号整型，将被映射到[0.0,1.0],如果是有符号整型，将被映射到[-1.0,1.0]
cudaReadModeElementType：那么将返回回原始类型，不做转换

纹理坐标归一化：
默认的，纹理坐标是使用浮点值进行访问:[0,N-1],N是纹理在相应维度上的值，例如，对于一个大小64X32的纹理，其x，y方向的坐标访问范围是：
[0,63],[0,31]
归一化的纹理坐标，将使相应的坐标访问在以下范围：[0,1-1/N],那么对于上面的64x32的纹理，其两个方向的坐标范围就分别是:[0,1-1/64],[0,1-1/32]

滤波模式：
基于输入的坐标值，给出其元素值的返回方式
cudaFilterModePoint:在该模式下，返回相应坐标下，最接近的元素值，类似与最近邻插值
cudaFilterModeLinear:只有当返回值的类型设定为浮点的时候才可以使用，其根据相应元素周围的元素值进行低精度的插值：一维的时候使用线性插值，二维 的时候使用双线性插值，三维的时候使用三线性插值

寻址模式：
寻址模式是一个有3个元素的数组，分别用于指定其相应维度上的寻址方式，用于定义当访问的坐标超出其坐标范围后的处理方式：
cudaAddressModeClamp:是默认的处理方式，对于未归一化的坐标，其截断到范围:[0,N),对于归一化后的坐标范围，截断到:[0.0,1.0)
cudaAddressModeBorder:对于超出范围的坐标，都返回0
cudaAddressModeWrap:只在归一化的坐标下有效，坐标变换公式如下:frac(x) = xfloor(x),其中floor(x)是不大于x的最大整数，
cudaAddressModeMirror:只在归一化坐标下有效，当floor(x)是奇数的时候，每一个坐标变换到frac(x)，当floor(x)是偶数的时候，每一个坐标变换到1-frac(x)

举例：
对于一维信号c(k),k的范围是k=0,1,…,M-1
对于cudaAddressModeClamp，就是将超出的值截断为最大最小值
对于cudaAddressModeBorder，就是将其截断为0
对于cudaAddressModeWrap,对于超出范围的值，就是寻找一个整数p(可以是正负、0)满足：c[(k+p*M)/M] = c[k/M],也就是将其变成周期信号，周期为M
对于cudaAddressModeMirror，也是将其变成周期信号，周期为2*M-2,c[1/M]=c[k/M],对于任何1和k,满足：(1+k)mod(2*M-2)=0
下面给出示例代码：

#include <stdio.h>

texture<float, 1, cudaReadModeElementType> texture_clamp;
texture<float, 1, cudaReadModeElementType> texture_border;
texture<float, 1, cudaReadModeElementType> texture_wrap;
texture<float, 1, cudaReadModeElementType> texture_mirror;

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

/******************************/
/* CUDA ADDRESS MODE CLAMPING */
/******************************/
__global__ void Test_texture_clamping(const int M) {

    printf("Texture clamping - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_clamp, -(float)threadIdx.x));
    printf("Texture clamping - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_clamp, (float)(M + threadIdx.x)));

}

/****************************/
/* CUDA ADDRESS MODE BORDER */
/****************************/
__global__ void Test_texture_border(const int M) {

    printf("Texture border - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_border, -(float)threadIdx.x));
    printf("Texture border - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_border, (float)(M + threadIdx.x)));

}

/**************************/
/* CUDA ADDRESS MODE WRAP */
/**************************/
__global__ void Test_texture_wrap(const int M) {

    printf("Texture wrap - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_wrap, -(float)threadIdx.x/(float)M));
    printf("Texture wrap - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_wrap, (float)(M + threadIdx.x)/(float)M));

}

/****************************/
/* CUDA ADDRESS MODE MIRROR */
/****************************/
__global__ void Test_texture_mirror(const int M) {

    printf("Texture mirror - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_mirror, -(float)threadIdx.x/(float)M));
    printf("Texture mirror - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_mirror, (float)(M + threadIdx.x)/(float)M));

}

/********/
/* MAIN */
/********/
void main(){

    const int M = 4;

    // --- Host side memory allocation and initialization
    float *h_data = (float*)malloc(M * sizeof(float));

    for (int i=0; i<M; i++) h_data[i] = (float)i;

    // --- Texture clamping
    cudaArray* d_data_clamping = NULL; gpuErrchk(cudaMallocArray(&d_data_clamping, &texture_clamp.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_clamping, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_clamp, d_data_clamping); 
    texture_clamp.normalized = false; 
    texture_clamp.addressMode[0] = cudaAddressModeClamp;

    dim3 dimBlock(2 * M, 1); dim3 dimGrid(1, 1);
    Test_texture_clamping<<<dimGrid,dimBlock>>>(M);

    printf("\n\n\n");

    // --- Texture border
    cudaArray* d_data_border = NULL; gpuErrchk(cudaMallocArray(&d_data_border, &texture_border.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_border, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_border, d_data_border); 
    texture_border.normalized = false; 
    texture_border.addressMode[0] = cudaAddressModeBorder;

    Test_texture_border<<<dimGrid,dimBlock>>>(M);

    printf("\n\n\n");

    // --- Texture wrap
    cudaArray* d_data_wrap = NULL; gpuErrchk(cudaMallocArray(&d_data_wrap, &texture_wrap.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_wrap, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_wrap, d_data_wrap); 
    texture_wrap.normalized = true; 
    texture_wrap.addressMode[0] = cudaAddressModeWrap;

    Test_texture_wrap<<<dimGrid,dimBlock>>>(M);

    printf("\n\n\n");

    // --- Texture mirror
    cudaArray* d_data_mirror = NULL; gpuErrchk(cudaMallocArray(&d_data_mirror, &texture_mirror.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_mirror, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_mirror, d_data_mirror); 
    texture_mirror.normalized = true ; 
    texture_mirror.addressMode[0] = cudaAddressModeMirror;

    Test_texture_mirror<<<dimGrid,dimBlock>>>(M);

    printf("\n\n\n");
}

最后的结果应该是：

index  -7  -6  -5  -4  -3  -2  -1  0  1  2  3  4  5  6  7  8  9  10  11
clamp   0   0   0   0   0   0   0  0  1  2  3  3  3  3  3  3  3   3   3
border  0   0   0   0   0   0   0  0  1  2  3  0  0  0  0  0  0   0   0
wrap    1   2   3   0   1   2   3  0  1  2  3  0  1  2  3  0  1   2   3
mirror  1   2   3   3   2   1   0  0  1  2  3  3  2  1  0  0  1   2   3