第一章：CUDA编程基础

各位同学，今天我们来聊聊CUDA编程的入门内容。说实话，我刚开始接触GPU编程时，也被那一堆概念搞得头晕。但别担心，咱们一步步来。

1.1 CUDA编程模型：Host与Device

CUDA的编程模型，说白了就是「CPU+GPU」协同工作。CPU这边叫Host，GPU那边叫Device。我习惯把Host想象成项目经理，Device就是干活的主力团队。

Host负责逻辑控制、串行任务，比如数据准备、任务调度。Device负责大规模并行计算，比如矩阵乘法、图像处理。两者通过PCIe总线通信，但要注意——数据传输是有开销的。

核心要点：Host端代码在CPU上运行，Device端代码在GPU上运行。用__global__修饰的函数叫kernel，它在Device上执行，但由Host调用。

我在项目中遇到过一个问题：有个同事把大量数据频繁在Host和Device之间拷贝，结果性能还不如纯CPU版本。嗯，这里要注意——尽量减少数据传输，能一次传完就别分多次。

1.2 线程层次结构：Grid、Block、Thread

GPU的并行执行单元分三层：Grid（网格）、Block（线程块）、Thread（线程）。你想想看，这就像军队编制：Grid是集团军，Block是师，Thread是士兵。

Thread：最小的执行单元，每个线程执行相同的kernel代码，但处理不同的数据
Block：一组线程的集合，同一Block内的线程可以共享数据（通过Shared Memory）
Grid：一组Block的集合，一个kernel对应一个Grid

为什么要有这种层次结构？因为GPU硬件有物理限制。每个SM（流多处理器）能承载的线程数和Block数都是有限的。我刚开始写CUDA时，总想把所有线程塞到一个Block里，结果发现性能反而下降——因为资源不够用了。

个人经验：Block大小一般设为128或256的倍数，这样能充分利用SM的warp调度能力。我习惯用256，兼容性最好。

线程的索引方式也很重要。CUDA提供了内置变量：threadIdx、blockIdx、blockDim、gridDim。举个例子：

// 获取全局线程ID
int tid = blockIdx.x * blockDim.x + threadIdx.x;

这段代码我写了不下千遍，但每次写都要确认一下维度对不对。你想想看，如果维度搞错了，数据就全乱套了。

1.3 内存层次结构：Global、Shared、Register

GPU的内存层次，是性能调优的关键。我见过太多人只盯着计算优化，却忽略了内存访问模式。其实，很多时候瓶颈在内存，不在计算。

内存类型	位置	访问速度	生命周期	作用域
Global Memory	显存	慢（~400周期）	程序运行期间	所有线程
Shared Memory	芯片内	快（~5周期）	Block生命周期	同一Block内
Register	寄存器	极快（~1周期）	线程生命周期	单个线程

Global Memory是最大的，但也是最慢的。Shared Memory是片上的，速度快得多，但容量有限（通常几十KB）。Register最快，但每个线程能用的数量有限（一般255个左右）。

避坑指南：我曾经在项目里过度使用Shared Memory，导致每个Block能承载的线程数大幅减少，最终性能反而下降。记住：Shared Memory是稀缺资源，要精打细算。

还有一个重要的概念——内存合并访问。当多个线程访问连续的Global Memory地址时，硬件会自动合并成一次大的内存事务。我建议你尽量让线程访问连续的数据，这样能充分利用内存带宽。

1.4 第一个CUDA程序：向量加法

好了，理论说完了，咱们写个实际的程序。向量加法是CUDA的「Hello World」，虽然简单，但包含了所有核心要素。

#include <cuda_runtime.h>
#include <stdio.h>

// 定义kernel函数
__global__ void vecAdd(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];
    }
}

int main() {
    int N = 1 << 20;  // 约100万个元素
    size_t size = N * sizeof(float);

    // 1. 在Host端分配内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

    // 2. 在Device端分配内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 3. 将数据从Host拷贝到Device
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 4. 配置kernel执行参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 调用kernel
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 6. 将结果从Device拷贝回Host
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 7. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 8. 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

这段代码有8个步骤，我建议你记住这个流程：分配Host内存 → 分配Device内存 → 数据拷贝到Device → 配置kernel参数 → 调用kernel → 结果拷回Host → 验证 → 释放内存。

小技巧：每次调用CUDA API后，最好检查一下返回值。我习惯写一个CHECK(cudaError)宏，这样出错了能立刻定位。别问我为什么——因为我在调试时吃过太多亏了。

关于kernel配置，blocksPerGrid的计算方式很关键。用(N + threadsPerBlock - 1) / threadsPerBlock可以保证覆盖所有元素，即使N不是threadsPerBlock的整数倍。kernel内部的if (i < N)检查就是处理边界情况的。

知识体系总览

下面这张图展示了本章的核心知识结构，我画了很久，希望能帮你理清思路：

这张图把三个核心概念串起来了。编程模型是骨架，线程层次是肌肉，内存层次是血液。第一个程序就是让它们协同工作的实战演练。

好了，第一章的内容就到这里。记住：多写代码，多调试，多思考。CUDA编程没有捷径，但只要你理解了这些基础，后面的路就好走了。

📚 课程目录

第一章：CUDA编程基础

1.1 CUDA编程模型：Host与Device

1.2 线程层次结构：Grid、Block、Thread

1.3 内存层次结构：Global、Shared、Register

1.4 第一个CUDA程序：向量加法

知识体系总览