10.4_纹理作为数据读取方式
10.4 纹理作为数据读取方式
当使用纹理作为读取数据的方式时，是为了利用硬件纹理单元避免繁锁的合并读取约束或利用纹理缓存，而不是为了利用线性插值等硬件特性。在这一方式下，许多纹理特性是不可用的。此类使用纹理的方式需注意如下要点：
·使用cudaBindTexture()或cuTexRefSetAddress()将纹理引用绑定到设备内存上；
·必须使用tex1Dfetch()指令，其接受一个27位的整型索引；[1]
· tex1Dfetch()具有将纹理内容转换为浮点型数值的可选项。整型值将转换成范围在 之间的浮点值，16位浮点值将增强为标准float值。
使用tex1Dfetch()读取设备内存的好处是双重的。首先，通过纹理读取内存不需要遵循针对读取全局内存的合并读取约束。其次，纹理缓存可以成为其他硬件资源甚至费米架构硬件上的二级缓存的有用补充，当传递给tex1Dfetch()的索引值越界，其返回值为0。
10.4.1 增加有效地址范围
由于27位的索引指定了待读取的纹理元素，而每个纹理元素最大为16字节，因此使用tex1Dfetch()读取的纹理最大能覆盖31位（ ）的内存。一个有效增加纹理覆盖到的数据量的方法是使用比数据实际大小更宽的纹理元素。例如，应用程序可以使用float4的纹理代替float，然后根据需要访问元素索引的最低的若干有效位，从float4中选择相应的元素。对于整型数，也可以使用该技术，尤其是对全局内存中的8位或16位的数据进行访问（因为它们的访问无法合并）。另一种方式，应用程序可以对设备内存中不同分段上的多个纹理设置别名，并对需要的纹理进行断定读取，以此种方式每次只有其中之一处于“活跃”状态。
程序演示：tex1dfetch/big.cu
该程序展示了如何使用tex1Dfetch()从大数组中读取包含多分量的单个纹理或多个纹理。该程序通过以下代码进行调用。
tex1dfetch/big
应用程序分配了指定MB大小的设备内存（若设备内存分配失败，将分配映射锁页主机内存），然后使用随机数对这块内存进行填充，接着使用单分量、双分量或四分量的纹理在这块数据区上计算校验和。每次最多可使用4个元素为int4类型的纹理，即应用程序最大可以在8192MB的内存上进行纹理处理。
为简单清晰起见，tex1dfetch/big.cu并没有执行任何精心设计的并行归约技术。每个线程只是写回一个局部的中间和，最终的校验和计算是在CPU端执行的。以下是程序定义27位的硬件限制的代码。
define CUDA_LG_MAXTEX1DFETCH_INDEX 27
#define CUDA_MAXTEX1DFETCH_INDEX
(( (size_t)1 << CUDA_LG_MAXTEX1DFETCH_INDEX) - 1)另外程序还定义了4个元素类型为int4的纹理。
texture<int4, 1,udaReadModeElementType> tex4_0;
texture<int4, 1,udaReadModeElementType> tex4_1;
texture<int4, 1,udaReadModeElementType> tex4_2;
texture<int4, 1,udaReadModeElementType> tex4_3;设备函数tex4Fetch()接受一个索引值，并将其分离成一个纹理序号和一个27位的索引值，传递给tex1Dfetch()。
device int4 tex4Fetch(size_t index) { int texID = (int) (index>>CUDA_LG_MAXTEX1DFETCH_INDEX); int i = (int) (index & (CUDA_MAXTEX1DFETCH_INDEX_SIZE_T-1)); int4 i4; if (texID == 0) { i4 = tex1Dfetch(tex4_0, i); } else if (texID == 1) { i4 = tex1Dfetch(tex4_1, i); } else if (texID == 2) { i4 = tex1Dfetch(tex4_2, i); } else if (texID == 3) { i4 = tex1Dfetch(tex4_3, i); } return i4; }该设备函数将编译成一小段使用了4个分支断定TEX指令的代码，其中4个纹理只有一个处于活跃状态。如果想要随机访问，应用程序也
可以利用分支断定技术从返回的int4的4个分量.x、.y、.z或.w中选择。
代码清单10-2中所示的绑定纹理操作，用到了少量的技巧。这段代码创建了两个小数组texSizes[]和texBases[]，并设定了它们能够覆盖到的设备内存的整个范围。无论需要映射到设备内存的纹理数目小于4还是等于4，“for循环”保证了所有纹理都进行了有效绑定。
代码清单10-2 texldfetch_large.cu（节选）
int iTexture;
CUDAChannelFormatDesc int4Desc $=$ CUDACreateChannelDesc<int4>();
size_t numInt4s $=$ numBytes / sizeof(int4);
int numTextures $\equiv$ (numInt4s+CUDA_MAX_TX1DFETCH_INDEX) $\gg$ CUDA_LG_MAX_TX1DFETCH_INDEX;
size_t Remainder $=$ numBytes & (CUDA_MAX Bytes_INT4-1);
if(!Remainder){ Remainder $=$ CUDA_MAX Bytes_INT4;
}
size_t texSizes[4];
char\*texBases[4];
for（iTexture $\equiv$ 0；iTexture $<$ numTextures；iTexture++）{ texBases[iTexture] $\equiv$ deviceTex+iTexture\*CUDA_MAX Bytes_INT4; texSizes[iTexture] $\equiv$ CUDA_MAX Bytes_INT4;
} texSizes[iTexture-1] $\equiv$ Remainder;
while（iTexture<4）{ texBases[iTexture] $\equiv$ texBases[iTexture-1]; texSizes[iTexture] $\equiv$ texSizes[iTexture-1]; iTexture++;
}
CUDABindTexture( NULL, tex4_0, texBases[0], int4Desc, texSizes[0]);
CUDABindTexture( NULL, tex4_1, texBases[1], int4Desc, texSizes[1]);
CUDABindTexture( NULL, tex4_2, texBases[2], int4Desc, texSizes[2]);
CUDABindTexture( NULL, tex4_3, texBases[3], int4Desc, texSizes[3]);当编译并运行之后，可以设定不同大小，调用程序以观察效果。在亚马逊EC2云计算环境的CG1实体上运行，分别以512M、768M、1280M以及8192M大小调用内核时的运行结果如下：
$ ./texldfetch_large 512
Expected checksum: 0x7b7c8cd3
tex1 checksum: 0x7b7c8cd3
tex2 checksum: 0x7b7c8cd3
tex4 checksum: 0x7b7c8cd3
$ ./texldfetch_large 768
Expected checksum: 0x559a1431
tex1 checksum: (not performed)
tex2 checksum: 0x559a1431
tex4 checksum: 0x559a1431
$ ./texldfetch_large 1280
Expected checksum: 0x66a4f9d9
tex1 checksum: (not performed)
tex2 checksum: (not performed)
tex4 checksum: 0x66a4f9d9
$ ./texldfetch_large 8192
Device alloc of 8192 Mb failed, trying mapped host memory
Expected checksum: 0xf049c607
tex1 checksum: (not performed)
tex2 checksum: (not performed)
tex4 checksum: 0xf049c607每个int4类型的纹理“只能”读取2GB的内存，因此，当启动程序时，大小超过8192MB时将失败。该应用程序强调了对带索引纹理的需要，此时待读取纹理可以指定为运行时的参数，然而，CUDA并没有提供该特性的支持。
10.4.2 主机内存纹理操作
此外，若将纹理作为一种读取方式使用，应用程序也可以通过分配映射锁页内存，获取设备指针，然后把该设备指针传给CUDABindAddress()或cuTexRefSetAddress()的方式从主机内存读取数据。虽然这种方式可行，但以纹理的形式读取主机内存的速度是很慢的。特斯拉架构的硬件设备可以以每秒2GB的速度从PCIe总线上获取纹理数据，费米架构的硬件设备的速度则要相对慢很多。因此，若要采用这种方式，最好有一定理由，例如，保持代码的简洁性。
程序演示：tex1dfetch_int2float.cu
这段代码将纹理作为一种读取方式使用，从主机内存获取数据，以确定TexPromoteToFloat()函数是否正常工作。即将使用的这个CUDA内核功能很简单，其实现了一个相当于无阻塞能力的memcpy函数，即从纹理中获取数据并写入到设备内存。
texture < signed char, 1, CUDAReadModeNormalizedFloat> tex;
extern "C" __global__ void
TexReadout(float *out, size_t N) {
for (size_t i = BlockIdx.x*blockDim.x + threadIdx.x; i < N; i += gridDim.x*blockDim.x)
out[i] = tex1Dfetch(tex, i);
}由于将整型值增强至浮点型数值只对8位及16位的整型值有效，我们可以通过分配一小块缓冲区，从这块区域获取纹理数据，并判定输出是否我们的预期，以此测试每个转换是否正确。代码清单10-3显示了从tex1dfetch_int2float.cu中节选的部分代码。该代码分配了两块主机内存，其中inHost保存了输入缓冲区中的256或65536个输入值，fOutHost保存了对应的浮点数输出。对应这些映射锁页主机内存的设备指针分别保存到了inDevice和foutDevice中。
代码清单10-3 texld_int2float.cu（节选）
template<class T>
void
CheckTexPromoteToFloat(size_t N)
{ T *inHost, *inDevice; float *fouthost, *foutDevice; CUDAError_t status; CUDART_CHECK(cudaHostAlloc( (void **) &inHost, N*sizeof(T), CUDAHostAllocated); CUDART_CHECK(cudaHostGetDevicePointer( (void **) &inDevice, inHost, 0)); CUDART_CHECK(cudaHostAlloc( (void **) &fouthost, N*sizeof(float), CUDAHostAllocated); CUDART_CHECK(cudaHostGetDevicePointer( (void **) &fouthost, foutHost, 0)); for ( int i = 0; i < N; i++) { inHost[i] = (T)i; } memset( fouthost, 0, N*sizeof(float) ); CUDART_CHECK(cudaBindTexture(NULL, tex, inDevice, CUDACreateChannelDesc<T>( ), N*sizeof(T) ); TexReadout<<2,384>>>(fouthost, N); CUDART_CHECK(cudaDeviceSynchronize()); for ( int i = 0; i < N; i++) { printf( "%.2f", fouthost[i]); assert( fouthost[i] == TexPromoteToFloat( (T)i ) ); } printf( "\n"); Error: CUDAFreeHost( inHost ); CUDAFreeHost( fouthost ); }输入值将被初始化为当前测试类型的任意值，然后调用 CUDABindTexture()将纹理引用与输入对应的设备指针进行绑定。接着启动TexReadOut()内核，读取输入纹理的每一个值，并将 tex1Dfetch()的返回值作为输出写到设备内存。这样，输入缓冲区与输出缓冲区均驻存在映射主机内存上了。由于内核是直接将结果写回到主机内存，因此需调用CUDADeviceSynchronize()以确保CPU与GPU之间不存在竞争条件。函数末尾调用了TexPromoteToFloat()，将当前测
试类型的值转换成浮点值，并与对应的内核返回值进行比较，看是否相等。若所有的比较都相等，函数则会返回true，若有任何API调用失败或比较不相等，则返回false。
[1] 所有支持CUDA的硬件均是27位的限制，因此目前没有查询当前设备上该限制的方法。