10.9 二维纹理操作：避免复制

最初引入CUDA时，CUDA内核只能通过纹理读取CUDA数组。应用程序只能通过内存复制将数据写入CUDA数组中。为了使CUDA内核能够将数据写入然后通过纹理读取，应用程序必须将数据写入设备内存中然后执行设备到数组之间的复制操作。后来，两种针对二维纹理的新机制的引入使其不再需要这些步骤。

·二维纹理能够绑定到一块等步长分配的线性设备内存上；
·表面内存加载及存储指令使得CUDA内核能够直接对CUDA数组执行写操作。

设备内存上的三维纹理操作以及三维表面内存的加载和存储操作不支持这两种机制。

对于那些以常规访问方式读取绝大部分或所有纹理内容的应用程序（例如视频编解码）或那些必须在特斯拉架构的硬件上运行的应用程序，数据最好存储在设备内存中。而对于那些在纹理操作时进行随机访问（但为局部访问）的应用程序，数据最好存储在CUDA数组中，并利用表面内存读写指令进行访问。

10.9.1 设备内存上的二维纹理操作

二维设备内存上的纹理操作并不享有“块线性”寻址所带来的益处，即存于纹理缓存的缓存行以水平的跨度拉取纹理元素（texel，也简称为纹素），而不是它们的二维或三维块。但是，除非程序是以随机方式对纹理进行访问，否则，免除将设备内存复制到CUDA数组所带来的优势，要远远超过因丧失块线性寻址而带来的损失。

调用CUDABindTexture2D(), 可以将二维纹理引用与一块设备内存进行绑定。

cudabindTexture2D( NULL, &tex, texDevice, &channelDesc, inWidth, inHeight, texPitch);

上述调用可以将纹理引用绑定到由texDevice/texPitch指定的二维设备内存区间上。基地址与步长必须符合硬件指定的对齐约束。

[1] 基地址必须符合CUDADeviceProp.texturePithAlignment的对齐限制，而步长则需要符合CUDADeviceProp.texturePitchAlignment的对齐限制。[2] 案例程序tex2d_addressing_device.cu与案例程序tex2d_addressing.cu几乎一模一样，只是前者使用的是设备内存保存的纹理数据。这两个程序的设计非常相似，因此我们只需要看它们不同的部分。这里是使用设备指针和步长的二元组代替CUDA数组。

< CUDAArray *texArray = 0;  
>T *texDevice = 0;  
> size_t texPitch;

调用CUDAAllocPitch()而不是CUDAAllocArray()。

CUDAMallocPitch()将委派驱动程序对基地址与步长进行选择，因此该代码仍可以在未来生产的硬件上执行（未来硬件会有增加对齐要求的趋势）。

< CUDART_CHECK(cudaMallocArray( &texArray, < &channelDesc, < inWidth, > CUDART_CHECK(cudaMallocPitch( &texDevice, > &texPitch, >inWidth\*sizeof(T), inHeight));

接着，调用)cudaTextureBind2D()而不是cudaBindTextureToArray()。

< CUDART_CHECK(cudaBindTextureToArray.tex, texArray));  
> CUDART_CHECK(cudaBindTexture2D(NULL, &tex, texDevice, &channelDesc, inWidth, inHeight, texPitch });

最后的差别在于CUDA数组的释放，此处，应对CUDA mallocPitch()返回的指针调用CUDAFree()进行释放。

< CUDAFreeArrayTEXArray>； > CUDAFree( texDevice）;

10.9.2 二维表面内存的读写

与一维表面内存的读写操作一样，费米架构的硬件允许内核直接使用表面内存读写内置函数对CUDA数组进行写操作。

template<class Type> Type surf2Dread(surface<void, 1> surfRef, int x, int y, boundaryMode =EMENTBoundryModeTrap);
template<class Type> Type surf2Dwrite(surface<void, 1> surfRef, Type data, int x, int y, boundaryMode =EMENTBoundryModeTrap);

Surf2Dmemset.cu给出了表面引用的声明以及对应的二维表面内存memset函数的CUDA内核实现。具体代码如下：

surface<void, 2> surf2D;
template<typename T>
    __global__void
surf2Dmemset_kernel{ T value,
        int xOffset, int yOffset,
        int Width, int Height }
{
    for (int row = BlockIdx.y*blockDim.y + threadIdx.y;
        row < Height;
        row += blockDim.y*gridDim.y)
    {
        for (int col = BlockIdx.x*blockDim.x + threadIdx.x;
            col < Width;
            col += blockDim.x*gridDim.x)
        {
            surf2Dwrite(value,
                surf2D,
                (xOffset+col)*sizeof(T),
                yOffset+row);
        }
    }
}

记住，surf2Dwrite()中的X偏移参数是以字节为单位计算的。

[1] CUDA数组也必须遵守相同的限制，但此时基地址与步长将由 CUDA管理，隐藏在内存层次结构中。
[2] 在驱动程序API中，对应设备属性查询为CU_DEVICE_ATTRIBUTE_TEXTURE Alignment 和

CU_DEVICE_ATTRIBUTE_TEXTTONPITCHalignment。