并行计算之路<2>——CUDA与纹理映射

纹理存储器是GPU常用的优化方法。

举个栗

参考《GPGPU编程技术——从GLSL、CUDA到OpenCL》代码6-2（离散卷积程序）。因为CUDA版本的不同，所以本人对原书的代码进行修改，修改后的代码如下。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>

using namespace std;

// Block 的尺寸为[BLOCK_X * BLOCK_Y]
#define BLOCK_X   16u
#define BLOCK_Y   16u

// 定义一个四维向量相加的宏
#define VectorAdd(a, b) a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w;

// 定义一个四维向量与标量除法的宏
#define VectorScalarDev(a, b, c) a.x = b.x / c; a.y = b.y / c; a.z = b.z / c; a.w = b.w / c;

// 纹理相关全局变量
texture<float4, 2, cudaReadModeElementType> refTex;
cudaArray* cuArray;

__global__ void convolution(int nWidth, int nHeight, int nRadius, float4* pfResult)
{
    // 计算线程号
    const int idxX = blockIdx.x * blockDim.x + threadIdx.x;
    const int idxY = blockIdx.y * blockDim.y + threadIdx.y;

    // 将线程号映射到全局存储器
    const int idxResult = idxY * nHeight + idxX;

    // 近邻元素的数量
    int nTotal = 0;
    // 近邻元素之和
    float4 f4Sum = { 0.0f, 0.0f, 0.0f, 0.0f };
    // 计算结果：输出四元向量
    float4 f4Result = { 0.0f, 0.0f, 0.0f, 0.0f };
    float4 f4Temp = { 0.0f, 0.0f, 0.0f, 0.0f };

    // 近邻元素求和
    for (int ii = idxX - nRadius; ii < idxX + nRadius; ii++)
    {
        for (int jj = idxY - nRadius; jj <= idxY + nRadius; jj++)
        {
            if (ii >= 0 && jj >= 0 && ii < nWidth && jj < nHeight)
            {
                f4Temp = tex2D(refTex, ii, jj);
                VectorAdd(f4Sum, f4Temp);
                nTotal++;
            }
        }
    }
    VectorScalarDev(f4Result, f4Sum, (float)nTotal);
    *(pfResult + idxResult) = f4Result;
}

int main()
{
    unsigned unWidth = 1024u;
    unsigned unHeight = 1024u;
    unsigned unSizeData = unWidth * unHeight;
    unsigned unRadius = 2u;

    // 准备输入数据
    unsigned unData = 0;
    float4* pf4Sampler;
    cudaMallocHost((void **)&pf4Sampler, unSizeData * sizeof(float4));

    for (unsigned i = 0; i < unSizeData; i++)
    {
        pf4Sampler[i].x = (float)(unData++);
        pf4Sampler[i].y = (float)(unData++);
        pf4Sampler[i].z = (float)(unData++);
        pf4Sampler[i].w = (float)(unData++);
    }

    // 准备纹理
    cudaChannelFormatDesc cuDesc = cudaCreateChannelDesc<float4>();
    cudaMallocArray(&cuArray, &cuDesc, unWidth, unHeight);
    cudaMemcpyToArray(cuArray, 0, 0, pf4Sampler, unSizeData * sizeof(float4), cudaMemcpyHostToDevice);

    // 分配全局存储器
    float4* pfResult;
    cudaMalloc((void **)&pfResult, unSizeData * sizeof(float4));

    // 配置线程并调用内核
    dim3 block(BLOCK_X, BLOCK_Y);
    dim3 grid(ceil((float)unWidth / BLOCK_X), ceil((float)unHeight / BLOCK_Y));
    convolution <<< grid, block >>>(unWidth, unHeight, unRadius, pfResult);

    // 结果
    cudaMemcpy(pf4Sampler, pfResult, unSizeData * sizeof(float4), cudaMemcpyDeviceToHost);

    cudaUnbindTexture(refTex);
    cudaFreeHost(pf4Sampler);
    cudaFreeArray(cuArray);
    cudaFree(pfResult);

    return 0;
}

砸开吃

一般使用纹理存储器处理图形（树上是这么说的），至于有什么好处，树上说了一大堆。天书啊( ╯□╰ )看不懂(*￣︶￣)y。不过好歹有个栗子可以砸开吃。

1. 声明纹理参照系 refTex以及CUDA数组cuArray
2. 准备输入数据 pf4Sampler。
3. （cudaArray类型本身并非模板，因此在使用cudaMallocArray() 分配存储器空间时调用） cudaChannelFormatDesc() 设置数组的数据类型。cudaMemcpyToArray() 初始化CUDA数组。
4. 纹理绑定cudaBindTextureToArray()。
5. 纹理拾取tex2D。
6. 数据操作。
7. 解除绑定。
8. 释放CUDA数组。
   经过8道工序就个享受纹理的美味了。

在CUDA的Samples中也提供了一个关于纹理的例子simpleTexture.cu。

参考：
《GPGPU编程技术——从GLSL、CUDA到OpenCL》♥♥♥♥♥
《数字图像处理高级应用——基于MATLAB与CUDA的实现》♥♥♥
《基于CUDA的并行程序设计》♥♥♥
《CUDA专家手册》♥♥♥♥♥
《高性能CUDA应用设计与开发》♥♥♥♥