CUDA -- Texture纹理存储器 示例程序

1、纹理存储器的特性

    纹理存储器中的数据以一维、二维或者三维数组的形式存储在显存中，可以通过缓存加速访问，并且可以声明大小比常数存储器要大的多。在kernel中访问纹理存储器的操作称为纹理拾取(texture fetching)。将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding).　显存中可以绑定到纹理的数据有两种，分别是普通的线性存储器和cuda数组。线性存储器是用cudaMalloc()声明的一段连续的线性内存空间，而CUDA数组则是通过cudaMallocArray()进行声明，并且是线性对齐的。CUDA数组对于纹理拾取访问进行了优化，但是设备端只能通过纹理拾取访问， 即我们不能直接操作CUDA数组。可以说，CUDA数组就是为纹理拾取而生的。

在通用计算中，纹理存储器十分适合用于实现图像处理或查找表，并且对数据量较大时的随机数据访问或者非对齐访问也有良好的加速效果。

注：线性存储器只能与一维或二维纹理绑定,采用整型纹理拾取坐标，坐标值与数据在存储器中的位置相同；

        CUDA数组可以与一维、二维、三维纹理绑定，纹理拾取坐标为归一化或者非归一化的浮点型，并且支持许多特殊功能。

属性设置：

纹理参照系中的其它属性可以不必声明， 并在运行时进行修改。这些参数规定了纹理的 寻址模式，是否进行 归一化，以及 纹理滤波。

•TextureReference是如下代码所描述的结构体

struct   textureReference{
int  normalized; //是否归一化 true or false

enum  cudaTextureFilterMode    filterMode;    //滤波模式：cudaFilterModepoint和cudaFilterModeLinear

enum  cudaTextureAddressMode    addressMode[3];    //寻址模式：cudaAddressModeClump和cudaAddressModeWarp

struct  cudaChannelFormatDesc    channelDesc;
}

例如：

2、纹理缓存的作用

1）纹理内存中的数据可以被重复利用，当需要的数据存在于纹理缓存中，就不用再去显存读取了

2）纹理拾取可以读取纹理坐标附近的几个象元，提高局部性的访问效率，实现滤波模式。换言之，对于图像滤波具有很好的性能。

3、一维纹理存储器的使用方法

1）在host端声明线性内存空间

分配线性内存空间的方法常见有cudaMalloc(), 由cudaMallocPitch()或者cudaMalloc3D()分配的线性空间是经过填充对齐的线性内存。

2）声明纹理参考系

texture<Type, Dim, ReadMode> texRef;
//Type指定数据类型，特别注意：不支持3元组
//Dim指定纹理参考系的维度，默认为1
//ReadMode可以是cudaReadModelNormalizedFloat或cudaReadModelElementType(默认)

注意：type仅限于基本整型、单精度浮点类型和CUDA运行时提供的1元组、2元组和四元祖。

texture<int, 1, cudaReaModeElementType> texRef;这样就声明了一个一维整数的纹理。

3）绑定数据到纹理

通过cudaBindTexture()函数将纹理参考连接到内存。将1维线性内存绑定到1维纹理。

cudaError_t cudaBindTexture(    size_t *offset, const struct textureReference *texref, const void *devPtr,const struct cudaChannelFormatDesc * desc,size_t  size = UINT_MAX )  
例如：cudaBindTexture(0, texRef, d_data);

4）设备端的纹理拾取

在kernel中对纹理存储器进行访问，要通过tex1Dfetch()函数,通过纹理参考和给出的坐标位置就可以取得数据了。

type tex1Dfetch(texture<type,1,ReadMode> texRef, int x);

通过以上四步，我们就可以在程序中利用纹理内存进行数据处理了。例子如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <stdio.h>
 
#define N 256
 
//声明纹理参考系
texture<int, 1, cudaReadModeElementType> texRef;
 
 
//CPU串行向量复制
void copyCPU(int* source, int* target, int size) {
    for (int i = 0; i < size; i++) {
        target[i] = source[i];
    }
}
 
//GPU使用全局内存实现向量复制
__global__ void kernel(int* source, int* target, int size) {
    int index = threadIdx.x + blockDim.x * blockIdx.x;
    if (index < size) {
        target[index] = source[index];
    }
}
 
void copyGPU(int* source, int* target, int size) {
    int Size = size * sizeof(int);
    int* d_data = NULL;
    //分配显存，并传递数据到显存上
    cudaMalloc((void**)&d_data, Size);
    cudaMemcpy(d_data, source, Size, cudaMemcpyHostToDevice);
 
    int* d_Result = NULL;
    cudaMalloc((void**)&d_Result, Size);
 
    //创建Block Grid size 
    dim3 blockNumber(16);
    dim3 gridNumber((size + blockNumber.x - 1) / blockNumber.x);
    kernel<< <gridNumber, blockNumber >> > (d_data, d_Result, size);
    //拷贝数据回主机端
    cudaMemcpy(target, d_Result, Size, cudaMemcpyDeviceToHost);
 
 
    cudaFree(d_data);
    cudaFree(d_Result);
}
 
 
//GPU使用纹理内存实现向量复制
__global__ void kernelTexture(int* target, int size) {
    int index = threadIdx.x + blockDim.x * blockIdx.x;
    if (index < size) {
        //纹理拾取 tex1Dfetch（纹理参考对象， 位置）
        target[index] = tex1Dfetch(texRef, index);
    }
}
 
void copyGPUTexture(int* source, int* target, int size) {
    int Size = N * sizeof(int);
    //locate Device memory and copy host data to device data
    int* d_data = NULL;
    cudaMalloc((void**)&d_data, Size);
    cudaMemcpy(d_data, source, Size, cudaMemcpyHostToDevice);
 
    int* d_Result = NULL;
    cudaMalloc((void**)&d_Result, Size);
 
 
    //绑定纹理内存的数据，从全局内存到纹理内存的关联
    cudaBindTexture(0, texRef, d_data);
 
    dim3 blockNumber(16);
    dim3 gridNumber((size + blockNumber.x - 1) / blockNumber.x);
    kernelTexture << <gridNumber, blockNumber >> > (d_Result, N);
    cudaMemcpy(target, d_Result, Size, cudaMemcpyDeviceToHost);
    
    //解除纹理绑定
    cudaUnbindTexture(texRef);
    cudaFree(d_data);
    cudaFree(d_Result);
}
 
int main()
{
    int *source, *target;
    source = (int *)malloc(sizeof(int) * N);
    target = (int *)malloc(sizeof(int) * N);
    for (int i = 0; i < N; i++) {
        source[i] = i + 1;
        target[i] = 0;
    }
 
    cudaEvent_t t1, t2;
    cudaEventCreate(&t1);
    cudaEventCreate(&t2);
    cudaEventRecord(t1, 0);
 
    int number = 2;
    switch (number) {
    case 0:
        copyCPU(source, target, N);
        break;
    case 1:
        copyGPU(source, target, N);
        break;
    case 2:
        copyGPUTexture(source, target, N);
        break;
    }
    
    cudaEventRecord(t2, 0);
    cudaEventSynchronize(t2);
    printf("
Source data:
");
    for (int i = 0; i < N; i++) {
        printf("%5d", source[i]);
    }
    printf("
Target data:
");
    for (int i = 0; i < N; i++) {
        printf("%5d", target[i]);
    }
    printf("
");
 
    float gpuTime = 0;
    cudaEventElapsedTime(&gpuTime, t1, t2);
    printf("GPU Time is：%f
", gpuTime);
    return 0;
}

4、二维纹理存储器的使用方法

1）声明CUDA数组

使用CUDA数组主要通过三个函数使用：cudaMallocArray(), cudaMemcpyToArray(),cudaFreeArray(). 在声明CUDA数组之前，必须先描述结构体cudaChannelFormatDes()的组件数量和数据类型。

结构体定义：

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

x, y, z和w分别是每个返回值成员的位数，而f是一个枚举变量，可以取一下几个值：

• cudaChannelFormatKindSigned，如果这些成员是有符号整型
• cudaChannelFormatKindUnsigned，如果这些成员是无符号整型
• cudaChannelFormatKindFloat，如果这些成员是浮点型

CUDA数组的创建方法：

cudaChannelFormatDesc  channelDesc = cudaCreateChannelDesc<float>();//声明数据类型
cudaArray *cuArray;
//分配大小为W*H的CUDA数组
cudaMallocArray(&cuArray, &channelDesc, Width, Height);

CUDA 数组的复制：

cudaMemcpyToArray(struct cuArray* dstArray, size_t dstX, size_t dstY,  const void *src, size_t  count, enum cudaMemcpyKind kind);
  
//函数功能是：把数据src复制到CUDA 数组dstArray中，复制的数据字节大小为count，从src的左上角(dstX,dstY)开始复制。
//cudaMemcpyKind 复制方向有：hostTohost, hostTodevice, deviceTohost,device todevice(简写方式).

2）声明纹理参考系

//声明一个float 类型的2维的纹理，读取模式为cudaReadModeElementType，也可以声明其他的读取模式
texture<float, 2, cudaReadModeElementType> texRef;

3）绑定CUDA数组到纹理

调用cudaBindTextureToArray()函数把CUDA数组和纹理连接起来。

4）设备端的纹理拾取

和一维纹理内存的tex1Dfetch()不同，需要使用tex1D()、tex2D()、tex3D()这三个函数，分别是用在1D/2D/3D的纹理。

二维纹理内存使步骤和一维区别不大，主要是利用CUDA数组来进行绑定。具体实例如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <stdio.h>
#define W 16
#define H 16
//声明一个float 类型的2维的纹理，读取模式为cudaReadModeElementType
texture<float, 2, cudaReadModeElementType> texRef;
 
__global__ void addKernel(float *array, const int w, const int h)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int idy = threadIdx.y + blockDim.y * blockIdx.y;
    
    if (idx < w && idy < h) {
        //二维纹理索取，根据坐标idx idy进行索引
        array[idx + idy*w] = tex2D(texRef, idx, idy);
    }
}
 
int main()
{
    float *h_data = (float*)malloc(W*H * sizeof(int));
    float *h_result = (float *)malloc(W*H * sizeof(int));
    for (int i = 0; i < H; i++) {
        for (int j = 0; j < W; j++) {
            h_data[i*W + j] = i*W + j;
        }
    }
   //分配CUDA数组
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();//声明数据类型
    cudaArray *cuArray;
    //分配大小为W*H的CUDA数组
    cudaMallocArray(&cuArray, &channelDesc, W, H);
    //从数据（0，0）开始复制数据，大小为W*H*sizeof(float)
    cudaMemcpyToArray(cuArray, 0, 0, h_data, W*H*sizeof(float), cudaMemcpyHostToDevice);
 
    //设置纹理属性
    texRef.addressMode[0] = cudaAddressModeWrap;//寻址方式
    texRef.addressMode[1] = cudaAddressModeWrap;//寻址方式 如果是三维数组则设置texRef.addressMode[2]
 
    texRef.normalized = false;//是否对纹理坐标归一化
    texRef.filterMode = cudaFilterModePoint;//纹理的滤波模式：最近点取样和线性滤波  cudaFilterModeLinear
 
    //纹理绑定，CUDA数组和纹理参考的连接
    cudaBindTextureToArray(&texRef, cuArray, &channelDesc);
 
    //设备内存结果
    float *d_data = NULL;
    cudaMalloc((void**)&d_data, W*H * sizeof(float));
    dim3 dimBlock(16, 16);
    dim3 dimGrid((W + dimBlock.x - 1) / dimBlock.x, (H + dimBlock.y - 1) / dimBlock.y);
    addKernel << <dimGrid, dimBlock >> > (d_data, W ,H);
    cudaMemcpy(h_result, d_data, W*H * sizeof(float), cudaMemcpyDeviceToHost);
 
    //解除绑定
    cudaUnbindTexture(texRef);
 
    //释放CDUA数组
    cudaFreeArray(cuArray);
    //释放显存
    cudaFree(d_data);
    
 
    printf("Origin data:
");
    for (int i = 0; i < W*H; i++) {
        printf("%8.1f", h_data[i]);
    }
 
    printf("
Result:
");
    for (int i = 0; i < W*H; i++) {
        printf("%8.1f", h_result[i]);
    }
 
    return 0;
}

参考：http://blog.csdn.net/augusdi/article/details/12187159

         《基于CUDA的并行程序设计》刘金硕

纹理存储器非常适合实现图像处理和查找表LUT，对大量数据的随机访问或非随机访问也有良好的加速效果。第一次接触纹理存储器，写了以下一个小程序。

/*
* Copyright 徐洪志(西北农林科技大学.信息工程学院).  All rights reserved.
* Data: 2012-4-20
*/
//
// 此程序是演示了1D和2D纹理存储器的使用
#include <stdio.h>
#include <cutil_inline.h>
#include <iostream>
using namespace std;
 
texture<float> texRef1D;     // 1D texture
texture<float, 2> texRef2D;  // 2D texture
 
// 1D 纹理操作函数
__global__ void Texture1D(float *dst, int w, int h)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
 
    if(x<w && y<h)
        dst[offset] = tex1Dfetch(texRef1D, offset);
}
// 2D 纹理操作函数
__global__ void Texture2D(float *dst, int w, int h)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
 
    dst[offset] = tex2D(texRef2D, x, y);
}
int main(int argc, char **argv)
{
    CUT_DEVICE_INIT(argc, argv);  // 启动 CUDA
    /// 1D 纹理内存
    cout << "1D texture" << endl;
    float *host1D = (float*)calloc(10, sizeof(float)); // 内存原数据
    float *hostRet1D = (float*)calloc(10, sizeof(float));// 内存保存返回数据
 
    float *dev1D, *devRet1D; // 显存数据
    int i;
    cout << " host1D:" << endl;
    for(i = 0; i < 10; ++i)  // 初始化内存原数据
    {
        host1D[i] = i * 2;
        cout << "  " << host1D[i] << " ";
    }
    cutilSafeCall( cudaMalloc((void**)&dev1D, sizeof(float)*10));  // 申请显存空间
    cutilSafeCall( cudaMalloc((void**)&devRet1D, sizeof(float)*10));
    cutilSafeCall( cudaMemcpy(dev1D, host1D, sizeof(float)*10, cudaMemcpyHostToDevice)); // 将内存数据拷贝入显存
    cutilSafeCall( cudaBindTexture(NULL, texRef1D, dev1D, sizeof(float)*10));  // 将显存数据和纹理绑定
    
    Texture1D<<<10, 1>>>(devRet1D, 10, 1);  // 运行1D纹理操作函数
 
    cutilSafeCall( cudaMemcpy(hostRet1D, devRet1D, sizeof(float)*10, cudaMemcpyDeviceToHost)); // 将显存数据拷贝入内存
    // 打印内存数据
    cout << endl << " hostRet1D:" << endl;
    for(i = 0; i < 10; ++i)
        cout << "  " << hostRet1D[i] << " ";
 
    cutilSafeCall( cudaUnbindTexture(texRef1D));  // 解绑定
    cutilSafeCall( cudaFree(dev1D));  // 释放显存空间
    cutilSafeCall( cudaFree(devRet1D)); 
    free(host1D);  // 释放内存空间
    free(hostRet1D);
    
    /// 2D 纹理内存    
    cout << endl << "2D texture" << endl;
    int width = 5, height = 3;
    float *host2D = (float*)calloc(width*height, sizeof(float));    // 内存原数据
    float *hostRet2D = (float*)calloc(width*height, sizeof(float)); // 内存返回数据 
 
    cudaArray *cuArray;  // CUDA数组
    float *devRet2D;     // 显存数据
    int row, col;
    cout << " host2D:" << endl;
    for(row = 0; row < height; ++row)  // 初始化内存原数据
    {
        for(col = 0; col < width; ++col)
        {
            host2D[row*width + col] = row + col;
            cout << "  " << host2D[row*width + col] << " ";
        }
        cout << endl;
    }
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    cutilSafeCall( cudaMallocArray(&cuArray, &channelDesc, width, height));  // 申请显存空间
    cutilSafeCall( cudaMalloc((void**) &devRet2D, sizeof(float)*width*height));
    cutilSafeCall( cudaBindTextureToArray(texRef2D, cuArray)); // 将显存数据和纹理绑定
    cutilSafeCall( cudaMemcpyToArray(cuArray, 0, 0, host2D, sizeof(float)*width*height, cudaMemcpyHostToDevice)); // 将内存数据拷贝入CUDA数组
 
    dim3 threads(width, height);
    Texture2D<<<1, threads>>>(devRet2D, width, height);  // 运行2D纹理操作函数
 
    cutilSafeCall( cudaMemcpy(hostRet2D, devRet2D, sizeof(float)*width*height, cudaMemcpyDeviceToHost)); // 将显存数据拷贝入内存
    // 打印内存数据
    cout << " hostRet2D:" << endl;
    for(row = 0; row < height; ++row)
    {
        for(col = 0; col < width; ++col)
            cout << "  " << hostRet2D[row*width + col] << " ";
        cout << endl;
    }
 
    cutilSafeCall( cudaUnbindTexture(texRef2D)); // 解绑定
    cutilSafeCall( cudaFreeArray(cuArray));  // 释放显存空间
    cutilSafeCall( cudaFree(devRet2D));
    free(host2D);  // 释放内存空间
    free(hostRet2D);
 
    CUT_EXIT(argc, argv);  // 退出CUDA
}