还记得当我第一次学习C语言编程的时候，写下的第一个程序便是Hello World!，相信读者也和我一样。CUDA之旅依旧从“Hello World!”开始，让“Hello World!”伴随我们开启新的篇章！

欢迎大家进入我的哔哩哔哩频道进行学习！
https://developer.nvidia.com/blog/even-easier-introduction-cuda/
https://developer.nvidia.com/blog/unified-memory-in-cuda-6/

1 CUDA程序结构

CUDA是由NVIDIA提出的异构计算编程模型，核心思想是通过扩展C/C++语言（如添加关键字__global__、__device__等），使开发者能够利用CPU（Host）和GPU（Device）的协同能力。CPU负责逻辑控制、内存分配及调用设备函数，而GPU专注于并行计算任务。开发者无需学习新语言，只需在现有C/C++代码中嵌入CUDA关键词即可实现异构环境下的并行执行。

说明：在今后的学习中，CPU、主机、Host可以理解为一个意思，GPU、设备、Device也是同理。

前面的章节我们已经知道CUDA的使用是基于CPU+GPU异构架构，GPU可以看作是CPU的协处理器，真正的大脑是CPU，想要GPU工作的话需要CPU给它下达命令。因此，一个CUDA程序应该既有主机（Host）代码，也有设备（Device）代码。主机对设备的调用是通过核函数（kernel function）来实现的。所以，一个典型的、简单的CUDA程序的结构具有下面的形式：

int main()
{
    Host代码
    核函数调用代码
    Host代码
    
    return0;
}

CUDA程序结构包含以下关键步骤：

代码分层：
- Host代码：运行在CPU上，负责串行逻辑（如数据初始化、资源管理）；
- Device代码：运行在GPU上，以核函数（Kernel）形式执行并行任务。核函数通过__global__修饰，由Host调用，并通过<<<grid, block>>>语法配置线程结构（如线程块数量和每块线程数）。
执行流程：
- Host分配并初始化数据，通过cudaMalloc分配设备内存，并用cudaMemcpy在Host与Device间传输数据；
- Host调用核函数启动GPU计算（如hello<<<1,1>>>()），GPU线程并行执行相同指令；
- 由于GPU异步执行，Host需调用cudaDeviceSynchronize()等待GPU完成，再通过cudaMemcpy将结果回传Host，最后释放内存。

2 核函数

CUDA核函数是GPU并行计算的核心，其设计与使用需遵循相应的规范，下面我们来介绍核函数的编码方式及注意事项。

2.1 核函数定义规范

CUDA中的核函数与C/C++中的函数调用总体上是相似的，CUDA可以看作是C/C++之上的扩展，因此会有一些额外的定义规范，程序设计者要严格遵守这些规范。

__global__修饰符：必须使用__global__修饰符声明（双下划线）；
核函数返回值：必须为void；

__global__修饰符和返回值类型void的编写次序是可以随意排序的，也就是说，下面两种定义核函数的方式都是正确的；

__global__ void kernel_fun();

void __global__ kernel_fun();

核函数禁用变长参数（如 printf 式可变参数），因线程独立性要求确定性的执行路径；
核函数禁用静态变量（Static Variables），避免线程间数据竞争（多线程并发访问同一内存）；
禁用函数指针，防止执行路径不可预测和编译优化失效。

2.2 调用语法

核函数就像普通C++函数一样，只有发生函数调用后，核函数才会执行，其调用语法如下所示：

kernel_fun<<<grid_dim, block_dim, shared_mem_size, stream>>>();

grid_dim：网格维度（线程块数量）
block_dim：线程块维度（每块线程数，上限1024）
shared_mem_size：动态共享内存大小（字节）
stream：执行流（默认0）

前面的课程只用到设置grid_dim和block_dim这两个参数，shared_mem_size参数将会在共享内存部分讲解，stream参数将会在CUDA流部分进行讲解，现在只需要忽略掉shared_mem_size和stream这两个参数。像下面这样调用和函数：

kernel_fun<<<grid_dim, block_dim>>>();

2.3 内存访问规则

内存隔离
- 核函数只能访问GPU显存（Device Memory），无法直接读写CPU主机内存；
- 主机-设备数据传输需显式调用cudaMemcpy（同步/阻塞）或cudaMemcpyAsync（异步/非阻塞）；

说明：cudaMemcpy和cudaMemcpyAsync作用都是用作数据同步，数据同步又分成同步方式和异步方式，同步方式即阻塞方式，异步非阻塞方式。
在学习CUDA流之前，主要使用cudaMemcpy（阻塞）进行数据同步，cudaMemcpyAsync会在涉及CUDA流的代码中大量应用。

参数传递限制
- 参数通过常量内存（Constant Memory）传递；
- 大数据量应传递设备指针而非完整结构体，如下所示：
```
__global__ void add(float *d_a, float *d_b, int n);
```

2.4 执行特性与同步

Host代码与Device代码是异步执行的，也就说说，在调用核函数后，Host代码立即继续执行，不会阻塞。如果Host代码在调用核函数后没有其他任务执行，那么Host代码会立即跳出程序，GPU只看做CPU的协处理器，CPU可以看作是GPU的领导，CPU中的程序维护GPU的程序，当Host中的代码退出程序，那么Device的代码也就不能继续执行了。因此CUDA中引入了同步机制。
CUDA同步API接口如下：

cudaDeviceSynchronize();  // 等待所有GPU操作完成
cudaStreamSynchronize(stream);  // 等待特定流完成

cudaStreamSynchronize(stream)是CUDA流的知识内容，在现如今阶段，只需要关心cudaDeviceSynchronize()函数。

3 从代码入手

我们来看一下下面这段代码，Host和Device分别打印“Hello World!”：

文件：hello_world.cu

#include <stdio.h>
#include<stdlib.h> 

__global__ void print_from_gpu(void) {
	printf("Hello World! from thread [%d,%d] \
		From device\n", threadIdx.x, blockIdx.x);
}

int main(void) {
	printf("Hello World from host!\n");
	print_from_gpu<<<1, 1 >>>();
	cudaDeviceSynchronize();
	return 0;
}

3.1 nvcc编译器

nvcc（NVIDIA CUDA Compiler）是NVIDIA专为CUDA编程设计的核心编译器驱动，负责将混合了主机（CPU）代码和设备（GPU）代码的源文件编译为可执行程序。nvcc专为协调CPU（Host）与GPU（Device）协同执行设计。

3.1.1 异构编译流程

nvcc分离CUDA源文件（.cu）中的主机代码和设备代码：

主机代码：交给标准C/C++编译器（如gcc、g++）处理;
设备代码：编译为中间表示（PTX）或二进制机器码（cubin），供GPU执行。

主机代码完整的支持C++语法，与编译过程与传统C++编译并无二样。设备代码通过多阶段编译适配不同GPU架构，nvcc编译大致流程如下：

主机代码（Host Code）：
- 分离后交由C/C++编译器处理
设备代码（Device Code）：
- 生成PTX伪汇编代码：使用虚拟架构选项（如-arch=compute_80）编译为PTX（Parallel Thread Execution）伪汇编代码，确保跨代兼容性；
- 编译为二进制cubin：通过真实架构选项（如-code=sm_86）将PTX编译为特定GPU的二进制机器码cubin;
合并链接：合并主机目标文件与设备代码，生成可执行文件。

3.1.2 虚拟架构计算能力和真实架构计算能力

现在我们来解释一下虚拟架构选项（-arch=compute_89）和真实架构选项（-code=sm_89）的作用。

类型	定义	代表参数	核心目的
虚拟架构计算能力	面向PTX（Parallel Thread Execution）中间码的指令集兼容性声明	-arch=compute_XY	定义代码可使用的CUDA功能范围，确保跨代兼容性
真实架构计算能力	面向具体GPU硬件的二进制机器码（cubin）生成目标，对应物理GPU的计算能力版本号	-code=sm_ZW	生成适配特定硬件的优化代码，发挥硬件性能

注意：真实架构计算能力必须大于等于虚拟架构计算能力！！！

如下编译选项：

-arch=compute_60 -code=sm_89

上述编译选项，真实架构计算能力大于虚拟架构计算能力，可以成功编译。

如下编译选项：

-arch=compute_89 -code=sm_60

上述编译选项，编译器会报错。

nvcc编译器通过两阶段分离实现兼容性与性能平衡：

设备代码编译为PTX：
- 使用虚拟架构参数（如compute_60）生成与硬件无关的PTX中间码，仅依赖声明的功能集；
- 低版本PTX（如compute_60）可在高算力GPU（如sm_89）上通过JIT编译运行，实现向后兼容。
PTX编译为cubin：
- 使用真实架构参数（如sm_89）将PTX转为GPU可执行的二进制码，充分利用硬件特性（如张量核心）。

如果仅针对一个GPU编译程序，一般情况下将上述的两种计算能力指定为所用GPU的计算能力，这样可以最大限度的发挥GPU的性能，同时也省去了编译兼容性可执行文件的麻烦。

本课程所使用的GPU为NVIDIA GeForce RTX 4070，基于Ada Lovelace架构，计算能力为8.9，因此本课程的编译选型全部设置为：

-arch=compute_89 -code=sm_89

大家在编译程序的时候，一定要按照自己显卡实际计算能力去设置编译选项。

使用如下指令编译hello_world.cu：

nvcc hello_world.cu -o hello_world -arch=compute_89 -code=sm_89

运行代码:

./hello_world

输出如下：

Hello World from host!
Hello World! from thread [0,0]          From device

3.2 代码解析

hello_world.cu实现了一个基础的CPU-GPU异构协作示例，主机端（CPU）打印一段内容，设备端（GPU）通过核函数print_from_gpu打印一段包含当前线程索引信息的内容。

__global__关键字：声明此函数为GPU核函数。
- 核函数由主机（Host）调用，运行在设备（device）；
- 设备函数的返回类型必须是void。
核函数调用配置<<<1,1>>>：这种方式是对设备（GPU）函数的调用，而不是对主机（Host）函数的调用。1,1参数决定了内核中要启动的线程数量。后面的章节我们还会详细介绍，在这里，1,1参数的设置意味着我们在GPU上只启动了一个线程，因此打印消息只有一行。
- 第一个1，网格（Grid）中仅有1个线程块（Block）；
- 第二个1：每个线程块包含1个线程（Thread），总计1线程。
内置变量
- threadIdx.x，线程在块内的索引（此处为0）；
- blockIdx.x：程块在网格中的索引（此处为0）。
cudaDeviceline()：是 CUDA 运行时 API 中的核心同步函数，用于协调主机（CPU）与设备（GPU）间的异步执行。cudaDeviceline()是相对粗粒度的同步方式，后面课程我们还会介绍更加细粒度的同步方式。
- 同步所有设备操作：阻塞调用它的主机线程，直到当前 GPU 设备上所有先前启动的核函数、数据传输或其他 CUDA 操作均完成；
- 强制刷新设备输出缓冲区（如GPU的printf输出），确保日志或调试信息完整显示。

4 线程层次

CUDA threads：CUDA线程在CUDA核心上执行。CUDA线程与CPU线程不同。CUDA线程非常轻量级，并提供快速上下文切换。快速上下文切换的原因是由于GPU和基于硬件的调度器中的大寄存器大小的可用性。与CPU相比，线程上下文存在于寄存器中，其中线程句柄驻留在较低的存储器层次结构中，例如高速缓存。因此，当一个线程空闲/等待时，另一个线程可以几乎没有延迟地开始执行。每个CUDA线程必须执行相同的内核，并独立处理不同的数据（SIMT）。

CUDA blocks：CUDA线程被分组到一个称为CUDA块的逻辑实体中。CUDA块在单个流多处理器（SM）上执行。一个块在单个SM上运行，也就是说，一个块内的所有线程只能在一个SM中的核上执行，而不能在其他SM的核上执行。每个GPU可能有一个或多个SM，因此要有效地利用整个GPU;用户需要将并行计算划分为块和线程。

GRID/kernel：CUDA块被组合成一个逻辑实体，称为CUDA GRID。然后在设备上执行CUDA GRID。

我们