CUDA实战笔记(12)——二维纹理内存

本系列代码托管于：https://github.com/chintsan-code/cuda_by_example

本篇使用的项目为：heat_2d

对于图形应用，由于通常为二维，因此二维空间内存时非常有用的。下面使用二维纹理内存对之前的热传导应用进行改进。

首先同样是声明纹理引用，这里声明为二维纹理：

texture<float, 2>  texConstSrc;
texture<float, 2>  texIn;
texture<float, 2>  texOut;

这里的参数”2″代表维度，表面这是一个二维纹理引用。

原先使用一维纹理内存时，我们需要使用offset配合tex1Dfetch()从纹理内存中读取数据：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int offset = y * gridDim.x * blockDim.x + x;

int left = offset - 1;
int right = offset + 1;
if (x == 0)
    left++;  // 边缘处理，下同
if (x == DIM - 1)
    right--;

int top = offset - DIM;
int bottom = offset + DIM;
if (y == 0)
    top += DIM;
if (y == DIM - 1)
    bottom -= DIM;

float t, l, c, r, b;
t = tex1Dfetch(texIn, top);     // top
l = tex1Dfetch(texIn, left);    // left
c = tex1Dfetch(texIn, offset);  // center
r = tex1Dfetch(texIn, right);   // right
b = tex1Dfetch(texIn, bottom);  // bottom

在使用二维纹理之后，可以简化我们的代码：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;

float t, l, c, r, b;
t = tex2D(texIn, x, y - 1); // top
l = tex2D(texIn, x - 1, y); // left
c = tex2D(texIn, x, y);     // center
r = tex2D(texIn, x + 1, y); // right
b = tex2D(texIn, x, y + 1); // bottom

二维纹理将简化访问，虽然需要将tex1Dfetch()改为tex2D()，但是不再需要使用线性化变量offset来计算偏移了，而是可以通过x, y来访问纹理。

而且，当使用tex2D()时，我们不需要再担心发生溢出问题了。如果x或y小于0，那么tex2D()将返回对应的0处的值；同理，如果x大于宽度，或y大于高度，也将会返回对应的宽度或高度处的值。

使用二维纹理还有一处需要修改的地方就是纹理绑定：

cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaBindTexture2D(NULL, texConstSrc, data.dev_constSrc, desc, DIM, DIM, sizeof(float) * DIM);
cudaBindTexture2D(NULL, texIn, data.dev_inSrc, desc, DIM, DIM, sizeof(float) * DIM);
cudaBindTexture2D(NULL, texOut, data.dev_outSrc, desc, DIM, DIM, sizeof(float) * DIM);

与使用一维纹理内存不同的是，当将缓冲区绑定到二维纹理时，需要提供一个cudaChannelFormatDesc给CUDA Runtime，这是一个对通道格式描述符(Channel Format Descriptor)的声明。

最后，最后，需要取消对纹理内存的绑定，这部分与一维纹理一致：

cudaUnbindTexture(texConstSrc);
cudaUnbindTexture(texIn);
cudaUnbindTexture(texOut);

无论使用一维纹理还是二维纹理，热传导模拟应用程序的性能基本相同。因此，基于性能来选择一维纹理还是二维纹理可能没有多大的意义。对于这个应用程序来说，当使用二维纹理时，代码可能更简单一些(尤其是可以自动处理边界)，因为模拟的问题刚好是二维的。当通常来说，如果问题不是二维的，那么究竟该选择一维纹理还是二维纹理要视具体情而定。

参考：

• 《GPU高性能编程 CUDA实战》