CUDA纹理内存的使用
2013-08-30 12:22
323 查看
首先,纹理存储器是一种只读存储器,如果一个kernel中更改了绑定到纹理存储器的数据,纹理缓存中的数据可能并没有被更新,后来读取的数据就可能是错误的,只能重新启动一个kernel,才会更新纹理缓存。对纹理引用的声明不能写成数组!
最简单的用法:
函数外声明:
texture<float> texConstSrc;
----------函数内-----
cudaBindTexture(NULL, texConstSrc,d_data, imageSize);
tex1Dfetch(texConstSrc, i);
cudaUnbindTexture(texConstSrc);
例子分析:Simpletexture
1.声明CUDA数组之前,必须先用结构体channelDesc描述CUDA数组中的数据类型。
其中,x,y,
z和w分别是每个返回值成员的位数,而f是一个枚举变量,可以取一下几个值:
n cudaChannelFormatKindSigned,如果这些成员是有符号整型;
n cudaChannelFormatKindUnsigned,如果这些成员是无符号整型;
n cudaChannelFormatKindFloat,如果这些成员是浮点型;
举个栗子:一个数据类型为char2型,宽×高×深为64×32×16的CUDA3D数组
cudaChannelFormatDescchannelDesc
= cudaCreateChannelDesc(8, 8, 0,0,cudaChannelFormatKindunsigned);//每个像元由两个char构成
2.
声明纹理参照系
纹理参照系中的一些属性必须在编译时之前被显示声明。因此要声明一个作用范围为全文件的texture型变量:texture<type,dim,
readmode> texRef;
其中,
Type 确定了由纹理拾取返回的数据类型;Type可以是B3.1节中描述的任意一种由基本整型或者单精度浮点型组成能的1-,2-或者4-元组向量类型。
Dim 确定了纹理参照系的维度,默认为1。
ReadMode 可以是cudaReadModeNormalizedFloat或者cudaReadModeElementType。如果ReadMode是cudaReadModeNomalizedFloat,并且Type是16-或者8-bit整型,那么返回的值将是一个浮点数。此时,原来整形的值域会被映射到[0.0,1.0](对无符号整型),或者[-1.0,1.0](对有符号整型)。例如,一个值为0xff的8-bit无符号整型会被映射为1.0f。如果使用cudaReadModeElementType,那么就不会对输出进行转换。ReadMode是一个可选参数,如果不写,那么默认就是cudaReadModeElementType。
例如,下面的代码声明了一个二维,像元数据为unsignedchar型,但将返回值转换为float型的纹理参照系:
texture texRef;
3.设置运行时纹理参照系属性
纹理参照系中的其它属性可以不必声明,并在运行时进行修改。这些参数规定了纹理的寻址模式,是否进行归一化,以及纹理滤波。有C和C++两种风格的接口:
Ø normalized 设置是否对纹理坐标是否进行归一化。如果normalized是一个非零值,那么就会使用归一化到[0,1)的坐标进行寻址,否则对尺寸为width, height, depth的纹理使用坐标[0,width-1], [0,height-1], [0,depth-1]寻址。例如,一个尺寸为64×32的纹理可以通过x维度范围为[0,63],y维度范围[0,31]的坐标寻址。如果采用归一化方式对尺寸为64×32的纹理进行寻址,在x和y维度上的坐标就都是[0.0,1.0)。这样就可以保证纹理的坐标与纹理的尺寸无关。
Ø filterMode用于设置纹理的滤波模式,即如何根据坐标计算返回的纹理值。滤波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。滤波模式为CudaFilterModePoint时,返回值是与坐标最接近的像元的值。CudaFilterModeLinear模式只能对返回值为浮点型的纹理使用,启用这一种模式时将拾取纹理坐标周围的像元,然后根据坐标与这些像元之间的距离进行插值计算。对一维纹理可以使用线性滤波,对二维纹理可以使用双线性滤波。返回值会是对最接近纹理坐标的两个像元(对一维纹理),四个像元(对二维纹理)或者八个像元(对三维纹理)进行插值后得到的值。
Ø addressmode说明了寻址模式,即如何处理超出寻址范围的纹理坐标;addressmode是一个大小为3的数组,三个元素分别说明对第一、二、三个纹理坐标的取址模式;取址模式可以是cudaAddressModeClamp或cudaAddressModeWrap中的一种,前者将超出寻址范围的纹理坐标”钳位”到寻址范围内的最大或最小值,后者将超出寻址范围的纹理坐标“折叠”进合理范围。cudaAddressModeWrap只支持归一化的纹理坐标。
对非归一化的坐标,如果寻址的坐标超过了范围[0,N],大于N的坐标将被钳位,设为N-1。
对归一化的坐标,有钳位和循环两种处理方式,在钳位方式下,超过[0.0,1.0)范围的坐标将被钳位到[0.0,1.0);循环方式一般用于周期循环纹理,它只使用了纹理坐标中有用的小数部分,例如1.25会被当作0.25处理,而-1.25则会被当成0.75处理。
Ø channelDesc描述纹理获取返回值类型,我们已经在3.2.4.1小节讲解CUDAarray时介绍过这个结构体。纹理参照系的返回值类型描述必须和与之绑定的CUDAarray的数据类型描述相同,或者和与之绑定的线性内存中的元素类型相同。
normalized, addressMode和filterMode可以直接在主机端代码中修改。它们只适用于与CUDA数组绑定的纹理参照系。
4.
纹理绑定
如果有CUDA数组想要使用纹理内存,需要用cudaBindTexture()或cudaBindTextureToArray()绑定到纹理上。cudaUnbindTexture()用于解除纹理参照系的绑定。
cudaMalloc((void**)&devPTr,w1*h1*sizeof(float));
以下代码示例绑定一个纹理参照系到devPtr指向的线性内存:
Ø 使用低级API:
Ø 使用高级API
以下代码示例绑定纹理参照系到一个CUDA数组cuArray:
Ø 使用低级API:
Ø 使用高级API
当绑定一个纹理到纹理参照系时,格式必须与声明纹理参照系时的参数匹配;否则,纹理获取的结果是undefined的。
5 纹理拾取
纹理拾取函数采用纹理坐标对纹理存储器进行访问。
对与线性内存绑定的纹理,使用texfetch1D函数访问,采用的纹理坐标是整型。由cudaMallocPitch或者cudaMalloc3D分配的线性空间实际上仍然是经过填充、对齐的一维线性空间,因此也用texfetch1D()函数访问。
对与一维、二维和三维CUDA数组绑定的问哪里,分别使用tex1D()、tex2D()和tex3D()函数访问,并且使用浮点型纹理坐标。
例如: output[y*
width + x] = tex2D(tex, tu, tv);
最简单的用法:
函数外声明:
texture<float> texConstSrc;
----------函数内-----
cudaBindTexture(NULL, texConstSrc,d_data, imageSize);
tex1Dfetch(texConstSrc, i);
cudaUnbindTexture(texConstSrc);
例子分析:Simpletexture
// 2D float texture texture<<span style="color: blue;">float, 2,cudaReadModeElementType> texRef; // Simple transformation kernel __global__ void transformKernel(float*output, int width, int height, float theta) { //根据tid bid计算归一化的拾取坐标 unsigned int x= blockIdx.x * blockDim.x + threadIdx.x; unsigned int y= blockIdx.y * blockDim.y + threadIdx.y; float u= x / (float)width; float v= y / (float)height; // 旋转拾取坐标 u-= 0.5f; v-= 0.5f; float tu= u * cosf(theta) –v * sinf(theta) + 0.5f; float tv= v * cosf(theta) + u * sinf(theta) + 0.5f; //从纹理存储器中拾取数据,并写入显存 output[y* width + x] = tex2D(tex, tu, tv); } // Host code int main() { //分配CUDA数组 cudaChannelFormatDescchannelDesc = cudaCreateChannelDesc(32, 0, 0,0,cudaChannelFormatKindFloat); cudaArray*cuArray; cudaMallocArray(&cuArray,&channelDesc, width, height); //Copy to device memory some data located at addressh_data //in host memory cudaMemcpyToArray(cuArray,0, 0, h_data, size, cudaMemcpyHostToDevice); //Set texture parameters texRef.addressMode[0]= cudaAddressModeWrap; //循环寻址方式 texRef.addressMode[1]= cudaAddressModeWrap; texRef.filterMode=cudaFilterModeLinear; //线性滤波,因为这里是一个图像。如果要保持原来的值则千万不要用线性滤波 texRef.normalized= true;//归一化坐标 //Bind the array to the texture cudaBindTextureToArray(texRef,cuArray, channelDesc); //Allocate result of transformation in devicememory float*output; cudaMalloc((void**)&output,width * height * sizeof(float)); //Invoke kernel dim3dimBlock(16, 16); dim3dimGrid((width + dimBlock.x –1) / dimBlock.x,(height + dimBlock.y–1) / dimBlock.y); transformKernel<<>>(output,width, height,angle); //Free device memory cudaFreeArray(cuArray); cudaFree(output); } |
struct cudaChannelFormatDesc{ int x,y, z, w; enum cudaChannelFormatKindf; }; |
z和w分别是每个返回值成员的位数,而f是一个枚举变量,可以取一下几个值:
n cudaChannelFormatKindSigned,如果这些成员是有符号整型;
n cudaChannelFormatKindUnsigned,如果这些成员是无符号整型;
n cudaChannelFormatKindFloat,如果这些成员是浮点型;
举个栗子:一个数据类型为char2型,宽×高×深为64×32×16的CUDA3D数组
cudaChannelFormatDescchannelDesc
= cudaCreateChannelDesc(8, 8, 0,0,cudaChannelFormatKindunsigned);//每个像元由两个char构成
2.
声明纹理参照系
纹理参照系中的一些属性必须在编译时之前被显示声明。因此要声明一个作用范围为全文件的texture型变量:texture<type,dim,
readmode> texRef;
其中,
Type 确定了由纹理拾取返回的数据类型;Type可以是B3.1节中描述的任意一种由基本整型或者单精度浮点型组成能的1-,2-或者4-元组向量类型。
Dim 确定了纹理参照系的维度,默认为1。
ReadMode 可以是cudaReadModeNormalizedFloat或者cudaReadModeElementType。如果ReadMode是cudaReadModeNomalizedFloat,并且Type是16-或者8-bit整型,那么返回的值将是一个浮点数。此时,原来整形的值域会被映射到[0.0,1.0](对无符号整型),或者[-1.0,1.0](对有符号整型)。例如,一个值为0xff的8-bit无符号整型会被映射为1.0f。如果使用cudaReadModeElementType,那么就不会对输出进行转换。ReadMode是一个可选参数,如果不写,那么默认就是cudaReadModeElementType。
例如,下面的代码声明了一个二维,像元数据为unsignedchar型,但将返回值转换为float型的纹理参照系:
texture texRef;
3.设置运行时纹理参照系属性
纹理参照系中的其它属性可以不必声明,并在运行时进行修改。这些参数规定了纹理的寻址模式,是否进行归一化,以及纹理滤波。有C和C++两种风格的接口:
struct textureReference{ int normalized; enum cudaTextureFilterModefilterMode; enum cudaTextureAddressModeaddressMode[3]; struct cudaChannelFormatDescchannelDesc; } |
Ø filterMode用于设置纹理的滤波模式,即如何根据坐标计算返回的纹理值。滤波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。滤波模式为CudaFilterModePoint时,返回值是与坐标最接近的像元的值。CudaFilterModeLinear模式只能对返回值为浮点型的纹理使用,启用这一种模式时将拾取纹理坐标周围的像元,然后根据坐标与这些像元之间的距离进行插值计算。对一维纹理可以使用线性滤波,对二维纹理可以使用双线性滤波。返回值会是对最接近纹理坐标的两个像元(对一维纹理),四个像元(对二维纹理)或者八个像元(对三维纹理)进行插值后得到的值。
Ø addressmode说明了寻址模式,即如何处理超出寻址范围的纹理坐标;addressmode是一个大小为3的数组,三个元素分别说明对第一、二、三个纹理坐标的取址模式;取址模式可以是cudaAddressModeClamp或cudaAddressModeWrap中的一种,前者将超出寻址范围的纹理坐标”钳位”到寻址范围内的最大或最小值,后者将超出寻址范围的纹理坐标“折叠”进合理范围。cudaAddressModeWrap只支持归一化的纹理坐标。
对非归一化的坐标,如果寻址的坐标超过了范围[0,N],大于N的坐标将被钳位,设为N-1。
对归一化的坐标,有钳位和循环两种处理方式,在钳位方式下,超过[0.0,1.0)范围的坐标将被钳位到[0.0,1.0);循环方式一般用于周期循环纹理,它只使用了纹理坐标中有用的小数部分,例如1.25会被当作0.25处理,而-1.25则会被当成0.75处理。
Ø channelDesc描述纹理获取返回值类型,我们已经在3.2.4.1小节讲解CUDAarray时介绍过这个结构体。纹理参照系的返回值类型描述必须和与之绑定的CUDAarray的数据类型描述相同,或者和与之绑定的线性内存中的元素类型相同。
normalized, addressMode和filterMode可以直接在主机端代码中修改。它们只适用于与CUDA数组绑定的纹理参照系。
4.
纹理绑定
如果有CUDA数组想要使用纹理内存,需要用cudaBindTexture()或cudaBindTextureToArray()绑定到纹理上。cudaUnbindTexture()用于解除纹理参照系的绑定。
cudaMalloc((void**)&devPTr,w1*h1*sizeof(float));
以下代码示例绑定一个纹理参照系到devPtr指向的线性内存:
Ø 使用低级API:
texture<<span style="color: blue;">float, 2,cudaReadModeElementType> texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr,“texRef”); cudaChannelFormatDesc channelDesc =cudaCreateChannelDesc<<spanstyle="color: blue;">float>(); cudaBindTexture2D(0, texRefPtr, devPtr,&channelDesc, width, height, pitch); |
texture<<span style="color: blue;">float, 2,cudaReadModeElementType> texRef; cudaChannelFormatDesc channelDesc =cudaCreateChannelDesc<<spanstyle="color: blue;">float>(); cudaBindTexture2D(0, texRef, devPtr,&channelDesc, width, height, pitch); |
Ø 使用低级API:
texture<<span style="color: blue;">float, 2,cudaReadModeElementType> texRef; textureReference* texRefPtr; cudaGetTextureReference(&texRefPtr,“texRef”); cudaChannelFormatDesc channelDesc; cudaGetChannelDesc(&channelDesc,cuArray); cudaBindTextureToArray(texRef, cuArray,&channelDesc); |
texture<<span style="color: blue;">float, 2,cudaReadModeElementType> texRef; cudaBindTextureToArray(texRef, cuArray); |
5 纹理拾取
纹理拾取函数采用纹理坐标对纹理存储器进行访问。
对与线性内存绑定的纹理,使用texfetch1D函数访问,采用的纹理坐标是整型。由cudaMallocPitch或者cudaMalloc3D分配的线性空间实际上仍然是经过填充、对齐的一维线性空间,因此也用texfetch1D()函数访问。
对与一维、二维和三维CUDA数组绑定的问哪里,分别使用tex1D()、tex2D()和tex3D()函数访问,并且使用浮点型纹理坐标。
例如: output[y*
width + x] = tex2D(tex, tu, tv);
相关文章推荐
- CUDA:纹理内存及其使用
- CUDA 图像处理使用纹理内存与不使用纹理内存对比
- cuda纹理内存使用
- cuda纹理内存的使用
- cuda纹理内存使用例子 vs2013 cuda7.5
- CUDA中多维数组以及多维纹理内存的使用
- CUDA内存--纹理内存的说明与使用
- CUDA使用纹理内存
- CUDA使用纹理内存
- (zz)CUDA中三维纹理的绑定和使用
- CUDA 纹理内存 Error:无法识别texture
- CUDA的内存结构,通过实例展示寄存器和共享内存的使用
- CUDA内存使用
- 基于纹理内存的CUDA热传导模拟
- CUDA By Examples 6 - 使用常量内存 Constant Memory
- CUDA Unified Memory统一内存使用注意
- CUDA学习日志:常量内存和纹理内存
- cuda共享内存,全局内存,纹理等的解释
- CUDA 纹理内存