GPU的核心结构它是是围绕一个流式多处理器SM的扩展阵列搭建的图中包含以下的几个核心部件LD/STLoad/Store单元负责内存加载和存储操作用于在寄存器和全局内存、共享内存之间传输数据。SFUSpecial Function Unit专门执行特殊数学运算比如三角函数、指数函数和倒数等提升计算复杂函数的效率。CUDA Core是NVIDIA GPU中的基本计算单元负责执行基本的算术逻辑运算如加法、乘法是实现并行计算的核心组件。Dispatch Unit分发单元负责将指令从线程束调度器发送到适当的执行单元例如 CUDA Cores、LD/ST 或 SFU。Warp Scheduler线程束调度器负责管理和调度一个SMStreaming Multiprocessor内的warp一组32个线程打包它选择准备好的warp并将它们的指令分派给相应的执行单元。NVIDIA GPU 中的 SMStreaming 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 程序#includecuda_runtime.h#includeiostream// 这段代码是在gpu当中执行的 函数称为核函数__global__voidhello_world(void){//被 global 修饰的函数是运行在 GPU 上的设备端函数在 CUDA 中被称为 kernel核函数核函数是实现并行计算的核心机制它将由多个线程同时执行从而实现高度并发的计算任务//每个线程答应线程idid为零的线程输出hello worldprintf(thread idx: %d\n,threadIdx.x);if(threadIdx.x0){printf(GPU: Hello world!\n);}}intmain(intargc,char**argv){printf(CPU: Hello world!\n);hello_world1,10();//这里三个尖括号相当于是启动程序启动GPU执行相关代码cudaDeviceSynchronize();// cpu要等待gpu执行结束if(cudaGetLastError()!cudaSuccess){std::cerrCUDA error: cudaGetErrorString(cudaGetLastError())std::endl;return1;}else{std::coutGPU: Hello world finished!std::endl;}std::coutCPU: Hello world finished!std::endl;return0;}1,101 个线程块block每个线程块中包含10 个线程thread。常用写法如下一般线程上线1024int threadsPerBlock 256;int blocksPerGrid (N threadsPerBlock - 1) / threadsPerBlock;kernelblocksPerGrid, 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当前线程块中第几个线程第几个线程块*每个线程块的线程个数intithreadIdx.xblockIdx.x*blockDim.x;if(iN){intresA[i]B[i];printf(A[i] %d\n,A[i]);C[i]res;}}cuda的内存分层模型和CPU的存储分成寄存器缓存主存等类似GPU的存储也有分层内存类型存储位置作用域可读写性生命周期访问速度申请方式寄存器片上on-chip单个线程可读写线程执行期间极快自动共享显存片上on-chip线程块内所有线程可读写线程块执行期间快显式声明全局显存片外off-chip所有线程可读写主机端配置最慢主机端分配也分为寄存器共享显存线程块内所有函数共享全局显存字面意思所有函数共享。根据共享显存访问速度比全局显存访问速度快的特点我们之后可以在计算中用共享显存对计算进行加速这个之后再做深入学习