CUDA高性能计算系列01：概述与GPU架构-洪萨配资

CUDA高性能计算系列01：概述与GPU架构

摘要：本系列的第一篇文章，我们将揭开 GPU 的神秘面纱。为什么深度学习需要 GPU？它的架构与 CPU 有何不同？我们将搭建 CUDA 开发环境，并手写第一个“Hello World”级别的 CUDA 程序——并行向量加法，体验 1000+ 核心同时工作的快感。

1. 引言：深度学习背后的算力引擎

在深度学习的浪潮中，我们习惯了model.to('cuda')这一行简单的代码带来的性能飞跃。但你是否思考过，为什么 GPU（图形处理器）会成为 AI 时代的算力霸主，而不是传统的 CPU？

当我们训练一个 ResNet 或者 Transformer 时，本质上是在进行海量的矩阵乘法和向量运算。这些计算具有高度的并行性：计算矩阵中一个元素的值，通常不依赖于其他元素的实时计算结果。

CUDA (Compute Unified Device Architecture) 是 NVIDIA 推出的并行计算平台和编程模型，它允许开发者直接访问 GPU 的虚拟指令集和并行计算元素。掌握 CUDA，意味着你不再局限于调用现成的 API，而是真正拥有了驾驭硬件底层的能力。

2. CPU vs GPU：设计哲学的差异

CPU 和 GPU 的设计目标截然不同，这也决定了它们擅长的领域不同。

2.1 延迟导向 (Latency-oriented) vs 吞吐量导向 (Throughput-oriented)

CPU (Central Processing Unit)：被设计为通用处理器。它拥有巨大的缓存（Cache）和复杂的控制逻辑（Control Unit），擅长处理复杂的逻辑分支、操作系统任务和串行指令。它的目标是最小化单个任务的延迟。
GPU (Graphics Processing Unit)：被设计为专用计算加速器。它牺牲了复杂的控制逻辑和部分缓存，换取了海量的算术逻辑单元 (ALU)。它的目标是最大化整体任务的吞吐量。

我们可以用一个 Mermaid 图表来直观对比两者的晶体管分配逻辑：

比喻：
CPU像一辆法拉利，运送少量货物（数据）非常快。
GPU像一辆重型卡车，虽然启动慢（延迟高），但一次能运送成吨的货物。

2.2 异构计算 (Heterogeneous Computing)

CUDA 采用异构计算模式。在这个模式中，我们有两个主要角色：

Host (主机)：CPU 及其内存（Host Memory）。负责复杂的逻辑控制、数据读取和环境初始化。
Device (设备)：GPU 及其显存（Device Memory）。负责大规模的数据并行计算。

一个典型的 CUDA 程序执行流程如下：

Copy: Host 将数据从 Host Memory 复制到 Device Memory。
Compute: Host 指挥 Device 启动核函数 (Kernel)，GPU 上的数千个线程并行计算。
Copy Back: Host 将计算结果从 Device Memory 取回 Host Memory。

3. CUDA 编程基础

3.1 核心概念

在 CUDA C++ 中，我们通过一些特殊的关键字来区分代码是在 CPU 上跑还是在 GPU 上跑：

关键字	执行位置	调用位置	备注
`__global__`	Device (GPU)	Host (CPU)	核函数 (Kernel)，这是并行计算的入口
`__device__`	Device (GPU)	Device (GPU)	GPU 内部调用的辅助函数
`__host__`	Host (CPU)	Host (CPU)	普通的 C++ 函数 (默认)

3.2 向量加法实战 (Vector Addition)

让我们通过一个经典的“向量加法”来演示 CUDA 程序的完整生命周期。
数学公式非常简单：对于长度为N NN的向量A AA和B BB，计算C = A + B C = A + BC=A+B，即：
C i = A i + B i , e x t f o r i = 0 , 1 , e x t d o t s , N − 1 C_i = A_i + B_i, ext{for } i = 0, 1, ext{dots}, N-1Ci=Ai+Bi,extfori=0,1,extdots,N−1

代码实现 (`vector_add.cu`)

#include<stdio.h>#include<cuda_runtime.h>// 错误检查宏，用于捕获 CUDA Runtime API 的错误#defineCHECK(call)\{\constcudaError_t error=call;\if(error!=cudaSuccess)\{\printf("Error: %s:%d, ",__FILE__,__LINE__);\printf("code:%d, reason: %s\n",error,cudaGetErrorString(error));\exit(1);\}\}// ---------------------------------------------------------// 1. Kernel Definition (核函数定义)// ---------------------------------------------------------// __global__ 表示该函数在 GPU 上运行，由 CPU 调用__global__voidvectorAdd(constfloat*A,constfloat*B,float*C,intnumElements){// 计算当前线程的全局索引// blockDim.x: 每个 Block 包含的线程数// blockIdx.x: 当前 Block 在 Grid 中的索引// threadIdx.x: 当前 Thread 在 Block 中的索引inti=blockDim.x*blockIdx.x+threadIdx.x;// 边界检查，防止越界访问if(i<numElements){C[i]=A[i]+B[i];}}// ---------------------------------------------------------// 2. Main Host Code// ---------------------------------------------------------intmain(void){// 向量大小：50000 个元素intnumElements=50000;size_t size=numElements*sizeof(float);printf("[Vector addition of %d elements]\n",numElements);// --- Host 侧内存分配与初始化 ---float*h_A=(float*)malloc(size);float*h_B=(float*)malloc(size);float*h_C=(float*)malloc(size);// 初始化向量 A 和 Bfor(inti=0;i<numElements;++i){h_A[i]=rand()/(float)RAND_MAX;h_B[i]=rand()/(float)RAND_MAX;}// --- Device 侧内存分配 ---float*d_A=NULL;float*d_B=NULL;float*d_C=NULL;CHECK(cudaMalloc((void**)&d_A,size));CHECK(cudaMalloc((void**)&d_B,size));CHECK(cudaMalloc((void**)&d_C,size));// --- Copy Data: Host -> Device ---printf("Copy input data from the host memory to the CUDA device\n");CHECK(cudaMemcpy(d_A,h_A,size,cudaMemcpyHostToDevice));CHECK(cudaMemcpy(d_B,h_B,size,cudaMemcpyHostToDevice));// --- Launch Kernel (启动核函数) ---// 设定执行配置：每个 Block 256 个线程intthreadsPerBlock=256;// 计算需要的 Block 数量，(N + 255) / 256 确保向上取整覆盖所有元素intblocksPerGrid=(numElements+threadsPerBlock-1)/threadsPerBlock;printf("CUDA kernel launch with %d blocks of %d threads\n",blocksPerGrid,threadsPerBlock);vectorAdd<<<blocksPerGrid,threadsPerBlock>>>(d_A,d_B,d_C,numElements);// 检查 Kernel 是否执行出错 (异步错误)CHECK(cudaGetLastError());// 同步主机与设备 (可选，用于计时或调试)CHECK(cudaDeviceSynchronize());// --- Copy Result: Device -> Host ---printf("Copy output data from the CUDA device to the host memory\n");CHECK(cudaMemcpy(h_C,d_C,size,cudaMemcpyDeviceToHost));// --- 验证结果 ---for(inti=0;i<numElements;++i){if(fabs(h_A[i]+h_B[i]-h_C[i])>1e-5){fprintf(stderr,"Result verification failed at element %d!\n",i);exit(EXIT_FAILURE);}}printf("Test PASSED\n");// --- 释放内存 ---CHECK(cudaFree(d_A));CHECK(cudaFree(d_B));CHECK(cudaFree(d_C));free(h_A);free(h_B);free(h_C);return0;}

3.3 编译与运行

CUDA 程序使用 NVIDIA 的编译器nvcc进行编译。它会自动将 Host 代码交给gcc/g++处理，将 Device 代码编译成 PTX (Parallel Thread Execution) 中间代码，最终生成 GPU 机器码。

假设保存为vector_add.cu：

# 编译nvcc vector_add.cu -o vector_add# 运行./vector_add

预期输出：

[Vector addition of 50000 elements] Copy input data from the host memory to the CUDA device CUDA kernel launch with 196 blocks of 256 threads Copy output data from the CUDA device to the host memory Test PASSED

4. 关键点解析

cudaMalloc&cudaMemcpy：这是 GPU 编程中最基础的内存管理 API。记住，GPU 的内存（显存）是独立的，CPU 无法直接读取显存中的指针，必须通过cudaMemcpy进行搬运。
<<<blocksPerGrid, threadsPerBlock>>>：这是 CUDA 独特的执行配置语法。它告诉 GPU：“请启动blocksPerGrid个线程块，每个块里有threadsPerBlock个线程，去执行这个函数”。
int i = blockDim.x * blockIdx.x + threadIdx.x：这行代码是 CUDA 编程的灵魂。它将并行的线程映射到线性的数据索引上。我们将在下一篇文章中详细图解这个计算过程。

5. 总结与下篇预告

今天我们成功迈出了 GPU 编程的第一步，理解了 CPU 与 GPU 架构的本质区别，并跑通了完整的向量加法流程。

虽然这个程序能跑，但你可能会问：为什么是 256 个线程？为什么会有 Block 和 Grid 的概念？能不能一个 Block 开 50000 个线程？

这涉及到 GPU 的硬件调度机制。下一篇CUDA系列02_线程模型与执行配置，我们将深入 Grid-Block-Thread 的三级架构，并揭示 Warp（线程束）的秘密，教你如何科学地设定线程数量以获得最佳性能。

参考文献

NVIDIA Corporation.CUDA C++ Programming Guide. 2024. https://docs.nvidia.com/cuda/cuda-c-programming-guide/
Sanders, J., & Kandrot, E.CUDA by Example: An Introduction to General-Purpose GPU Programming. Addison-Wesley Professional, 2010.

CUDA高性能计算系列01：概述与GPU架构