GPU的核心结构：它是是围绕一个流式多处理器（SM）的扩展阵列搭建的
图中包含以下的几个核心部件：
LD/ST（Load/Store）单元：负责内存加载和存储操作，用于在寄存器和全局内存、共享内存之间传输数据。
SFU（Special Function Unit）：专门执行特殊数学运算，比如三角函数、指数函数和倒数等，提升计算复杂函数的效率。
CUDA Core：是NVIDIA GPU中的基本计算单元，负责执行基本的算术逻辑运算（如加法、乘法），是实现并行计算的核心组件。
Dispatch Unit（分发单元）：负责将指令从线程束调度器发送到适当的执行单元（例如 CUDA Cores、LD/ST 或 SFU）。
Warp Scheduler（线程束调度器）：负责管理和调度一个SM（Streaming Multiprocessor）内的warp（一组32个线程打包）；它选择准备好的warp并将它们的指令分派给相应的执行单元。
NVIDIA GPU 中的 SM（Streaming Multiprocessor，流式多处理器） 是 GPU 的核心执行单元。
常说的CUDA目前一般说的是用nvcc cuda编译器编译的扩展后的类C/C++程序 .cu后缀
CUDA C/C++是在标准 C/C++ 基础上进行扩展的一种语言，新增了用于编写设备端代码的关键字和语法结构，例如
global、device和host等，这些扩展使得开发者可以在同一份代码中灵活地划分主机端与设备端的执行逻辑。CUDA C/C++ 还提供了丰富的运行时 API 和库函数（如
cudaMalloc、cudaMemcpy、cudaLaunchKernel等），用于管理 GPU 内存、启动核函数以及实现 CPU 与 GPU 之间的数据传输与同步。
第一个cuda 程序
#include<cuda_runtime.h>#include<iostream>// 这段代码是在gpu当中执行的 函数称为核函数__global__voidhello_world(void){//被 `global` 修饰的函数是运行在 GPU 上的设备端函数，在 CUDA 中被称为 kernel（核函数）核函数是实现并行计算的核心机制，它将由多个线程同时执行，从而实现高度并发的计算任务//每个线程答应线程id，id为零的线程输出hello worldprintf("thread idx: %d\n",threadIdx.x);if(threadIdx.x==0){printf("GPU: Hello world!\n");}}intmain(intargc,char**argv){printf("CPU: Hello world!\n");hello_world<<<1,10>>>();//这里三个尖括号相当于是启动程序，启动GPU执行相关代码cudaDeviceSynchronize();// cpu要等待gpu执行结束if(cudaGetLastError()!=cudaSuccess){std::cerr<<"CUDA error: "<<cudaGetErrorString(cudaGetLastError())<<std::endl;return1;}else{std::cout<<"GPU: Hello world finished!"<<std::endl;}std::cout<<"CPU: Hello world finished!"<<std::endl;return0;}<<<1,10>>>1 个线程块（block），每个线程块中包含10 个线程（thread）。
常用写法如下，一般线程上线1024
int threadsPerBlock = 256;int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;kernel<<<blocksPerGrid, threadsPerBlock>>>();除了thread和block的概念，还有grid![[Pasted image 20260702172518.png]]
图中的箭头可以理解为线程 —— 这是 CUDA 中最小的执行单元。多个线程可以组织成一个线程块（Thread Block），而多个线程块又组成一个网格（Grid）。整个核函数所启动的所有线程都包含在这个网格之中。
当启动一个核函数时，其网格中的线程块会被动态分配到 GPU 上各个可用的 SM 中执行。==一旦某个线程块被调度到某个 SM 上，其中的所有线程将始终在该 SM 上并发执行，不会迁移到其他 SM 上。==由于 GPU 包含多个 SM，多个线程块可以在不同的 SM 上同时运行，实现真正的物理级并行，而不是像 CPU 那样依赖时间片轮转。SM充当了硬件级的实时控制中枢，在硬件资源允许范围内驱动大量线程并发执行，最终依靠内部的束调度器（Warp Scheduler）精准调度指令，确保海量计算任务高效、有序地在核心上推进。多个线程块可以被分配到同一个SM上。
此外，线程块中的线程会以32 个线程为一组进行调度（调度的时候受限一个SM中的可用资源），这组线程被称为线程束（Warp）。每个线程束中的线程会同时执行相同的指令，但各自使用不同的数据。每个线程拥有独立的程序计数器和寄存器状态，以实现数据级别的并行。SM 将线程块划分为多个 warp，并根据硬件资源调度它们执行，Warp Scheduler 将选中的 warp 的指令发送给对应的 Dispatch Unit（调度单元），随后Dispatch Unit 再将该指令分发到相应的功能单元：CUDA Cores、LD/ST 或 SFU。
这些 warp 将被 SM 实际调度和执行。也就是说，GPU 并不是以单个线程或整个线程块为单位进行调度，而是将线程块划分为多个 warp，每个 warp 包含 32 个线程，并由 SM 按照硬件资源进行调度执行。
所有线程运行相同的核函数代码： 无论是 100 个还是 100 万个线程，它们读入的代码指令都是完全一样的。但是不同的线程处理不同的数据： 虽然代码一样，但每个线程会根据自己的“身份”去处理数据集中不同的部分（例如：线程 A 处理像素点 1，线程 B 处理像素点 2）。
cuda线程的组织形式
虽然有 thread -> block ->grid 这个形式但是从下面这个函数中，我们可以看出cuda线程本质上还是一维形式组织的，block只是逻辑上的概念
__global__voidvecAdd(int*A,int*B,int*C,intN){// i=当前线程块中第几个线程+第几个线程块*每个线程块的线程个数inti=threadIdx.x+blockIdx.x*blockDim.x;if(i<N){intres=A[i]+B[i];printf("A[i] = %d\n",A[i]);C[i]=res;}}cuda的内存分层模型
和CPU的存储分成寄存器，缓存，主存等类似，GPU的存储也有分层
| 内存类型 | 存储位置 | 作用域 | 可读写性 | 生命周期 | 访问速度 | 申请方式 |
|---|---|---|---|---|---|---|
| 寄存器 | 片上（on-chip） | 单个线程 | 可读写 | 线程执行期间 | 极快 | 自动 |
| 共享显存 | 片上（on-chip） | 线程块内所有线程 | 可读写 | 线程块执行期间 | 快 | 显式声明 |
| 全局显存 | 片外（off-chip） | 所有线程 | 可读写 | 主机端配置 | 最慢 | 主机端分配 |
也分为寄存器，共享显存（线程块内所有函数共享），全局显存，字面意思所有函数共享。
根据共享显存访问速度比全局显存访问速度快的特点，我们之后可以在计算中用共享显存对计算进行加速，这个之后再做深入学习
