cuda学习笔记(4)

上篇文章介绍cuda常量内存的特点及使用方法。本篇记录cuda纹理内存( texture )的特点及使用方法。

1. 纹理内存

cuda编程经常要对二维或者三维数据进行操作，为了加速内存读写，需要使用纹理内存。纹理内存不能单独使用，必读绑定到全局内存上，纹理内存（ texture memory ）实质上是全局内存的一个特殊形态，全局内存被绑定为纹理内存，对其的读写操作将通过专门的 texture cache（纹理缓存）进行，其实称为纹理缓存更加贴切。

纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性插值。纹理缓存针对二维空间的局部性访问进行了优化，所以通过纹理缓存访问二维矩阵的邻域会获得加速。纹理缓存不需要满足全局内存的合并访问条件。此外，纹理内存访问会自动处理边界条件，对于越界有自动处理机制。

纹理可以是一段连续的设备内存，也可以是一个CUDA数组。但是CUDA数组对局部寻址有优化，称为“块线性”，原理是将邻域元素缓存在同一条cache线上，这将加快邻域内的寻址，但是对于设备内存，并没有“块线性”。所以，选择采用CUDA数组，还是设备内存，需要根据实际情况决定，将数据copy至CUDA数组是很耗时的。纹理的一个元素称为texels。纹理最多支持三维，width，x方向的维数；height，y方向的维数；depth，z方向的维数。

纹理的创建通过texture object或者texture reference，注意只有计算能力3.0以上的设备支持texture object。texture object在运行时创建，texture reference在编译期创建，运行时绑定到实际的纹理。texture reference必须定义为全局变量，不能作为函数的参数。

2. 纹理内存的使用

纹理内存使用流程如下：

a. 声明纹理参考系(注：纹理参照系必须定义在所有函数体外)

texture<uchar4,cudaTextureType2D,cudaReadModeNormalizedFloat> tex;

b.设置纹理通道格式等，并以此申请global内存

cudaChannelFormatDesc channelDesc=cudaCreateChannelDesc<uchar4>();
cudaArray *cuArray;
cudaMallocArray(&cuArray,&channelDesc,width,height);

c.初始化申请的cuda global内存

cudaMemcpyToArray(cuArray,0,0,src.data,size,cudaMemcpyHostToDevice);

d.设置纹理参考系的属性，并绑定到cuda global 内存上

纹理参考系的结构体如下：

struct textureReference
{
    int normalized;
    enum cudaTextureFilterMode filterMode;
    enum cudaTextureAddressMode addressMode[3];
    struct cudaChannelFormatDesc channelDesc;
}

normalized设置读取模式：是否对纹理坐标归一化
1）归一化浮点模式，fetch将返回【0.0,1.0】的归一化之后的数值，输入数据类型为8位、16位整型或浮点数；（2）元素类型模式，fetch返回原始数值。 纹理支持归一化的浮点坐标，计算方式如下：[0.0，1-1/N]，N为当前维度的维数。归一化坐标在某些情况下非常方便，如放大缩小操作。

filterMode用于设置纹理的滤波模式
定义了fetch返回结果的计算方式。有两种模式：cudaFilterModePoint or cudaFilterModeLinear 。 cudaFilterModePoint：点模式，返回最接近的一个点，即最近邻插值。插值公式 tex(x) = T(i)，i=floor(x)，注意是对坐标向下取整，所以一般对输入坐标值+0.5，避免无法精确表示的某些数值出现错误取值，如 x=3，实际是2.99999，此时实际获取的是x=2的元素。 cudaFilterModeLinear：线性模式，即线性插值，对于一维纹理，两点插值；对于二维纹理，四点插值；对于三维纹理，八点插值。线性模式只有在fetch返回浮点类型数据（注意并非指read mode的归一化浮点模式）下才有效。

addressMode说明了寻址方式
寻找模式，定义超出坐标范围的取值（越界情况）。一共四种模式：cudaAddressModeBorder, cudaAddressModeClamp（默认模式）, cudaAddressModeWrap, and cudaAddressModeMirror，后两个只能在归一化坐标时使用。addressing mode定义为一个三维向量，分别代表纹理各个方向上的寻址模式。
cudaAddressModeClamp：超出范围就用边界值代替，示意： AA | ABCDE | EE
cudaAddressModeBorder：超出范围就用零代替，示意： 00 | ABCDE | 00
cudaAddressModeWrap：重叠模式（循环），示意： DE | ABCDE || AB
cudaAddressModeMirror：镜像模式，示意： BA | ABCDE | ED

绑定示例：

tex.addressMode[0]=cudaAddressModeWrap;
tex.addressMode[1]=cudaAddressModeWrap;
tex.filterMode = cudaFilterModeLinear;
tex.normalized =false;
cudaBindTextureToArray(tex,cuArray,channelDesc)

e.使用内存（纹理拾取）

tex2D(tex,x,y);

f.解开绑定并释放cuda global内存。

cudaUnbindTexture(tex);
cudaFree(cuArray);

下面用一个展示CUDA二维纹理内存+OpenCV图像滤波示例（注意：由于纹理内存使用浮点型4字节，对于opencv读取RGB三通道，应使用 cv::cvtColor(src, src, CV_BGR2BGRA) 转换为RGBA三通道格）：

#include "cuda.h"
#include "cuda_runtime.h"
#include "opencv2/core/core.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/opencv.hpp"
#include "stdio.h"
using namespace std;
using namespace cv;

texture<uchar4,cudaTextureType2D,cudaReadModeNormalizedFloat> tex;
//cudaReadModeNormalizedFloat 为了让tex2D读取，格式可转换。

__global__ void smooth_kernel(char *img,int width,int heigth,int channels)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int offset = x + y*blockDim.x+gridDim.x;

    //若使用归一化
    float u = x/(float)width;
    float v = y/(float)heigth;

    //如果使用cudaReadModeElementType，则读取uchar4不能转为float
    //图像边界越界自动处理！！
    float4 pixel    = tex2D(tex,x,y);
    float4 left     = tex2D(tex,x-1,y);
    float4 right    = tex2D(tex,x+1,y);
    float4 top      = tex2D(tex,x,y-1);
    float4 botton   = tex2D(tex,x,y+1);

    img[(y*width+x)*channels+0] = (left.x+right.x+top.x+botton.x)/4*255;
    img[(y*width+x)*channels+1] = (left.y+right.y+top.y+botton.y)/4*255;
    img[(y*width+x)*channels+2] = (left.z+right.z+top.z+botton.z)/4*255;
    img[(y*width+x)*channels+3] = 0;
}


#define IMAGE_DIR "1.jpg"

int main(int argc,char **argv)
{
        Mat src = imread(IMAGE_DIR,IMREAD_COLOR);

    //注意：纹理内存绑定限制每行应该为256字节，也有非256字节掉对齐方法
    //  这里为了方便，我们将图片resize位256*256大小
        resize(src, src, Size(256, 256));

        //为了使用float的纹理，将RGB三字节的格式改为BGRA四字节掉存储方式
        cvtColor(src, src, CV_BGR2BGRA);

        int rows=src.rows;
        int cols=src.cols;
        int channels=src.channels();
        int width=cols,height=rows,size=rows*cols*channels;

        cudaChannelFormatDesc channelDesc=cudaCreateChannelDesc<uchar4>();
        cudaArray *cuArray;
        cudaMallocArray(&cuArray,&channelDesc,width,height);
        cudaMemcpyToArray(cuArray,0,0,src.data,size,cudaMemcpyHostToDevice);

        tex.addressMode[0]=cudaAddressModeWrap; 
        tex.addressMode[1]=cudaAddressModeWrap;
        tex.filterMode = cudaFilterModeLinear;  
        tex.normalized =false;          //不使用归一化

        cudaBindTextureToArray(tex,cuArray,channelDesc);


        Mat out=Mat::zeros(width, height, CV_8UC4);
        char *dev_out=NULL;
        cudaMalloc((void**)&dev_out, size);

        dim3 dimBlock(16, 16);
        dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y);
        smooth_kernel<<<dimGrid,dimBlock,0>>>(dev_out,width,height,channels);

        cudaMemcpy(out.data,dev_out,size,cudaMemcpyDeviceToHost);

        imshow("orignal",src);
        imshow("smooth_image",out);
        waitKey(0);

        cudaFree(dev_out);
        cudaFree(cuArray);
        cudaUnbindTexture(tex);
        return 0;

}

引用

[1] : https://blog.csdn.net/kelvin_yan/article/details/54019017
[2] : https://www.cnblogs.com/traceorigin/archive/2013/04/11/3015755.html
[3] : https://www.jianshu.com/p/6e2c50c5e0a6
[4] : https://www.cnblogs.com/riddick/p/7892663.html