一、CUDA编程模型概述：GPU架构演进、编程模型与核函数

大家好，我是你们的老朋友。今天咱们正式开始CUDA性能优化的第一课。说实话，每次带新人入门，我都要先问一句：你了解GPU的“脾气”吗？不了解硬件，写出来的核函数就像闭着眼睛开车——能跑，但绝对跑不快。

这一章，我会带你从GPU架构的演进讲起，再到CUDA编程模型的核心概念。嗯，都是基础，但基础不牢，后面优化会踩很多坑。我自己就吃过这个亏。

1.1 GPU架构演进：从图形卡到通用计算利器

先聊聊GPU是怎么“进化”成今天这个样子的。你想想看，最早的GPU就是专门画图的，处理顶点、像素、纹理。那时候没人拿它做科学计算。

转折点出现在2006年。NVIDIA发布了CUDA架构，也就是G80系列。我记得当时在实验室第一次看到CUDA的demo，心里就一个想法：这玩意儿要改变计算格局了。

从架构演进来看，有几个关键节点：

G80（2006年）：统一着色器架构，首次支持CUDA。说白了，就是把顶点和像素着色器合并成通用的流处理器。
Fermi（2010年）：引入真正的L1/L2缓存层次，支持ECC内存。我当年做金融风险计算，Fermi的ECC让我少了很多数据校验的麻烦。
Kepler（2012年）：引入动态并行和Hyper-Q。这个特性让GPU能自己“生”出新的核函数，不用每次都跟CPU打招呼。
Maxwell（2014年）：优化功耗，提升共享内存带宽。做图像处理的朋友应该深有体会。
Pascal（2016年）：统一内存寻址，NVLink高速互联。多卡通信不再是瓶颈。
Volta（2017年）：Tensor Core横空出世。AI训练从此起飞。
Turing（2018年）：RT Core加入，实时光线追踪成为现实。
Ampere（2020年）：第三代Tensor Core，稀疏计算支持。
Hopper（2022年）：Transformer引擎，专门为大模型优化。

你可能会问：了解这些有什么用？我告诉你，知道架构才能知道瓶颈在哪。比如Fermi时代，全局内存访问延迟很高，所以大家拼命用共享内存。到了Maxwell，L2缓存大了很多，有些场景直接用全局内存反而更快。

核心观点：GPU架构的每一次演进，都在解决一个核心矛盾——计算单元越来越多，但数据搬运越来越难。说白了，性能优化的本质就是“让数据离计算更近”。

1.2 CUDA编程模型：主机与设备的协作艺术

CUDA编程模型，说白了就是一套“CPU指挥，GPU干活”的协作框架。CPU叫主机（Host），GPU叫设备（Device）。主机负责逻辑控制、数据准备、任务调度；设备负责大规模并行计算。

我刚开始学CUDA时，犯过一个低级错误：在核函数里调用了printf，结果发现控制台啥也没输出。后来才明白，核函数在设备上执行，它的输出不能直接到主机控制台。嗯，这就是主机与设备的隔离性。

CUDA编程模型的核心要素：

主机代码：在CPU上执行，负责内存分配、数据拷贝、核函数启动。
设备代码：在GPU上执行，即核函数（Kernel），负责并行计算。
内存管理：主机内存和设备内存是分离的，需要显式拷贝（cudaMemcpy）。
执行模型：核函数以网格（Grid）→线程块（Block）→线程（Thread）的层次结构组织。

来看一个最简单的例子：

// 主机代码
int main() {
    int *d_a;
    cudaMalloc(&d_a, N * sizeof(int));  // 设备内存分配
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice); // 数据拷贝到设备
    
    myKernel<<<gridDim, blockDim>>>(d_a);  // 启动核函数
    
    cudaMemcpy(h_a, d_a, N * sizeof(int), cudaMemcpyDeviceToHost); // 结果拷回主机
    cudaFree(d_a);  // 释放设备内存
    return 0;
}

// 设备代码（核函数）
__global__ void myKernel(int *a) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    a[idx] = a[idx] * 2;
}

这个例子虽然简单，但包含了CUDA编程的全部要素。你注意看，主机和设备之间的数据流动是显式的，这也是性能优化的关键点之一。

个人经验：我建议新手先把这个流程刻在脑子里——分配、拷贝、启动、同步、拷回、释放。少一步都可能出问题。我曾经在项目里忘了cudaDeviceSynchronize，结果主机代码先执行完了，设备还在算，数据全是错的。

1.3 核函数概念：GPU上的并行执行单元

核函数（Kernel）是CUDA编程的核心。它用__global__修饰符声明，在设备上执行，由主机调用。一个核函数会被成千上万个线程同时执行，每个线程处理不同的数据。

核函数有几个重要特性：

并行执行：所有线程执行相同的指令，但处理不同的数据（SIMT模型）。
层次结构：线程组织成线程块（Block），线程块组织成网格（Grid）。
内置变量：threadIdx、blockIdx、blockDim、gridDim，用于定位线程身份。
同步机制：__syncthreads()用于线程块内同步。

为什么要有Block这个概念？直接所有线程平铺不行吗？嗯，这里有个硬件限制。GPU的流多处理器（SM）只能同时调度有限数量的线程。Block就是给SM一个“任务包”，SM负责调度包内的线程。Block内的线程可以共享内存、可以同步，这为很多算法提供了便利。

来看一个矩阵加法的核函数：

__global__ void matrixAdd(float *A, float *B, float *C, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < width && col < width) {
        int idx = row * width + col;
        C[idx] = A[idx] + B[idx];
    }
}

这个核函数里，每个线程负责计算一个元素。通过blockIdx和threadIdx的组合，每个线程都能找到自己对应的数据位置。边界检查if (row < width && col < width)是必须的，因为数据可能不是线程数的整数倍。

避坑指南：我曾经在项目里没做边界检查，结果数组越界，读到了脏数据，算出来的结果忽大忽小。排查了整整两天才发现。所以，核函数里的边界检查，千万别省。

1.4 本章知识体系总览

为了让你对本章内容有个整体把握，我画了一张结构图。它展示了GPU架构演进、CUDA编程模型、核函数概念三者之间的关系。

这张图把三者的关系说得很清楚：GPU架构决定了硬件能力，CUDA编程模型提供了软件框架，核函数则是具体的计算实现。三者缺一不可。

本章小结：

GPU架构从G80到Hopper，核心趋势是计算单元越来越多，数据搬运越来越智能。
CUDA编程模型是主机-设备协作模型，内存管理是显式的，需要开发者手动控制。
核函数是GPU上的并行执行单元，通过线程层次结构实现大规模并行。
边界检查、同步机制、内存拷贝是核函数开发中的常见陷阱。

好了，第一章就到这里。内容不多，但都是后面优化的基础。你把这些概念吃透了，后面讲内存优化、线程束调度、指令级并行的时候，才能跟得上节奏。

想要更多资料，联系可可

📚 课程目录

一、CUDA编程模型概述：GPU架构演进、编程模型与核函数

1.1 GPU架构演进：从图形卡到通用计算利器

1.2 CUDA编程模型：主机与设备的协作艺术

1.3 核函数概念：GPU上的并行执行单元

1.4 本章知识体系总览