10.8 二维纹理操作

多数情况下，二维纹理操作与之前介绍的一维纹理操作类似。应用程序可以选择性地将整型值的纹理元素增强为浮点数，其同样可以使用归一化或非归一化的坐标。当支持线性过滤时，可以以纹理坐标小数部分位作为权重，对四个纹理值进行双线性过滤。此外，硬件可以针对每个维度采用不同的寻址模式。例如，X轴可以采用夹取寻址模式，而Y轴采用重叠寻址模式。

程序演示：TEX2D_OPENGL.CU

该程序演示利用图形展示不同纹理操作模式的效果。程序使用了具有可移植性的OpenGL以及GL应用包（GLUT）以减少辅助代码量。为了突出该程序演示的目的，减少其他因素的干扰，程序中将不会使用CUDA与OpenGL的互操作函数。相反，程序将分配映射主机内存，然后使用glDrawPixels()将其渲染到帧缓冲区。对OpenGL而言，该数据就如同来自CPU。

程序支持归一化与非归一化的坐标，并支持分别在X轴方向和Y轴方向执行夹取、重叠、镜像以及边界这四种寻址模式的纹理操作。对非归一化的坐标，以下内核被用来将纹理内容写到输出缓冲区中。

global void RenderTextureUnnormalized( uchar4 *out, int width, int height ) { for ( int row = blockIdx.x; row < height; row += gridDim.x ) { out = (uchar4 *) (((char *) out) + row*4*width); for ( int col = threadIdx.x; col < width; col += blockDim.x ) { out [col] = tex2D (tex2d, (float) col, (float) row); } }

内核主要根据纹理坐标从纹理中的相应位置读出像素值，并将该值相应填充到一个具有width $\times$ height个像素的长方形块中。使用这种方法就能很容易地看出夹取寻址模式和边界寻址模式对超出范围的像素的处理方式。

下面这个内核使用了归一化的坐标将纹理的内容写到输出缓冲区中。

global void RenderTextureNormalized( uchar4 *out, int width, int height, int scale) {
    for (int j = blockIdx.x; j < height; j += gridDim.x) {
        int row = height - j - 1;
        out = (uchar4 *) (((char *) out) + row * 4 * width);
        float texRow = scale * (float) row / (float) height;
        float invWidth = scale / (float) width;
        for (int col = threadIdx.x; col < width; col += blockDim.x) {
            float texCol = col * invWidth;
            out[col] = tex2D.texCol, texRow);
        }
    }
}

参数scale指定了纹理平铺到输出缓冲区中的次数，默认值为1.0，即纹理只会平铺到缓冲区中一次。执行该程序时，我们可以按下键盘的1到9号数字键来选择纹理复制的次数。C、W、M以及B键设置了当前方向的寻址模式，X与Y键指定当前的方向。

读者最好运行该程序，并变换使用不同的参数模式，以观察不同的纹理操作设置所产生的效果。图10-8显示了该程序产生的X方向重叠与镜像以及Y方向重叠与镜像的四个输出，其中纹理复制了5次。

图10-8 重叠与镜像寻址模式