📚 课程目录第4章第一个CUDA图像处理程序

4. 第一个CUDA图像处理程序

灰度化与颜色空间转换是图像处理中最基础、也最适合作为第一个CUDA实战项目的任务。它计算规则、访存模式清晰，同时又能完整暴露GPU编程中的核心问题：线程如何映射到像素、内存如何组织、CPU与GPU之间如何传输数据、以及性能瓶颈究竟出现在哪里。本节以一张1920×1080的RGB图像为例，实现RGB转灰度与BGR转RGB两个核函数，并用Nsight Systems观察传输与计算的时间分布。

一、任务定义与数据布局

图像在内存中通常以行优先的连续数组存储。对三通道图像，每个像素占3个字节，排列顺序为R、G、B（或B、G、R）。设图像宽为W、高为H，则第y行第x列像素的起始下标为：

idx = (y * W + x) * 3
R = data[idx + 0]
G = data[idx + 1]
B = data[idx + 2]

灰度化采用ITU-R BT.601亮度公式，这是工程中最常用的加权方式：

Gray = 0.299 * R + 0.587 * G + 0.114 * B

三个权重之和为1，保证灰度值仍在0～255范围内。BGR转RGB则是通道交换：把每个像素的0号与2号字节互换，1号字节保持不变。这两个操作都是逐像素独立计算，天然适合并行。

二、CPU基准实现

先写一个CPU版本作为性能基准。它用单层循环遍历所有像素，逐像素计算灰度并写入输出缓冲区：

void rgb2gray_cpu(const unsigned char* in, unsigned char* out, int pixels) {
    for (int i = 0; i < pixels; ++i) {
        int r = in[i * 3 + 0];
        int g = in[i * 3 + 1];
        int b = in[i * 3 + 2];
        out[i] = (unsigned char)(0.299f * r + 0.587f * g + 0.114f * b);
    }
}

对1920×1080图像，pixels = 2073600。CPU版本在单线程下通常需要数毫秒到十几毫秒，具体取决于编译优化与CPU型号。这个数字就是后续对比的参照。

三、GPU核函数与线程映射

GPU版本的核心是把“每个像素”映射到“每个线程”。最直接的方式是一维网格：线程总数覆盖像素总数，每个线程处理一个像素。核函数如下：

__global__ void rgb2gray_kernel(const uchar3* in, unsigned char* out, int pixels) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < pixels) {
        uchar3 p = in[i];
        out[i] = (unsigned char)(0.299f * p.x + 0.587f * p.y + 0.114f * p.z);
    }
}

这里用uchar3结构体表示一个像素，它占3字节，但实际对齐到4字节。这种“结构体数组”（AoS）布局在GPU上并不理想：每个线程读取3个连续字节，相邻线程的访问地址间隔为3，导致合并访存效率下降。更优的做法是“数组结构体”（SoA）或平面布局，即把R、G、B分别存成三个连续数组，这样每个线程读取的地址是连续的，能获得更好的内存合并。

启动配置上，每个block取256个线程是常见选择。对2073600个像素，需要的block数为：

int threads = 256;
int blocks = (pixels + threads - 1) / threads;  // 8100
rgb2gray_kernel<<<blocks, threads>>>(d_in, d_out, pixels);

BGR转RGB的核函数结构相同，只是把输出像素的0号与2号字节交换。由于两个操作都是逐像素的，也可以合并成一个核函数，一次读取同时完成灰度化和通道交换，减少一次全局内存往返。

四、主机端内存管理与传输

GPU不能直接访问主机内存，必须显式分配设备内存并拷贝数据。典型流程是：

size_t bytes_in  = pixels * 3;
size_t bytes_out = pixels;
unsigned char *d_in, *d_out;
cudaMalloc(&d_in,  bytes_in);
cudaMalloc(&d_out, bytes_out);
cudaMemcpy(d_in, h_in, bytes_in, cudaMemcpyHostToDevice);
rgb2gray_kernel<<<blocks, threads>>>(d_in, d_out, pixels);
cudaMemcpy(h_out, d_out, bytes_out, cudaMemcpyDeviceToHost);

这里有一个容易被忽视的细节：cudaMemcpy是同步的，它会阻塞主机直到拷贝完成。如果按“拷贝—计算—拷贝”的顺序写，GPU在拷贝期间是空闲的，计算期间PCIe总线又是空闲的。对于单张图像，这种串行方式简单可靠；但如果要处理视频流，就应该用流（stream）和异步拷贝把传输与计算重叠起来。

五、用Nsight Systems定位瓶颈

Nsight Systems的作用是把时间轴摊开，让你看到每一段时间到底花在哪里。采集命令如下：

nsys profile --stats=true -o gray_profile ./gray_demo

采集完成后，在时间轴上通常能看到三段明显的区间：HtoD拷贝、核函数执行、DtoH拷贝。对1920×1080图像，输入约5.9 MB，输出约2.0 MB。如果PCIe带宽按6 GB/s估算，HtoD约1 ms，DtoH约0.3 ms；而灰度核函数本身在GPU上可能只需要0.1～0.3 ms。也就是说，传输时间远大于计算时间，瓶颈在PCIe，而不是在核函数。

这个结论非常重要：它意味着继续优化核函数（比如调block大小、改SoA布局）对总时间的改善有限，真正有效的方向是减少传输量或隐藏传输延迟。常见手段包括：把灰度化与通道交换合并，避免中间结果回传；使用cudaMallocHost分配页锁定内存提高拷贝带宽；用多个CUDA流把不同图像的传输与计算重叠。

六、结果验证与性能对比

验证GPU结果是否正确，最直接的方法是把GPU输出拷回主机，与CPU输出逐字节比较。由于浮点运算顺序不同，灰度值可能有±1的舍入差异，因此比较时应允许1的容差：

int max_diff = 0;
for (int i = 0; i < pixels; ++i)
    max_diff = max(max_diff, abs(h_cpu[i] - h_gpu[i]));
printf("max_diff = %d\n", max_diff);  // 期望 0 或 1

性能对比时要注意计时口径。CPU版本用std::chrono包住整个循环；GPU版本则应分别记录“仅核函数”和“含传输”两个时间。前者反映计算能力，后者反映端到端真实开销。对1920×1080图像，一个合理的示例结果是：CPU单线程约8 ms，GPU核函数约0.2 ms，但含传输的端到端时间约1.5 ms。GPU核函数比CPU快约40倍，但端到端只快约5倍——差距全部来自PCIe传输。

判断优化是否有效的标准不是“核函数快了多少”，而是“端到端时间降了多少”。当传输占比超过50%时，优先优化传输；当计算占比超过70%时，才值得深入调优核函数。

七、从单图到批处理的自然延伸

单张图像的传输开销无法被计算掩盖，但批处理可以。假设一次处理100张图像，可以把它们拼成一个大缓冲区一次性传输，或者用多个CUDA流让第i张的传输与第i-1张的计算重叠。Nsight Systems的时间轴上会看到拷贝段与计算段交错排列，GPU利用率显著上升。这也是后续章节中视频处理和实时渲染场景的基本优化模式。

本节的两个核函数虽然简单，却完整覆盖了CUDA编程的基本要素：线程索引、边界判断、内存分配、同步拷贝、结果验证和性能分析。把这条链路走通，后续更复杂的卷积、滤波、直方图等算法，都只是在核函数内部增加计算量，而主机端的管理框架和性能分析方法是一致的。