第一章 CUDA入门:GPU架构概述、CUDA编程模型、开发环境搭建

各位同学好,我是你们这门课的老朋友。今天咱们正式开篇,聊聊CUDA入门那些事儿。说实话,我当年刚接触GPU编程时,也是一头雾水——CPU玩得挺溜,一上GPU就懵了。后来踩了不少坑,才慢慢摸清门道。这一章,我就把最基础、最核心的东西掰开揉碎讲给你听。

1.1 GPU架构概述:为什么GPU能“并行”到飞起?

先问个问题:CPU和GPU到底有啥区别?你想想看,CPU就像个全能博士,啥都能干,但一次只能处理一两件事。GPU呢?它更像一支建筑队——每个工人只会搬砖,但几千人一起干,效率惊人。

GPU的核心设计理念就是“大量简单计算单元并行干活”。以NVIDIA的架构为例,从Fermi到Ampere,再到Hopper,核心逻辑没变:

  • SM(流多处理器):GPU的基本调度单元。每个SM里有一堆CUDA核心、共享内存、寄存器等。
  • Warp(线程束):32个线程为一组,一起执行指令。这是硬件调度的最小单位。
  • Global Memory(全局内存):显存,所有线程都能访问,但延迟高(几百个周期)。
  • Shared Memory(共享内存):每个SM内部的高速缓存,延迟低(几个周期),但容量小。

我在项目中遇到过一个问题:有个同事写了个核函数,数据全从全局内存读,结果性能惨不忍睹。后来我把热点数据搬到共享内存里,速度直接翻了5倍。嗯,这里要注意——访存优化是CUDA调优的重中之重。

核心概念速览:

  • GPU适合数据并行任务(图像处理、矩阵运算、深度学习)
  • CPU适合控制密集型任务(分支多、逻辑复杂)
  • Warp divergence(线程束分歧)会严重拖慢性能——同一个warp里别搞太多if-else

下面这张图是我自己画的,帮你理清GPU的层级结构:

GPU(设备) SM 0 Shared Memory (48KB) Register File CUDA Core ×32 CUDA Core ×32 Warp Scheduler SM 1 (结构同SM 0) SM 2 (结构同SM 0) Global Memory(显存) SM(流多处理器) CUDA Core Global Memory

1.2 CUDA编程模型:你写的代码到底怎么跑?

说白了,CUDA编程就是“CPU管控制,GPU管计算”。CPU叫Host(主机),GPU叫Device(设备)。你写的核函数(Kernel)在GPU上跑,剩下的逻辑在CPU上跑。

我习惯把CUDA程序分成三步走:

  1. 数据准备:在CPU上分配内存,把数据从CPU拷贝到GPU
  2. 执行计算:启动核函数,GPU上的成千上万个线程并行干活
  3. 结果回传:把计算结果从GPU拷回CPU,释放资源

来看个最简单的例子——两个数组相加:

// 核函数定义:每个线程处理一个元素
__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    float *d_a, *d_b, *d_c;
    
    // 1. 在GPU上分配显存
    cudaMalloc(&d_a, n * sizeof(float));
    cudaMalloc(&d_b, n * sizeof(float));
    cudaMalloc(&d_c, n * sizeof(float));
    
    // 2. 把数据从CPU拷到GPU(省略了CPU端数据初始化)
    cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, n * sizeof(float), cudaMemcpyHostToDevice);
    
    // 3. 启动核函数:256个线程一个块,一共4个块
    vecAdd<<<4, 256>>>(d_a, d_b, d_c, n);
    
    // 4. 结果拷回CPU
    cudaMemcpy(h_c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);
    
    // 5. 释放显存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

小技巧:我刚开始写CUDA时,经常忘了检查cudaMalloc和cudaMemcpy的返回值。后来养成习惯,每次调用CUDA API后都加一句cudaGetLastError()检查,省了好多调试时间。

这里有个关键概念——线程层次结构:

  • Thread(线程):最小的执行单元
  • Block(线程块):一组线程,共享shared memory,可以同步
  • Grid(网格):一组block,构成整个核函数

你想想看,<<<gridDim, blockDim>>>这个语法,gridDim是网格维度,blockDim是块维度。我一般用一维索引处理数组,二维索引处理图像,三维索引处理体数据。

1.3 开发环境搭建:别让环境卡住你

说实话,环境搭建是很多新手的第一道坎。我当年装CUDA Toolkit时,因为驱动版本不对,折腾了一整天。下面我把步骤拆清楚,你照着来就行。

1.3.1 NVIDIA驱动安装

先确认你的显卡型号。在终端里跑:

lspci | grep -i nvidia   # Linux
wmic path win32_videocontroller get name  # Windows

然后去NVIDIA官网下载对应驱动。我个人建议装Game Ready驱动或者Studio驱动都行,关键是版本别太老。

避坑指南:我曾经在Ubuntu上装了驱动后黑屏,原因是nouveau开源驱动没禁用。记得在/etc/modprobe.d/blacklist-nouveau.conf里加上blacklist nouveau,然后重启。

1.3.2 CUDA Toolkit安装

CUDA Toolkit包含了编译器nvcc、CUDA库、调试工具等。推荐用runfile(本地安装包)方式安装,可控性高:

wget https://developer.download.nvidia.com/compute/cuda/12.3.0/local_installers/cuda_12.3.0_linux.run
sudo sh cuda_12.3.0_linux.run

安装时注意:

  • 不要勾选Driver(如果你已经装好了驱动)
  • 记住安装路径,默认是/usr/local/cuda-12.3
  • 配置环境变量:export PATH=/usr/local/cuda-12.3/bin:$PATH
  • 配置库路径:export LD_LIBRARY_PATH=/usr/local/cuda-12.3/lib64:$LD_LIBRARY_PATH

验证安装:

nvcc --version
# 应该输出类似:Cuda compilation tools, release 12.3, V12.3.107

1.3.3 cuDNN安装

cuDNN是深度神经网络的加速库。如果你做AI相关开发,这个必装。步骤:

  1. 去NVIDIA官网下载cuDNN(需要注册账号)
  2. 解压后把文件拷贝到CUDA目录:
tar -xzvf cudnn-linux-x86_64-8.9.7.29_cuda12-archive.tar.xz
sudo cp cudnn-*-archive/include/cudnn*.h /usr/local/cuda/include
sudo cp -P cudnn-*-archive/lib/libcudnn* /usr/local/cuda/lib64
sudo chmod a+r /usr/local/cuda/include/cudnn*.h /usr/local/cuda/lib64/libcudnn*

验证cuDNN版本:

cat /usr/local/cuda/include/cudnn_version.h | grep CUDNN_MAJOR -A 2

版本兼容性速查表(以CUDA 12.x为例):

组件 推荐版本 备注
NVIDIA驱动 ≥ 545.23 向下兼容
CUDA Toolkit 12.3 LTS版本,稳定
cuDNN 8.9.7 匹配CUDA 12.x

嗯,到这里,你的CUDA开发环境就搭好了。我建议你写个简单的hello.cu跑一下,确认一切正常:

// hello.cu
#include <stdio.h>

__global__ void hello() {
    printf("Hello from thread %d in block %d\n", 
           threadIdx.x, blockIdx.x);
}

int main() {
    hello<<<2, 4>>>();
    cudaDeviceSynchronize();
    return 0;
}

编译运行:

nvcc hello.cu -o hello
./hello

看到每个线程打印自己的ID,说明环境没问题了。记住,环境搭建只是起点,真正的挑战在后面。下一章我们会深入线程模型和内存模型,那才是CUDA性能调优的核心。


想要更多资料,联系可可