📚 课程目录第2章 CUDA编程模型与图像数据映射

2. CUDA编程模型与图像数据映射

在图像处理中，一张图像天然是二维甚至三维的数据结构：宽、高、通道数。而CUDA的线程组织是三维的：grid、block、thread。把图像映射到CUDA线程层次，本质上就是建立"像素坐标 ↔ 线程索引"的对应关系。理解这个映射，是写出高效图像kernel的第一步。

线程层次：从grid到thread的三级结构

CUDA的线程组织分为三级：grid（网格）包含若干block（线程块），每个block包含若干thread（线程）。grid和block都可以是一维、二维或三维的。每个线程通过内置变量唯一标识自己：

threadIdx.x/y/z：线程在block内的索引
blockIdx.x/y/z：block在grid内的索引
blockDim.x/y/z：block的维度大小
gridDim.x/y/z：grid的维度大小

全局线程索引的计算公式为：globalIdx = blockIdx.x * blockDim.x + threadIdx.x。对于二维图像，通常用x维度对应列（宽），y维度对应行（高）。

图像到线程的映射策略

假设处理一张1920×1080的RGB图像，最自然的映射方式是：每个线程处理一个像素。此时可以设置block为16×16（256个线程），grid为(1920/16)×(1080/16) = 120×68（向上取整）。这样每个线程的(threadIdx.x + blockIdx.x*blockDim.x, threadIdx.y + blockIdx.y*blockDim.y)就对应一个像素坐标。

但图像处理中还有几种常见变体：

一像素一线程：适合逐像素点操作（如灰度化、阈值化），映射最直接。
一通道一线程：当通道数较多时，可将z维度用于通道索引，即threadIdx.z对应通道。
多像素一线程：每个线程处理多个像素（如4个），可提高指令级并行和内存访问效率，适合计算密集型滤波。
ROI映射：只处理感兴趣区域时，线程索引需要加上ROI的偏移量，即imgX = roi.x + threadIdx.x + blockIdx.x*blockDim.x。

内存层次与图像数据的存放

CUDA的内存层次直接影响图像处理的性能。从快到慢依次为：

寄存器：每个线程私有，速度最快，适合存放局部变量。
共享内存：block内共享，速度接近寄存器，适合存放图像块（tile），用于卷积、滤波等需要邻域访问的操作。
全局内存：所有线程可访问，容量大但延迟高，图像数据通常存放在这里。
常量内存：只读，适合存放卷积核、查找表等。
纹理内存：针对二维空间局部性优化，适合图像插值和邻域访问。

对于图像处理，一个关键优化是合并访问：相邻线程访问相邻内存地址。当线程按x方向连续排列，且图像按行优先存储时，threadIdx.x对应的像素在内存中也是连续的，这样就能实现合并访问，大幅提升带宽利用率。

kernel启动配置实战

下面是一个灰度化kernel的启动配置示例：

// 图像尺寸：1920x1080，RGB三通道
// 每个线程处理一个像素
dim3 blockSize(16, 16);  // 256线程/block
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
              (height + blockSize.y - 1) / blockSize.y);

grayScaleKernel<<<gridSize, blockSize>>>(d_input, d_output, width, height);

// kernel内部
__global__ void grayScaleKernel(uchar3* input, uchar* output, int w, int h) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < w && y < h) {
        int idx = y * w + x;
        uchar3 pixel = input[idx];
        output[idx] = 0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z;
    }
}

这里有几个关键点：边界检查（if (x < w && y < h)）必不可少，因为grid尺寸是向上取整的，可能超出图像边界；索引计算使用y * w + x，与行优先存储一致；block尺寸选择16×16是常见经验值，兼顾了合并访问和占用率。

对于ROI处理，只需在计算x、y后加上ROI偏移：int imgX = roi.x + x; int imgY = roi.y + y;，然后判断imgX < roi.x + roi.width && imgY < roi.y + roi.height即可。

线程层次与内存层次的对应关系

理解线程层次和内存层次的配合，是优化图像kernel的核心。下图展示了从图像数据到线程映射，再到内存访问的完整关系：

启动配置的常见陷阱与调优

在实际工程中，有几个容易踩坑的地方：

block尺寸不是越大越好：256线程/block通常是甜点区，但具体取决于kernel的寄存器使用量和共享内存需求。可用cudaOccupancyMaxPotentialBlockSize自动计算。
grid尺寸要向上取整：(width + blockSize.x - 1) / blockSize.x，否则会漏掉边缘像素。
二维block的x维度应连续：让threadIdx.x对应图像列，这样相邻线程访问相邻内存，实现合并访问。
共享内存的bank冲突：当使用共享内存做tile时，注意padding避免bank冲突，例如将tile宽度设为blockDim.x + 1。

对于大图像（如4K），grid尺寸可能超过65535的限制（旧架构）。此时需要将grid拆分为多个维度，或在kernel内用循环处理多个像素。现代架构（Compute Capability 3.0+）的gridDim.x可达2^31-1，通常不是问题。

一个完整的ROI处理示例

假设要从1920×1080图像中提取ROI(100, 200, 640, 480)并做高斯模糊，启动配置如下：

// ROI: x=100, y=200, w=640, h=480
dim3 block(16, 16);
dim3 grid((640 + 15) / 16, (480 + 15) / 16);  // 40 × 30

gaussianBlurROIKernel<<<grid, block>>>(d_img, d_out, 1920, 1080, roi);

// kernel内部
__global__ void gaussianBlurROIKernel(uchar* img, uchar* out, int w, int h, ROI roi) {
    int lx = blockIdx.x * blockDim.x + threadIdx.x;
    int ly = blockIdx.y * blockDim.y + threadIdx.y;
    if (lx >= roi.width || ly >= roi.height) return;
    int imgX = roi.x + lx;
    int imgY = roi.y + ly;
    // 高斯模糊：读取3×3邻域
    float sum = 0;
    for (int dy = -1; dy <= 1; dy++)
        for (int dx = -1; dx <= 1; dx++) {
            int nx = imgX + dx, ny = imgY + dy;
            if (nx >= 0 && nx < w && ny >= 0 && ny < h)
                sum += img[ny * w + nx] * kernel[dy+1][dx+1];
        }
    out[imgY * w + imgX] = (uchar)sum;
}

这个例子展示了ROI映射、边界检查、邻域访问的完整流程。注意lx和ly是ROI内的局部坐标，imgX和imgY是原图坐标，两者不能混淆。

小结

图像到CUDA线程的映射，核心是建立"像素坐标 ↔ 线程索引"的对应关系。一像素一线程是最直接的策略，配合二维block和grid，能自然表达图像的二维结构。内存层次的选择则决定了性能上限：全局内存负责数据存放，共享内存负责邻域复用，寄存器负责局部计算。启动配置时，block尺寸、grid取整、边界检查、合并访问是四个必须关注的要点。掌握这些，就能为后续的卷积、滤波、直方图等算法打下坚实的并行化基础。