CUDA——纹理内存（纹理参考）

纹理内存（一）

1.1纹理类型

CUDA支持GPU上的一部分纹理硬件（它们原本是为图形处理而设计的）。从纹理内存而非全局内存读取数据。

在内核中 ，使用名为纹理拾取的设备函数读取纹理内存。纹理拾取的第一个参数是一个名为纹理参考的对象。

1.2纹理内存的使用

纹理内存的API分为两种：1.纹理参考 2.纹理对象

(纹理对象相较于纹理参考，为更高级的API。)

1.2.1 纹理参考API

纹理参考定义要拾取哪部分纹理内存。它必须通过宿主运行时函数绑定要一些内存区域（称为纹理），然后 才能供内核使用 。几个不同的纹理参考可以绑定到同一纹理或在内存中有所交叠。

纹理参考的API主要包括纹理参考声明、纹理参考绑定、纹理参考使用和纹理参考解绑。

纹理参考声明

1)纹理参考的一些属性不可变（如:Type,Dim,ReadMode），而且必须在编译时已知（即纹理声明必须为全局变量）；它们在声明纹理参考时被指定，纹理参考在文件范围内被声明为texture类型的变量。

texture<Type,Dim,ReadMode> texRef;
/*
Type指定拾取纹理内存时返回的数据类型；Type可为基本的整数int和浮点数float。
*/

/*
Dim指定纹理参考的维度，等于1或2；Dim是可选参数，缺省值为1.
*/

/*
ReadMode指定读取(访问)纹理内存的模式.ReadMode等于cudaReadModeNormalizedFloat或cudaReadModeElementType；ReadMode是可选参数，默认为cudaReadModeElementType.
1）当ReadMode=cudaReadModeNormalizedFloat时，且Type不为float时，其返回值会全范围归一化映射到[0.0,1.0]区间(对于无符号类型)或[-1.0,1.0]区间(对于有符号)
2）当ReadMode=cudaReadModeElementType时，则不执行任何转换。
*/

2)运行时纹理参考属性
纹理参考的其它属性是可变的，可以在执行时通过宿主运行时(即运行时API)进行更改。它们指定寻址模式、纹理筛选以及纹理坐标是否归一化。

a.纹理坐标
默认情况下，使用[0,N)区间的浮点坐标实现纹理参考(N是与坐标相对应的维度的纹理大小)。例如，大小为64x32的纹理将分别在x和y维度上使用区间[0,63]和[0,31]进行寻址引用。
归一化的纹理坐标将区间映射为[0.0,1.0)区间。如果纹理坐标独立于纹理大小(即当图片的每一个像素点都不要求全索引到的话，则坐标跟纹理大小无关系，例如:resize图片)，则一般选用归一化坐标。

b.寻址模式(数组的领域处理)
寻址模式定义纹理坐标超出范围时要执行的操作。该模式有: clamp模式（默认）、wrap模式、Border模式、Mirror模式
clamp寻址模式：纹理坐标小于0.0设置为0.0，大于1.0的设置为1.0(非归一化时为N-1)。
wrap寻址模式(仅用于归一化模式)：该模式仅使用纹理坐标的小数部分；例如，1.25当做0.25处理，-1.25当做0.75处理。(该模式主要处理当纹理包含周期性信息时)
Border寻址模式:超过范围的纹理坐标所对应的纹理元素用0表示。
Mirror寻址模式: abs(纹理坐标)，超出范围的纹理坐标取其绝对值。

c.纹理筛选
仅为设置为返回浮点数数据的纹理执行线性纹理过滤，即声明时Type应为float。启动线性纹理过滤时，将读取纹理拾取位置周围的纹理元素。并基于纹理坐标落入纹理元素之间的位置插值生成纹理拾取的返回值。对于一维纹理执行简单线性插值，对于二维纹理执行双线性插值。

纹理参考绑定

纹理参考声明后，需要将绑定内存才能使用。线性内存(cudaMalloc)，使用cudaBindTexture()进行绑定;CUDA数组(cudaMallocArray),使用cudaBindTextureToArray()进行绑定。

1)线性纹理内存绑定

/*
	线性内存中分配的纹理:
	1.维度只能等于1；
	2.不支持纹理过滤；
	3.只能使用非归一化的整数纹理坐标寻址；
	4.寻址模式单一：越界的纹理访问将返回0值。
*/

//声明纹理参考
texture<float,cudaTextureType2D,cudaReadModeElementType> texRef;
//声明线性内存
float *d_data;
size_t pitch;
size_t offset;
cudaMallocPitch((void**)&d_data, &pitch, width * sizeof(float), height);//拷贝为cudaMemcpy2D()

//将线性内存绑定为纹理
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaBindTexture2D(&offset, &texRef, d_data, &channelDesc, width, height, pitch);

2)CUDA数组纹理内存绑定

texture<float,cudaTextureType2D,cudaReadModeElementType> texRef;

cudaChannelFormatDesc channelDesc=cudaCreatChannelDesc<float>();
cudaArray *dev_array;
cudaMallocArray(&dev_array,&channelDesc,width,height);//拷贝时为cudaMemcpyToArray()

cudaBindTextureToArray(&texRef,dev_array,&channelDesc);

纹理参考使用

1）从线性内存(设备内存)取纹理
一维使用tex1Dfetch()函数访问纹理内存
二维使用tex2D()

template<class Type>
Type tex1Dfetch(
	texture<Type,1,cudaReadModeElementType> texRef;
	int x);
)
//x为非归一化纹理坐标，为int类型

2)从CUDA数组取纹理
从CUDA数组取纹理时，使用tex1D()或tex2D()。

template<class Type,enum cudaTextureReadMode readMode>
Type tex1D(texture<Type,1,readMode> texRef,float x);

template<class Type,enum cudaTextureReadMode readMode>
Type tex2D(texture<Type,2,readMode> texRef,float x,float y);

纹理参考解绑

使用cudaUnbindTexture()进行解绑。

示例代码（CUDA数组）

#include <cuda.h>

#include <cuda_runtime.h>
#include <device_functions.h>

#include <device_launch_parameters.h>
#include "texture_fetch_functions.h"
#include"cuda_texture_types.h"


#include<iostream>
#include <stdio.h>
#include<memory.h>

using namespace std;

//纹理声明
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;


int iDivUp(int a, int b)
{
	return (a%b != 0) ? (a / b + 1) : (a / b);
}

__device__ size_t flatten_2d_index(size_t idx, size_t idy, size_t width)
{
	return idx + idy * width;
}

__global__ void Texture_test(float *output,size_t width,size_t height)
{
	size_t ix = threadIdx.x + blockIdx.x*blockDim.x;
	size_t iy = threadIdx.y + blockIdx.y*blockDim.y;
	
	if ((ix < width) && (iy < height))
	{
		size_t out_id = flatten_2d_index(ix, iy, width);

		//采用线性纹理过滤时，纹理坐标 为 输入参数-0.5；固这里的参数需要加上0.5f的偏移量(使用线性内存绑定纹理参考时，纹理坐标不用加0.5f？？？？)
		output[out_id] = tex2D<float>(texRef, ix + 0.5f, iy + 0.5f);
	}

}

int main()
{
	//实验数据
	const int width = 16;
	const int height = 8;
	size_t size = sizeof(float)*width*height;

	//float(*h_data)[width] = new float[height][width]();

	float *h_data = new float[width*height]();

	for (int i = 0; i < height; ++i)
	{
		for (int j = 0; j < width; ++j)
		{
			h_data[j + i * width] = rand() % 100;	//产生0-99的随机数
			cout << h_data[j + i * width] << " ";
		}
		cout << endl;
	}
	cout << endl;

	//声明CUDA数组
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();//指定CUDA数组类型；CUDA API的cudaChannelFormatDesc类型不能传指针，传引用形参
	cudaArray* cuArray;
	
	cudaMallocArray(&cuArray, &channelDesc, width, height);//申请CUDA数组空间
	cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);//数据从Host->Device

	//声明运行时纹理参考可变属性
	texRef.addressMode[0] = cudaAddressModeBorder;
	texRef.addressMode[1] = cudaAddressModeBorder;
	texRef.filterMode = cudaFilterModeLinear;
	texRef.normalized = 0;	//坐标不做归一化

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

	//启动内核

	float *d_output;
	cudaMalloc((void**)&d_output, size);

	dim3 block(8, 8);
	dim3 grid(iDivUp(width, block.x), iDivUp(height, block.y));

	Texture_test << <grid, block >> > (d_output, width, height);

	cudaMemcpy(h_data, d_output, size, cudaMemcpyDeviceToHost);
	
	//纹理解绑和内存释放
	cudaUnbindTexture(&texRef);
	cudaFree(d_output);
	cudaFreeArray(cuArray);

	for (int i = 0; i < height; ++i)
	{
		for (int j = 0; j < width; ++j)
		{
			cout << h_data[j + i * width] << " ";
		}
		cout << endl;
	}

	system("pause");

	return 0;
}

参考

https://blog.csdn.net/u012348774/article/details/79138901
https://blog.csdn.net/kelvin_yan/article/details/54019017