CUDA纹理处理与流式工作负载全解析

### CUDA纹理处理与流式工作负载全解析 #### 1. 2D纹理处理：避免复制 早期CUDA引入时，CUDA内核只能通过纹理从CUDA数组读取数据，应用程序只能通过内存复制写入CUDA数组。为了让CUDA内核写入后续通过纹理读取的数据，必须先写入设备内存，再进行设备 - 数组的内存复制。不过，现在针对2D纹理有了两种避免此步骤的机制： - 2D纹理可以绑定到按间距分配的线性设备内存范围。 - 表面加载/存储内建函数使CUDA内核能够直接写入CUDA数组。 需要注意的是，不支持从设备内存进行3D纹理处理以及3D表面加载/存储。对于大多数或全部以规则访问模式读取纹理内容的应用程序（如视频编解码器），或者必须在Tesla级硬件上运行的应用程序，最好将数据保留在设备内存中。而对于在纹理处理时进行随机（但局部）访问的应用程序，最好将数据保留在CUDA数组中，并使用表面读写内建函数。 #### 2. 从设备内存进行2D纹理处理 从2D设备内存进行纹理处理不具备“块线性”寻址的优势，即纹理缓存中的缓存行填充只会拉入水平方向的纹理元素，而非2D或3D块。但除非应用程序对纹理进行随机访问，否则避免从设备内存复制到CUDA数组的好处可能会超过失去块线性寻址的代价。 要将2D纹理引用绑定到设备内存范围，可以调用`cudaBindTexture2D()`： ```c cudaBindTexture2D( NULL, &tex, texDevice, &channelDesc, inWidth, inHeight, texPitch ); ``` 上述调用将纹理引用`tex`绑定到由`texDevice / texPitch`指定的2D设备内存范围。基地址和间距必须符合特定硬件的对齐约束。基地址必须相对于`cudaDeviceProp.textureAlignment`对齐，间距必须相对于`cudaDeviceProp.texturePitchAlignment`对齐。 以下是使用设备内存和CUDA数组的代码对比： | 操作 | 使用CUDA数组 | 使用设备内存 | | ---- | ---- | ---- | | 声明 | `cudaArray *texArray = 0;` | `T *texDevice = 0; size_t texPitch;` | | 分配 | `CUDART_CHECK(cudaMallocArray( &texArray, &channelDesc, inWidth, inHeight));` | `CUDART_CHECK(cudaMallocPitch( &texDevice, &texPitch, inWidth*sizeof(T), inHeight));` | | 绑定 | `CUDART_CHECK(cudaBindTextureToArray(tex, texArray));` | `CUDART_CHECK(cudaBindTexture2D( NULL, &tex, texDevice, &channelDesc, inWidth, inHeight, texPitch ));` | | 释放 | `cudaFreeArray( texArray );` | `cudaFree( texDevice );` | #### 3. 2D表面读写 与1D表面读写类似，Fermi级硬件允许内核使用内建表面读写函数直接写入CUDA数组。相关函数如下： ```c template<class Type> Type surf2Dread(surface<void, 1> surfRef, int x, int y, boundaryMode = cudaBoundaryModeTrap); template<class Type> Type surf2Dwrite(surface<void, 1> surfRef, Type data, int x, int y, boundaryMode = cudaBoundaryModeTrap); ``` 在`surf2Dmemset.cu`中给出的2D表面`memset`的表面引用声明和相应的CUDA内核如下： ```c 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偏移参数是以字节为单位的。 #### 4. 3D纹理处理 从3D纹理读取数据与从2D纹理读取类似，但有更多限制： - 3D纹理的大小限制较小（2048x2048x2048，而不是65536x32768）。 - 没有避免复制的策略，CUDA不支持从设备内存进行3D纹理处理或对3D CUDA数组进行表面加载/存储。 除此之外，差异很明显：内核可以使用`tex3D()`内建函数从3D纹理读取数据，该函数接受3个浮点参数，底层的3D CUDA数组必须通过3D内存复制来填充。支持三线性过滤，根据纹理坐标读取并插值8个纹理元素，精度限制与1D和2D纹理处理相同，为9位。 可以通过调用`cuDeviceGetAttribute()`并传入`CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH`、`CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT`和`CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH`，或者调用`cudaGetDeviceProperties()`并检查`cudaDeviceProp.maxTexture3D`来查询3D纹理的大小限制。 创建3D CUDA数组时，`cudaMalloc3DArray()`函数接受`cudaExtent`结构，而不是宽度和高度参数： ```c cudaError_t cudaMalloc3DArray(struct cudaArray** array, const struct cudaChannelFormatDesc* desc, struct cudaExtent extent, unsigned int flags __dv(0)); ``` `cudaExtent`定义如下： ```c struct cudaExtent { size_t width; size_t height; size_t depth; }; ``` 描述3D内存复制操作非常复杂，CUDA运行时和驱动API都使用结构来指定参数。运行时API使用`cudaMemcpy3DParams`结构，声明如下： ```c struct cudaMemcpy3DParms { struct cudaArray *srcArray; struct cudaPos srcPos; struct cudaPitchedPtr srcPtr; struct cudaArray *dstArray; struct cudaPos dstPos; struct cudaPitchedPtr dstPtr; struct cudaExtent extent; enum cudaMemcpyKind kind; }; ``` 其中大部分结构成员本身也是结构，`extent`给出复制的宽度、高度和深度，`srcPos`和`dstPos`成员是`cudaPos`结构，指定复制源和目标的起始点： ```c struct cudaPos { size_t x; size_t y; size_t z; }; ``` `cudaPitchedPtr`是为3D内存复制添加的结构，用于包含指针/间距元组： ```c struct cudaPitchedPtr { void *ptr; /**< Pointer to allocated memory */ size_t pitch; /**< Pitch of allocated memory in bytes */ size_t xsize; /**< Logical width of allocation in elements */ size_t ysize; /**< Logical height of allocation in elements */ }; ``` 可以使用`make_cudaPitchedPtr`函数创建`cudaPitchedPtr`结构： ```c struct cudaPitchedPtr make_cudaPitchedPtr(void *d, size_t p, size_t xsz, size_t ysz) { struct cudaPitchedPtr s; s.ptr = d; s.pitch = p; s.xsize = xsz; s.ysize = ysz; return s; } ``` SDK中的`simpleTexture3D`示例展示了如何使用CUDA进行3D纹理处理。 #### 5. 分层纹理 分层纹理在图形领域被称为纹理数组，因为它们允许将1D或2D纹理排列成可通过整数索引访问的数组。与普通的2D或3D纹理相比，分层纹理的主要优势是它们支持切片内更大的范围，但使用分层纹理没有性能优势。 分层纹理在内存中的布局与2D或3D纹理不同，如果2D或3D纹理使用为分层纹理优化的布局，性能会不佳。因此，创建CUDA数组时，必须在`cudaMallo