• CUDA 纹理内存


    原文链接

    1、概述

      纹理存储器中的数据以一维、二维或者三维数组的形式存储在显存中,可以通过缓存加速访问,并且可以声明大小比常数存储器要大的多。

      在kernel中访问纹理存储器的操作称为纹理拾取(texture fetching)。将显存中的数据与纹理参照系关联的操作,称为将数据与纹理绑定(texture binding).

      显存中可以绑定到纹理的数据有两种,分别是普通的线性存储器和cuda数组。

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

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

    2、纹理缓存

      (1)、纹理缓存中的数据可以被重复利用

      (2)、纹理缓存一次预取拾取坐标对应位置附近的几个象元,可以实现滤波模式。

    3、纹理存储器的特殊功能

    4、纹理存储器的使用

      使用纹理存储器时,首先要在主机端声明要绑定到纹理的线性存储器或CUDA数组

      (1)声明纹理参考系

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

        注:纹理参照系必须定义在所有函数体外

      (2) 声明CUDA数组,分配空间

        CUDA数组可以通过cudaMalloc3DArray()或者cudaMallocArray()函数分配。前者可以分配1D、2D、3D的数组,后者一般用于分配2D的CUDA数组。使用完毕,要用              cudaFreeArray()函数释放显存。 

    复制代码
     1 //1数组
     2 cudaMalloc((void**)&dev_A, data_size);
     3 cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice);
     4 cudaFree(dev_A);
     5 
     6 //2维数组
     7 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>()
     8 cudaArray *cuArray;
     9 cudaMallocArray(&cuArray, &channelDesc, 64, 32); //64x32
    10 cudaMemcpyToArray(cuArray, 0, 0, h_data, sizeof(float)*width*height, cudaMemcpyHostToDevice);
    11 cudaFreeArray(cuArray);
    12 
    13 //3维数组 64x32x16
    14 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
    15 cudaArray *d_volumeArray;
    16 cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumSize);
    17 
    18 cudaMemcpy3DParms copyParams = {0};
    19 copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
    20 copyParams.dstArray = d_volumeArray;
    21 copyParams.extent   = volumeSize;
    22 copyParams.kind     = cudaMemcpyHostToDevice;
    23 cudaMemcpy3D(&copyParams);
    24 
    25 tex.normalized = true;
    26 tex.filterMode = cudaFilterModeLinear;
    27 tex.addressMode[0] = cudaAddressModeWrap;
    28 tex.addressMode[1] = cudaAddressModeWrap;
    29 tex.addressMode[2] = cudaAddressModeWrap;
    复制代码

    (3)设置运行时纹理参照系属性

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

      normalized设置是否对纹理坐标归一化

      filterMode用于设置纹理的滤波模式

      addressMode说明了寻址方式

     (4)纹理绑定

      通过cudaBindTexture() 或 cudaBindTextureToArray()将数据与纹理绑定。

      通过cudaUnbindTexture()用于解除纹理参照系的绑定

      注:与纹理绑定的数据的类型必须与声明纹理参照系时的参数匹配

      (I).cudaBindTexture() //将1维线性内存绑定到1维纹理

    复制代码
    1 cudaError_t cudaBindTexture(    
    2     size_t *     offset,
    3     const struct textureReference *     texref,
    4     const void *     devPtr,
    5     const struct cudaChannelFormatDesc *     desc,
    6     size_t     size = UINT_MAX     
    7 )    
    复制代码

    例:

    1 cudaMalloc((void**)&data.dev_inSrc, imageSize);
    2 cudaBindTexture(NULL, tex, data.dev_inSrc, imageSize);

      (II).cudaBindTexture2D //将1维线性内存绑定到2维纹理

    复制代码
    1 cudaError_t cudaBindTexture2D(    
    2     size_t *     offset,
    3     const struct textureReference *     texref,
    4     const void *     devPtr,
    5     const struct cudaChannelFormatDesc *     desc,
    6     size_t     width,
    7     size_t     height,
    8     size_t     pitch     
    9 )    
    复制代码

     例:

    1 cudaMalloc((void**)&data.dev_inSrc, imageSize);
    2 cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    3 cudaBindTexture2D(NULL, tex, data.dev_inSrc, desc, DIM, DIM, sizeof(float)*DIM);

      (III). cudaBindTextureToArray() //将cuda数组绑定到纹理

    1 cudaError_t cudaBindTextureToArray    (    
    2     const struct textureReference *     texref,
    3     const struct cudaArray *     array,
    4     const struct cudaChannelFormatDesc *     desc     
    5 )

    例:

    复制代码
     1 void initCudaTexture(const uchar *h_volume, cudaExtent volumeSize)
     2 {
     3     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
     4 
     5     cutilSafeCall(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize));
     6 
     7     cudaMemcpy3DParms copyParams = {0};
     8     copyParams.srcPtr = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
     9     copyParams.dstArray = d_volumeArray;
    10     copyParams.extent   = volumeSize;
    11     copyParams.kind     = cudaMemcpyHostToDevice;
    12     cutilSafeCall(cudaMemcpy3D(&copyParams));
    13 
    14     tex.normalized = true;
    15     tex.filterMode = cudaFilterModeLinear;
    16     tex.addressMode[0] = cudaAddressModeWrap;
    17     tex.addressMode[1] = cudaAddressModeWrap;
    18     tex.addressMode[2] = cudaAddressModeWrap;
    19 
    20     cutilSafeCall(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));
    21 }
    复制代码

     (5)纹理拾取

      对于线性存储器绑定的纹理,使用tex1Dfetch()访问,采用的纹理坐标是整型。由cudaMallocPitch() 或者 cudaMalloc3D()分配的线性空间实际上仍然是经过填充、对齐的一维线性空间,因此也用tex1Dfetch()

      对与一维、二维、三维cuda数组绑定的纹理,分别使用tex1D(), tex2D() 和 tex3D()函数访问,并且使用浮点型纹理坐标。

  • 相关阅读:
    前端基础进阶(四)-让你一分钟就看懂的作用域和作用域链
    前端基础进阶(三)-史上最详细的变量对象详解
    前端基础进阶(二)-知识点解析最精炼最详细
    前端基础进阶(一):内存空间详解-月薪5万
    知道这20个前端正则表达式,能让你做项目时少写1000行甚至一万行,真的
    学习web前端的免费12个学习网站,等你来撩
    一个老牌程序员推荐的JavaScript的书籍,看了真的不后悔!
    零基础的同学看过来,如何系统学习前端
    这是那些大佬程序员常用的学习java网站,这就是别人薪资上万的原因
    Debug outlook add-in (office.js) 小技巧
  • 原文地址:https://www.cnblogs.com/liangliangdetianxia/p/4198838.html
Copyright © 2020-2023  润新知