4、GPU并行计算：GPU架构简介、CUDA/OpenCL基础、利用GPU加速渲染任务

好，咱们进入第四讲。说实话，每次讲到GPU架构，我都会想起自己第一次用CUDA写像素处理程序时的场景——那时候我天真地以为，只要把循环扔给GPU就能快100倍。结果呢？跑出来比CPU还慢。嗯，这里面的门道，咱们今天好好聊聊。

4.1 GPU架构简介：它为什么这么快？

先说说GPU和CPU的根本区别。CPU是「精打细算型」——它只有几个核心，但每个核心都特别强，能处理复杂的逻辑分支、乱序执行、大缓存。GPU呢？它是「人海战术型」——成百上千个简单核心，每个核心能力一般，但架不住人多。

我打个比方。CPU就像几个顶尖外科医生，能做高难度手术。GPU就像一千个护士，每个人只会打针，但一千个人同时打针，效率就上来了。渲染任务说白了就是大量像素的重复计算，太适合GPU了。

关键概念：SIMT（单指令多线程）

GPU的核心执行模式。一条指令，多个线程同时执行。但注意，如果线程之间走了不同的分支（if-else），性能会急剧下降。我在项目中遇到过这种情况，一个简单的条件判断让帧率直接腰斩。

现代GPU的架构大致分三层：

  • 流式多处理器（SM）：GPU的基本计算单元，每个SM包含多个CUDA核心
  • CUDA核心：最基础的执行单元，负责整数和浮点运算
  • 共享内存：每个SM内部的高速缓存，比全局内存快两个数量级

你想想看，一个RTX 4090有128个SM，每个SM有128个CUDA核心，总共16384个核心。这什么概念？你CPU才16个核心，差了1000倍。但别高兴太早——这些核心不能独立工作，它们必须集体行动。

4.2 CUDA/OpenCL基础：怎么写GPU程序？

我个人习惯用CUDA，毕竟NVIDIA的生态最成熟。OpenCL虽然跨平台，但说实话，用起来没那么顺手。咱们以CUDA为例，看看GPU程序长什么样。

先看一个最简单的例子——把两个数组相加：

// CPU版本
void add_cpu(float* a, float* b, float* c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

// GPU版本（CUDA kernel）
__global__ void add_gpu(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

看到区别了吗？CPU版本用for循环遍历，GPU版本用线程索引直接定位。每个线程处理一个元素，成千上万个线程同时干活。

我的经验之谈：写CUDA kernel时，一定要检查线程索引是否越界。我曾经因为少写了一个if判断，程序跑起来直接黑屏，查了半天才发现是内存越界把显存写坏了。

调用GPU kernel的代码长这样：

// 分配显存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, n * sizeof(float));
cudaMalloc(&d_b, n * sizeof(float));
cudaMalloc(&d_c, n * sizeof(float));

// 拷贝数据到显存
cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, n * sizeof(float), cudaMemcpyHostToDevice);

// 启动kernel
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
add_gpu<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

// 拷贝结果回内存
cudaMemcpy(h_c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);

这里有个关键点：数据传输。PCIe带宽再快，也比不上显存带宽。所以我的原则是：尽量少传数据，尽量在GPU上完成所有计算。

4.3 利用GPU加速渲染任务

好了，理论说完了，咱们来点实际的。渲染任务中，哪些适合GPU加速？说白了，只要是对每个像素独立操作的任务，都适合。

渲染任务 GPU加速效果 我的建议
颜色校正（亮度/对比度） 极好（10-50倍） 无脑上GPU
模糊/锐化 好（5-20倍） 注意边界处理
图像拼接 中等（2-5倍） 内存访问模式要优化
光流计算 好（10-30倍） 共享内存是关键

举个具体的例子——多屏拼接中的颜色校正。假设你有4个屏幕，每个屏幕的亮度和色温都不一样，需要统一校正。CPU做法：遍历每个像素，计算校正值。GPU做法：每个线程处理一个像素，同时计算。

避坑指南：我曾经在做一个8K多屏渲染项目时，发现GPU利用率只有30%。查了半天，原来是kernel启动次数太多，每次启动都有开销。解决方案：把多个小kernel合并成一个，减少启动次数。帧率直接从15fps飙到55fps。

再聊聊共享内存的使用。做图像处理时，经常需要访问像素的邻居（比如模糊操作）。如果每次都从全局内存读，速度很慢。正确做法：先把数据块加载到共享内存，然后从共享内存读取邻居像素。

__global__ void blur_gpu(unsigned char* input, unsigned char* output, 
                         int width, int height) {
    // 声明共享内存，每个block共享
    __shared__ unsigned char shared[BLOCK_SIZE + 2][BLOCK_SIZE + 2];
    
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    // 加载数据到共享内存（包括边界）
    // ... 这里省略边界加载代码
    
    __syncthreads();  // 同步，确保所有线程都加载完毕
    
    // 从共享内存读取邻居像素
    if (x < width && y < height) {
        float sum = 0;
        for (int dy = -1; dy <= 1; dy++) {
            for (int dx = -1; dx <= 1; dx++) {
                sum += shared[threadIdx.y + 1 + dy][threadIdx.x + 1 + dx];
            }
        }
        output[y * width + x] = sum / 9;
    }
}

看到__syncthreads()了吗？这行代码特别重要。它保证所有线程都加载完数据后，才开始计算。没有它，有的线程可能还在读数据，有的已经开始算了，结果全乱套。

性能调优三原则：

  1. 最大化并行度：让GPU始终有活干，别让核心闲着
  2. 最小化数据传输：能算的别传，能一次传的别分多次
  3. 优化内存访问：尽量用共享内存，全局内存要连续访问

最后说一句，GPU编程不是银弹。有些任务（比如大量分支判断、小数据量计算）在GPU上反而更慢。我的经验是：先分析任务特性，再决定用CPU还是GPU。别一听GPU就往上冲，那跟用牛刀杀鸡没啥区别。

下一讲咱们聊聊多屏同步的具体实现，包括帧缓冲管理和垂直同步的坑。到时候我会分享一个我在展会现场遇到的惨痛教训——4块屏幕不同步，画面撕裂得像抽象画。嗯，那故事可有意思了。

← 上一章 📖 返回目录 下一章 →