Cuda三维纹理的使用

1、代码说明

1.1 代码 

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#ifndef __CUDACC__
#define __CUDACC__
#endif // !__CUDACC__
#include <cuda_texture_types.h>

texture<short, cudaTextureType3D, cudaReadModeNormalizedFloat>	my_texture;
int main()
{
    int width = 128;
    int height = 128;
    int depth = 128;
    short* buffer = new short[width * height * depth];
    cudaArray* cuda_array = nullptr;
    cudaExtent extent = make_cudaExtent(width, height, depth);
    cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc<short>();
    cudaMalloc3DArray(&cuda_array, &channel_desc, extent);

    // 将数据拷贝到纹理数组
    cudaMemcpy3DParms copy_params = { 0 };
    copy_params.srcPtr = make_cudaPitchedPtr(buffer, width * sizeof(short), width, height);
    copy_params.dstArray = cuda_array;
    copy_params.extent = extent;
    copy_params.kind = cudaMemcpyHostToDevice;
    cudaMemcpy3D(&copy_params);

    my_texture.normalized = true;
    my_texture.filterMode = cudaFilterModeLinear;
    my_texture.addressMode[0] = cudaAddressModeClamp;
    my_texture.addressMode[1] = cudaAddressModeClamp;
    my_texture.addressMode[2] = cudaAddressModeClamp;

    // 绑定纹理对象到纹理数组
    cudaBindTextureToArray(&my_texture, cuda_array, &channel_desc);

    cudaDeviceSynchronize();

    return 0;
}

1.2 说明 

        纹理内存实质上是全局内存的一个特殊形态，支持自动内存对齐和边界处理、具备硬件插值特性，可以实现最近邻插值和线性插值，而且还支持归一化功能，将数据范围映射到[0, 1]之间，方便进行数据处理和映射操作。纹理缓存针对二维空间的局部性访问进行了优化，所以通过纹理缓存访问二维矩阵的邻域会获得加速。

        下面是使用三维纹理的基本步骤：

1.2.1 定义 Texture 对象并分配内存：

        1、使用cudaExtent创建长宽高一个三维参数

        2、使用cudaChannelFormatDesc设置数据通道类型

               定义：__host__ struct cudaChannelFormatDesc CUDARTAPI         cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f);

                cudaCreateChannelDesc<short>() = cudaCreateChannelDesc(16, 0, 0, 0,         cudaChannelFormatKindSigned)

                x,y,z,w为通道分量，f为数据类型。

        3、使用cudaMalloc3DArray分配三维数组类型

1.2.2 将数据拷贝到 Texture 对象中：

        1、使用cudaMemcpy3DParms定义数据传输的结构体

        2、使用make_cudaPitchedPtr创建一个带有封装偏移量的结构体

                定义：cudaPitchedPtr make_cudaPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz); 

      d：指向内存的指针。

      p：每行的字节数。

      xsz：内存块在x方向上的大小。

                ysz：内存块在y方向上的大小。

        3、使用cudaMemcpy3D根据以上设置的参数实行主机到设备内存拷贝

3. 定义并绑定纹理：

        1、声明纹理texture<short, cudaTextureType3D, cudaReadModeNormalizedFloat>

                定义：texture<Type, Dim, ReadMode> VarName;
                Type：基本的整型和浮点类型，以及其它的对齐类型
                Dim： 纹理数组的维度，值为1或2或3，默认缺省为1
                ReadMode：cudaReadModelNormalizedFloat 或 cudaReadModelElementType(默认)

                cudaReadModelNormalizedFloat：读取数据时会自动将整型数据归一化为浮点数。具体地，如果是无符号整型，则转化为[0 1]之间的浮点数；如果是有符号整型，则转化为[-1 1]之间的浮点数。
                cudaReadModelElementType：默认值，不进行任何转换

        2、设置参数

                1、normalized：是否归一化

                2、filterMode：滤波模式

                       cudaFilterModePoint：点模式，返回最接近的一个点，即最近邻插值

                       cudaFilterModeLinear：线性模式，即线性插值

                3、addressMode：寻址模式

                        cudaAddressModeClamp：超出范围就用边界值代替

                        cudaAddressModeBorder：超出范围就用零代替

                        cudaAddressModeWrap：重叠模式（循环） 归一化使用

                        cudaAddressModeMirror：镜像模式 归一化使用

        3、使用cudaBindTextureToArray绑定纹理对象到纹理数组