cuda从代码到GPU执行，再到返回到host

要弄清楚代码之后的东西

---

前言

重要性极高，属于基础中的精华。此处内容是基于H100的。

---

一、代码编写阶段

代码可以按照在CPU/GPU上运行而分为host代码和device代码
写cuda代码主要就是写这两部分的代码

二、NVCC编译阶段

在代码编写完成后，使用NVCC编译器进行编译，此过程会将host代码编译成CPU可执行代码，将device代码编译成GPU可理解的指令集PTX代码（GPU的中间代码，类似于汇编语言）

三、GPU启动阶段

在NVCC编译阶段完成之后，主机CPU通过CUDA API向CUDA驱动程序发送命令，请求GPU执行kernel函数。此时，内核函数、数据大小、线程布局等信息被封装为任务队列中的一个任务。

CUDA驱动程序在接收到CPU的请求命令之后，会根据不同的硬件架构（不同的GPU型号）选择最合适的指令，并将上一阶段中的PTX代码编译成SASS（Streaming Assembler—适用于特定NVIDIA架构的二进制代码）

GPU接收到上面传递的执行任务的时候，会通过线程块调度器将GPU中的一个个Block根据线程块调度机制动态分配到对应的SM上（一个SM处理一组Block，可以能是一个Block，也可能是多个Block）

在SM中的warp调度器会将在代码中指定的线程分配到相应的warp上

四、GPU执行阶段

GPU以warp为单位调度线程，先根据Blockdim来确定一个Block需要分配多少个warp，然后乘以块数，得到所需warp的数量。

每个warp中的线程会执行相同的内核函数代码，但处理不同的输入数据。内核中的指令被送到GPU的执行单元，这些一个个执行单元就是所谓的SP。

指令级并行（ILP）（指令阻塞时执行其他指令）和线程级并行（TLP）（线程阻塞时执行其他线程）是 GPU 提高执行效率的关键。GPU 会同时执行多个 warp，当某个 warp 等待数据时，其他 warp 可以继续执行，以隐藏访问延迟，达到延迟隐藏的效果。

当内核需要访问设备内存时，GPU 会尝试从 缓存 或 共享内存 中读取数据。如果缓存未命中，则会发起对全局内存（device memory）的请求。全局内存是 GPU 上的显存，访问延迟较高。为减少延迟，GPU使用了多级缓存（如L1和L2缓存），并且使用协同访存（coalesced memory access）机制来优化内存带宽。数据从显存加载到线程的寄存器中，在寄存器中执行计算操作

线程执行计算后，将结果写回到 全局内存 或 共享内存。
每个线程处理自己负责的那一部分数据，因此数据在执行过程中会存储在GPU的寄存器、共享内存、L1缓存或L2缓存等存储介质中。

五、其他补充，并行（Parallelism）

在GPU中的并行指的是多个线程同时在不同的计算单元上执行。这意味着：

真正的同时执行：GPU硬件设计允许多个计算核心（例如CUDA核心）在同一时刻执行计算操作。因此，当一个warp中的32个线程执行同一个指令时，它们实际上是同时在硬件的不同计算单元上并行进行。例如，多个ALU（算术逻辑单元）同时计算加法或乘法操作。

SIMT（Single Instruction, Multiple Threads）模型：在GPU中，多个线程可以同时执行同样的指令，操作不同的数据。以warp为单位，所有线程会同步地执行同样的指令，并行计算不同的值。例如，假设我们要对32个矩阵元素进行乘法操作，那么warp中的32个线程将同时执行乘法操作，每个线程负责计算不同的矩阵元素。

这种并行是硬件级的并行，即在同一个时钟周期内，多个计算单元并行工作，不是“依次运行”，而是同时运行。

六、共享代码机制

再解释一下，为什么只有一个代码，没做4个副本，但是四个warp却能同时执行。
在CUDA的架构设计中，并不需要为每个warp创建代码的多个副本。这是因为CUDA使用了一种指令流共享的机制，让多个线程组（warp）可以从同一份代码中提取指令，并独立执行。这是由GPU的硬件架构和SIMT（Single Instruction, Multiple Threads）模型共同实现的。

6.1.共享代码的机制

在CUDA中，代码只需存储一份，而多个线程（包括多个warp）可以共享这一份代码。这是通过以下方式实现的：

内核代码加载到指令缓存：当内核（kernel）启动时，内核代码会被加载到GPU的指令缓存中。这些指令缓存对于所有的线程和warp是共享的，无论是在同一个block中的warp，还是跨block的warp，所有线程都会从同一份指令缓存中获取指令。

SIMT模型的特点：CUDA的SIMT模型允许多个线程执行相同的指令流。因此，每个warp中的32个线程会根据同一指令集来执行相同的指令，而并不需要为每个warp复制代码。指令的调度和分发由GPU的硬件来管理，不同的warp虽然执行相同的代码，但它们操作不同的数据。

6.2. 每个线程有独立的上下文

虽然所有的warp和线程都共享同一份代码，但每个线程都有自己独立的上下文（上下文包括寄存器、局部变量等），这些上下文确保了每个线程执行相同代码时能够使用独立的数据和变量：

寄存器：每个线程有自己独立的寄存器空间，存储自己的局部变量和中间计算结果。即使所有线程执行相同的指令，每个线程在执行时访问的寄存器和局部变量都是独立的。

线程索引：每个线程通过threadIdx.x等内置变量区分自己在warp和block中的位置。因此，即使所有线程执行同样的指令，例如矩阵乘加操作，row和col等索引是每个线程自己计算出来的，确保它们操作不同的数据。

6.3. 指令调度与执行并行

在实际的执行过程中，GPU会利用硬件中的warp调度器来并行调度多个warp执行同一指令：

指令级并行：在GPU中，一个**Streaming Multiprocessor（SM）**可以调度多个warp同时执行。例如，在一个SM中，多个warp可以同时执行相同的加法、乘法、访存等操作。虽然代码只有一份，但GPU硬件允许多个warp在不同的执行单元上并行操作。

warp调度器的作用：每个SM中有多个warp调度器，这些调度器可以并行调度多个warp执行相同的指令集。warp调度器会确保各个warp能够高效地并行执行相同的内核代码。

6.4. 为什么不用复制代码？

在现代GPU架构中，为每个warp复制代码是没有必要的，原因如下：

硬件支持并行执行：GPU硬件专为并行设计，多个warp可以在同一时刻执行相同的代码。指令缓存和调度器能够处理并行执行的需求。

资源节省：只需要一份代码副本可以大幅减少对指令存储的需求。每个warp和线程只需要独立的数据上下文，而不需要独立的代码副本。

SIMT模型的设计：CUDA中的SIMT架构设计就是为了让多个线程在执行同样的指令时，只用一份指令流，从而减少代码冗余并提高执行效率。

---

总结

将Block调入SM时，是根据SM的资源来尽可能多的放入一个SM中，当此SM中资源不够的时候，再调其他SM。
将Thread分配到warp中时，是根据Blockdim的大小分配的，即根据一个Block中的Thread总量
每个执行单元最多可以执行一个warp的指令，也就是一次最多可以处理32个线程
每个 thread block 内的 thread（线程）都会从 0 开始编号，但这个编号实际上是通过 线程索引 来实现的。

以H100为例，一个H100的SM有四个warp调度器，即在每个时钟周期内，SM中的四个warp调度器最多可以从活跃的warp中挑选出四个warp执行。这意味着每个时钟周期内SM可以并行调度四个warp中的32*4这么多个线程。
然后可以在一个CLK之后调度其他的warp，warp的上限取决于这个SM的执行单元资源和存储资源。由于SM可以管理更多的warp，调度器会在下一个CLK周期调度其他的warp。时间片轮转调度机制确保多个warp可以有效地利用硬件资源，特别是在处理访存延迟时，通过切换warp来最大化执行效率。

warp上限取决于资源：

• 执行单元资源：SM中的算术逻辑单元（ALU）、浮点运算单元（FP）等执行单元决定了每个时钟周期可以并行执行的计算任务。
• 存储资源：SM的寄存器文件、共享内存和其他存储资源决定了可以支持的活跃warp数量上限。例如，寄存器数量限制了每个线程可以使用的寄存器数量，从而影响SM中可以并行执行的warp和线程数量。

warp调度器在CUDA执行模型中起到了重要作用。它不仅负责调度warp的执行，还负责管理SM中的资源，包括寄存器分配、共享内存分配等

CUDA中的每个线程的寄存器和存储资源（例如寄存器文件）实际上都是位于SM（Streaming Multiprocessor）上，并由warp调度器进行调度和管理。

线程的存储器（例如局部变量）主要是通过寄存器文件实现的。每个线程的局部变量和中间结果会存储在SM的寄存器文件中，寄存器是每个线程独有的，不能被其他线程访问。

warp调度器负责调度和分配这些寄存器资源。当多个warp同时执行时，warp调度器负责从寄存器文件中为每个warp中的线程分配寄存器，并确保这些寄存器能高效地支持线程的并行计算。

当线程的寄存器需求超过SM能够提供的寄存器数量时，可能会发生以下情况：
寄存器溢出：当寄存器不足时，线程会使用本地内存（通常存储在全局内存中）来存储溢出的数据。这种情况会大大降低性能，因为全局内存的访问延迟远高于寄存器。
减少warp数量：为了避免寄存器溢出，CUDA运行时可能会减少每个SM中同时执行的warp数量。这意味着减少了占用率，影响了GPU的计算效率。

为什么不需要复制代码？
代码共享（指令缓存）：CUDA内核代码只需要存储一份，所有线程都从同一个指令流中提取指令，保证代码不会被重复复制。

独立上下文：每个线程有自己独立的寄存器和局部变量，虽然执行相同的指令，操作的数据却不同。

硬件支持并行执行：GPU的设计使得多个warp能够在同一时刻从同一个代码中提取指令并并行执行。