CUDA优化实例（四）纹理内存

CUDA优化实例（四）纹理内存

本文参考：http://www.cnblogs.com/cuancuancuanhao/p/7809713.html
本节的内容可能和标题不服，本节主要将纹理内存的使用，它到底有什么有什么速度的提升，优化体现在哪里，我下节会写，本节主要写一个纹理内存的例子。

引言

• 纹理内存与全局内存一样，都在DRAM上，所以纹理内存的容量是很大的
• 纹理内存是专为图像而设计的，图像处理的特点是局部内存访问较为频繁，即其内存的请求有局部性，如取图像的像素一般不一个一个取，一般取或处理一片。全局内存只有在合并的情况下高效，所以纹理内存专门为图像处理设计，其内存访问有空间相干性，其具体是怎么访问纹理内存的，是怎么优化的，官方网站上没有介绍，书也很少，暂时知道，warp中所有的内存访问如果是区域性的，使用纹理内存最好。
• 纹理内对有关图像的一些操作提供了方便，如边界问题，插值问题等。

实验

纹理内存的使用有两种方式，分别是Texture Object和Texture Reference，前者可在程序中动态生成，后者的部分在编译期间静态生成（确定了），有点像静态数组和动态数组的意思。如当事先不知道要处理的数据的类型时，使用Texture Object是个不错的选择。本文以Texture Reference为例。

使用纹理内存步骤较复杂简单介绍：

1. 声明Texture Referece：
   texture<uchar, cudaTextureType2D, cudaReadModeElementType> texRef;
2. 申请设备内存，可以是线性内存或CUDA内存，其不同[CUDA纹理笔记]
   (http://blog.youkuaiyun.com/fb_help/article/details/79548209)
   这里以2维CUDA Array为例：

    // 申请 cuda 数组并拷贝数据
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsigned);
    cudaArray* cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);
    cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);

cudaCreateChannelDesc 函数中的返回的是”Channel format kind “，其参数的含义是每维数据的字数，如果是float,int等为32位，如果是char，unsigned char等为8位。最后一个参数cudaChannelFormatKindUnsigned 是数据的类型。它与textture reference里一样的datatype要对应。
cudaMemcpyToArray 是将host段的数据拷贝到CUDA Array。
3. 设置Texture Reference参数：

    // 指定纹理引用参数
    //一维的超越边界取0
    texRef.addressMode[0] = cudaAddressModeBorder;
    //二维的超越边界取0
    texRef.addressMode[1] = cudaAddressModeBorder;
    //不插值，取整
    texRef.filterMode = cudaFilterModePoint;
    //不使用归一化纹理坐标
    texRef.normalized = 0;

1. 绑定纹理内存

    // 绑定纹理引用
    cudaBindTextureToArray(texRef, cuArray, channelDesc);

1. 运行核函数。

代码：

texture memory.cu：

#include <stdio.h>
#include <stdlib.h>
#include <malloc.h>
#include <cuda_runtime_api.h>
#include "device_launch_parameters.h"

#define DEGRE_TO_RADIAN(x) ((x) * 3.1416f / 180)
#define CEIL(x,y) ((((x) + (y) - 1))/ (y) )

typedef unsigned char uchar;
// 声明纹理引用
texture<uchar, cudaTextureType2D, cudaReadModeElementType> texRef;

// 简单的线性变换
__global__ void transformKernel(uchar* output, int width, int height, float theta)
{
    // 计算正规化纹理坐标
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idy = blockIdx.y * blockDim.y + threadIdx.y;

    // 正规化和平移
    float u = idx - width/2;
    float v = idy - height/2;

    //printf("u:  %f, v: %f  \n",u,v);

    // 旋转

    float tu = u * __cosf(theta) - v * __sinf(theta) + width/2;
    float tv = v * __cosf(theta) + u * __sinf(theta) + height/2;

    //tu = idx;
    //tv = idy;

    //printf("\n(%2d,%2d,%2d,%2d)->(%f,%f,%d)", 
    //    blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y, tu, tv,tex2D(texRef, tu, tv));

    // 纹理内存写入全局内存
    output[idy * width + idx] = tex2D(texRef, tu, tv);
}

int main()
{
    // 基本数据
    int i;
    uchar *h_data, *d_data;
    int width = 16;
    int height = 16;
    float angle = DEGRE_TO_RADIAN(0);

    int size = sizeof(uchar)*width*height;
    h_data = (uchar *)malloc(size);
    cudaMalloc((void **)&d_data, size);

    for (i = 0; i < width*height; i++)
        h_data[i] = i;

    printf("\n\n");
    for (i = 0; i < width*height; i++)
    {
        printf("%d ", h_data[i]);
        if ((i + 1) % width == 0)
            printf("\n");
    }

    // 申请 cuda 数组并拷贝数据
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsigned);
    cudaArray* cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);
    cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);

    // 指定纹理引用参数，注意与纹理对象的使用不一样
    texRef.addressMode[0] = cudaAddressModeBorder;
    texRef.addressMode[1] = cudaAddressModeBorder;
    texRef.filterMode = cudaFilterModePoint;
    texRef.normalized = 0;

    // 绑定纹理引用
    cudaBindTextureToArray(texRef, cuArray, channelDesc);

    // 运行核函数
    dim3 dimBlock(16, 16);
    dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
    printf("x:  %d , y:  %d  \n",dimGrid.x,dimGrid.y);
    transformKernel << <dimGrid, dimBlock >> > (d_data, width, height, angle);
    cudaDeviceSynchronize();

    // 结果回收和检查结果
    cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);

    printf("\n\n");
    for (i = 0; i < width*height; i++)
    {
        printf("%d ", h_data[i]);
        if ((i + 1) % width == 0)
            printf("\n");
    }

    // 回收工作
    cudaFreeArray(cuArray);
    cudaFree(d_data);

    return 0;
}

---

结果

当angle = 180时

当angle = 30时

当angle = 0时

分析

180旋转出现0是因为texRef.addressMode[1] = cudaAddressModeBorder;和#define DEGRE_TO_RADIAN(x) ((x) * 3.1415926 / 180)pi精度不够，
#define DEGRE_TO_RADIAN(x) ((x) * acos(-1) / 180) 会好一些。

结论

本节介绍个纹理内存的例子，下节会介绍纹理内存比全局内存的优越性。