文档教程人工智能【免费下载链接】AISystemAISystem 主要是指AI系统包括AI芯片、AI编译器、AI推理和训练框架等AI全栈底层技术项目地址https://gitcode.com/GitHub_Trending/ai/AISystem点击查看免费下载本篇文章聚焦于 AI 芯片中最核心的两种并行计算模型——SIMD单指令多数据与 SIMT单指令多线程从计算本质、硬件结构、流水线调度到 NVIDIA CUDA 的实际实现层层展开。读者学完本篇后将能准确区分 SIMD 与 SIMT 的执行差异理解 Warp、Thread Block、Grid 与 CUDA Core、SM 之间的软硬件对应关系并看懂基于blockIdx/blockDim编写 CUDA Kernel如矩阵乘的底层原理为后续理解 AI 编译器的向量化、算子优化与 NPU/GPU 架构设计打下基础。本篇内容以 02SIMTSIMD.md 为主体骨架并结合 07Thought 系列课程中 03SPMT.mdSPMD/CUDA 关系与 04NVSIMT.mdCUDA 硬件关系以及code/目录下的 CUDA 示例进行纵深扩充。SIMD 计算本质一条指令并行处理多个数据SIMDSingle Instruction Multiple Data是指对多个执行同样操作的处理元素在同一时刻执行完全相同的计算操作。它利用的是数据级并行Data-Level Parallelism而不是并发性即有多个计算但只有一个进程在运行。一条单一命令可以同时对多个数据值进行操作因此 SIMD 是提升 CPU 计算能力、实现数据并行最经典的手段其硬件代价是更宽位数的计算单元 ALU 与较小的控制逻辑。需要特别强调的是SIMD 依然是单线程操作硬件上只需要一个计算核心只不过一次操作多个数据。这与 GPU 的多线程并行有本质区别——SIMD 的计算本质是在多个数据上并行执行相同操作的硬件部分。以两个向量相乘为例假设向量 A 与向量 B 各有 4 个元素逐元素相乘得到结果 C$$C[0: 3] A[0: 3] × B[0: 3]$$为了使一次乘法能够同时完成多个元素的计算硬件上需要增加 ALU 单元的数量形成多个处理单元Process Unit同时增加功能单元的数据通路数量由控制单元Control Unit将数据传送给处理单元从而在单一时钟周期内整体提升硬件的计算吞吐量。在图片中可以看到SIMD 硬件由控制单元内含 PC 程序计数器与 IR 指令寄存器、多个处理单元每个处理单元包含 ALU 与寄存器文件 Reg File以及存储器构成指令由控制单元统一译码数据通过多条数据通路并行送入各处理单元执行这正是单指令、多数据通路的硬件体现。SIMD 的优缺点在实际计算过程中SIMD 有其鲜明的优缺点缺点SIMD 使用独立线程该线程能同时进行多个数据元素的计算但由于 ALU 宽度的限制计算时要求数据类型、格式、大小必须严格对齐无法对任意散乱的数据进行灵活操作。优点在一定程度上可以提升计算性能充分利用内存数据总线带宽多个数据可以同时从内存读出和写入。例如 $C[0: 3] A[0: 3] × B[0: 3]$ 操作在使用 SIMD 之后代码量为原来的 1/4执行周期也相应降为原来的 1/4。从非 SIMD 到 SIMD指令数降为 1/4以 $C[0: 3] A[0: 3] × B[0: 3]$ 计算为例计算机在没有使用 SIMD 时实际执行的指令如下总共有 4 组读-加-存序列对 4 个元素进行逐元素相加或相乘t1 LD B, i t2 LD C, i t3 t1 t2 ST A, i, t3 t1 LD B, i1 t2 LD C, i1 t3 t1 t2 ST A, i1, t3 t1 LD B, i2 t2 LD C, i2 r3 t1 t2 ST A, i2, t3 t1 LD B, i3 t2 LD C, i3 r3 t1 t2 ST A, i3, t3SIMD 本身是对指令的控制在使用 SIMD 之后只需要一组指令即可每条指令后面的4表示该指令执行时同时对 4 个元素进行操作。编译器会将下面的代码编译成硬件能够识别的 SIMD 指令v1 LD B, i, 4 v2 LD C, i, 4 v3 v1 v2, 4 ST A, i, 4, v3对比可见指令代码量变为原来的 1/4执行周期也相应降为原来的 1/4执行效率得到显著提升。指令集落地Intel MMX/SSE/AVX 与 ARM NEONSIMD 能力的落地依赖指令集扩展。Intel 从MMXMultiMedia eXtensions开始支持 SIMD它是 Intel 于 1996 年推出的一种 SIMD 指令集扩展用于对多个数据元素同时执行相同操作包含数据移动指令、整数运算指令、逻辑运算指令等可加速多媒体处理、图像处理等应用。随着技术发展Intel 后续推出 SSEStreaming SIMD Extensions、AVXAdvanced Vector Extensions等指令集扩展进一步提高了处理器对 SIMD 计算的支持和性能。ARM 则通过NEON将 SIMD 扩展引入 ARM-Cortex 架构。ARM NEON 是 2004 年推出的 SIMD 扩展技术其指令集包括数据加载/存储指令、整数运算指令、浮点运算指令等同样可以对多个数据元素同时执行相同操作为移动端多媒体、图像处理应用提供更高的计算性能。NEON SIMD 单元位宽为128-bit包含 16 个 128-bit 寄存器能够被用作 32 个 64-bit 寄存器。这些寄存器可被当作同等数据类型的 vector 使用此时数据是对齐的、元素格式也是相同的因此可以用一个进程对多个数据进行计算一个寄存器位宽 128 bit可存放 4 个 32-bit 元素。相关代码如下//对四个数据同时进行乘法操作 C[0:3] A[0:3]*B[0:3] //一个寄存器 128bit可以存放 4x32bits15 寄存器存放向量 B vldmia.32 r0!, {s15} //通过 s14 寄存器存放向量 A vldmia.32 r1!, {s14} // s15 s15*s14 vmul.f32 s15, s15, s14 //保存 s15 的计算结果 vstmia.32 r2!, {s15}综上SIMD 最重要且最本质的改变是增加了硬件计算单元的数量与数据读取通路的数量同时对上层提供更多指令集。在实际编程中程序员很少会直接操作 SIMD 指令通常由编译器自动向量化完成。SIMT 计算本质单指令多线程的 GPU 并行模型SIMTSingle Instruction Multiple Threads单指令多线程是英伟达NVIDIA提出的、基于 GPU 的新概念。与 SIMD 相比二者都通过将同样的指令广播给多个执行单元来实现数据并行和计算主要的不同在于SIMD 要求所有向量元素在统一的同步组里一个线程内同步执行而 SIMT 允许多个线程在一个 Warp 中独立执行。SIMT 类似 CPU 上的多线程有多个计算核心每一个核心中有独立的寄存器文件Register FileRF和计算单元Arithmetic Logic UnitALU但没有独立的指令缓存Instruction Cache、解码器和程序计数器PC命令从统一的指令缓存广播给多个 SIMT 核心。因此 SIMT 的所有核心各自独立在不同的数据上执行相同的计算操作——即执行命令相同、多个线程各有各的处理单元而 SIMD 则是共用同一个 ALU。还是以数组相乘 $C[0: 3] A[0: 3] × B[0: 3]$ 为例SIMT 给每个元素分配一个线程一个线程只需完成一个元素的乘法所有线程并行执行完成后两个数组的相乘就完成了。SIMT 硬件结构SIMT Core Cluster具体到 SIMT 的硬件结构SIMT 提供一个多核系统SIMT Core ClusterCPU 负责将算子Kernel加载到 SIMT Core Cluster 中。每个 SIMT 核SIMT Core有独立的 RFRegister File、ALU、Data Cache但只有一个程序计数器Program Counter和一个指令译码器指令被同时广播给所有的 SIMT 核执行具体计算。GPU 由多个 SIMT Core Cluster 组成每个 SIMT Core Cluster 由多个 SIMT Core 构成SIMT Core 中运行着多个 Thread Block。从图中可以看到完整的 SIMT 硬件组成外部 CPU 通过 Kernel Launch 将任务加载进来SIMT Core Cluster 内部包含 SIMT Stacks线程栈、Shared Memory、多个 Thread Block、Register File、Constant Cache、Texture Cache、Data Cache 与 Memory Port集群通过互联网络Network与内存分区Memory Partition含原子操作单元、末级缓存组与片外 DRAM交互数据。SIMT 硬件核心流水三个调度循环GPU 的 SIMT 可以看作一个特殊的 SIMD 结构SIMT 硬件核心流水可被分为 SIMT 前端SIMT front-end和 SIMD 后端SIMD back-end。流水线中存在三个调度循环取指循环包含 Fetch取指、I-Cache指令缓存、Decode译码和 I-Buffer指令缓冲四个阶段指令发射循环包含 I-Buffer、Score Board记分板、Issue发射和 SIMT-StackSIMT 栈四个阶段寄存器访问循环包含 Operand Collector操作数收集器、ALU 和 Memory 三个阶段。流水线中的三个调度循环共同组成 SIMT 硬件核心流水其中取指是将具体指令放在堆栈SIMT Stack中堆栈在运行时就会把所有的线程分发到具体的 ALU 中具体执行时采用 SIMD 的方式而 SIMT 主要完成具体线程的前端控制。换句话说SIMT 负责线程的前端控制SIMD 负责数据的后端执行这正是SIMT 前端 SIMD 后端流水线设计的精髓。SIMD 与 SIMT 的区别和联系结合上述内容SIMD 和 SIMT 的主要区别和联系可归纳如下SIMT 与 SIMD 的基本原理相同都是采用单指令多数据的思想SIMT 形式上是多线程但本质上在硬件端执行的还是单线程使用多个核心来实现多线程并行SIMT 比 SIMD 更灵活允许一条指令对数据分开寻址可以实现每个线程独立寻址SIMD 必须连续取址要求数据在类型、格式和大小方面严格对齐。因此SIMT 是 SIMD 的一种推广在编程模式上更加灵活对开发者更友好。NVIDIA CUDA 实现从 Grid、Block 到 Warp 的线程分级回顾 GPU 的线程分级在图形图像处理中会将图像进行切分网格Grid表示要执行的任务大的网格会被分成多个小的网格每个网格中包含很多相同线程Threads数量的块Blocks。此时线程分层执行块中的线程独立执行对像素数据进行处理和计算可以共享数据、同步数据交换。CUDA 并行编程模型基于单程序多数据Single Program Multiple DataSPMD模式SPMD 与 SIMT 的联系和区别在 03SPMT.md 中重点讲解。在 CUDA 编程中grid是线程块block的阵列集合线程块映射到 SM 上进行计算处理一个线程块可包含多个线程束Warp线程块的大小影响 CUDA Kernel 程序的性能GPU 执行时的最小单位是线程thread一个 block 中的线程可存取同一块共享内存并且可以快速进行同步。与 SIMD 不同的是SIMT 允许程序员为独立、标量线程编写线程级的并行代码也允许为协同线程编写数据并行代码。为了确保正确性开发者甚至可以忽略 SIMT 行为很少需要维护一个 warp 块内的线程分支只需维护相关代码即可获得硬件并行带来的显著性能提升。在一个线程块Thread Block中所有线程执行同一段代码在英伟达 GPU 中这段代码称为Kernel。每一个线程有一个自己的线程索引threadIdx.x用于计算内存地址和执行控制决策每个线程在执行时被分配唯一标识符因此可以通过程序准确控制每一个线程。块索引与线程内索引将多个线程块组合在一起就组成一个 Grid 线程组因此线程块是 SM 的基本调度单元SM 对应具体的硬件单元而线程块是编程抽象出来的概念。因为有多个线程块组合同时存在硬件计算单元在横向和纵向两个维度的排布因此线程索引通常由**块索引Block Index和线程内索引Thread Index Within Block**组成块索引用于标识当前线程所在的块Block线程内索引用于标识当前线程在所属块中的位置。在 CUDA 中使用blockIdx.x和blockDim.x访问块索引和块维度中的 x 分量blockIdx.x当前线程所在块的 x 方向索引。块索引是三维向量包括 x、y、z 三个分量blockDim.x当前块的 x 方向维度大小即该方向上线程数量。块维度同样是三维向量包括 x、y、z 三个分量。通过blockIdx.x和blockDim.x可以方便地获取当前线程所在块的 x 方向索引和当前块在 x 方向上的线程数量从而进行相应的计算和操作。结合threadIdx.x线程的全局一维 ID 通常写作int tid blockIdx.x * blockDim.x threadIdx.x;这一经典公式在仓库示例 cuda1.cu 中也有体现——它将 CPU 上的 for 循环for (int i 0; i10000; i) C[i] A[i] B[i];改写为 GPU Kernelint tid blockDim.x * blockIdx.x threadIdx.x;后每个线程独立取a[tid]、b[tid]并相加写入c[tid]。软件线程与硬件单元的对应关系回顾英伟达 GPU 软件和硬件之间的对应关系线程Thread对应 CUDA Core线程以线程块为单位被分配到 SM 上线程块Block对应 SMSM 维护线程块和线程 IDSM 管理和调度线程执行每个线程块又按照每个 Warp 中共32 个线程执行Warp 是 SM 的调度单位Warp 里的线程执行 SIMDBlock 线程块只在一个 SM 上通过 Warp 进行调度一旦在 SM 上调用了 Block 线程块就会一直保留到执行完 KernelSM 可以同时保存多个 Block 线程块块间并行执行。从图中可以清晰看到软硬件层级映射软件侧 Thread → Thread Block → Grid 分别对应硬件侧 CUDA Core → SM → Device完整 GPU 设备。这正是线程对应 CUDA Core、线程块对应 SM、网格对应 GPU的直观表达。需要补充的是Warp 之所以成为调度单位是因为SIMT 架构通过细粒度多线程Fine-Grained Multi-ThreadingFGMT调度将执行流水线细分使不同线程的指令交错执行从而以访存和计算并行来隐藏延迟详见 04NVSIMT.mdWarp 支持乱序执行以隐藏访存延迟线程寄存器值都保留在 RF 中英伟达通过 Warp Scheduler 硬件调度使 Warp 先访存完毕后交给 SIMD Pipeline 执行尽可能多的指令同时隐藏其它 Warp 的访存时间。实战CUDA 矩阵乘 Kernel在 AI 框架的开发流程方面首先会按照编程思想定义神经网络然后根据 AI 框架编写对应的程序AI 框架会自动构建计算正向图并根据自动微分原理构建反向图。其中在神经网络中比较重要的算子是矩阵乘。以 CUDA 代码为例实现 $C A × B$使用blockIdx.x和blockDim.x来访问块索引和块维度#include stdio.h #define N 4 // 矩阵大小 // 矩阵乘法的 CUDA 核函数 __global__ void matrixMultiplication(int *a, int *b, int *c) { // 使用 blockIdx.x 和 blockDim.x 来访问块索引和块维度 int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; int sum 0; for (int k 0; k N; k) { sum a[row * N k] * b[k * N col]; } c[row * N col] sum; } int main() { int a[N][N], b[N][N], c[N][N]; int *dev_a, *dev_b, *dev_c; // 分配内存 cudaMalloc((void**)dev_a, N * N * sizeof(int)); cudaMalloc((void**)dev_b, N * N * sizeof(int)); cudaMalloc((void**)dev_c, N * N * sizeof(int)); // 初始化矩阵 a 和 b for (int i 0; i N; i) { for (int j 0; j N; j) { a[i][j] i * N j; b[i][j] j * N i; } } // 将矩阵 a 和 b 传输到设备 cudaMemcpy(dev_a, a, N * N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * N * sizeof(int), cudaMemcpyHostToDevice); // 定义块大小和网格大小 dim3 blockSize(2, 2); dim3 gridSize(N / blockSize.x, N / blockSize.y); // 调用核函数 matrixMultiplicationgridSize, blockSize(dev_a, dev_b, dev_c); // 将结果传回主机 cudaMemcpy(c, dev_c, N * N * sizeof(int), cudaMemcpyDeviceToHost); // 打印结果 for (int i 0; i N; i) { for (int j 0; j N; j) { printf(%d , c[i][j]); } printf(\n); } // 释放内存 cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; }这段代码的要点在于线程索引计算row blockIdx.y * blockDim.y threadIdx.ycol blockIdx.x * blockDim.x threadIdx.x通过块索引、块维度与线程内索引的组合精确映射到输出矩阵的每个元素位置二维网格配置dim3 blockSize(2, 2)定义每个块为 2×2 共 4 个线程dim3 gridSize(N / blockSize.x, N / blockSize.y)定义网格为 2×2 共 4 个块恰好覆盖 4×4 矩阵的全部 16 个元素数据搬运流程cudaMalloc分配设备内存 →cudaMemcpyHostToDevice上传输入 → 启动 Kernel →cudaMemcpyDeviceToHost取回结果 →cudaFree释放这是 CUDA 编程的标准生命周期。仓库中的 cuda2.cu 与 cuda3.cu 展示了矩阵加法的CPU 串行版 → CUDA 并行版演进CPU 版本用双重 for 循环按index i j * N逐元素相加CUDA 版本则通过int i blockIdx.x * blockDim.x threadIdx.x;、int j blockIdx.y * blockDim.y threadIdx.y;让每个线程独立计算一个(i, j)位置的元素并用if (i N j N)做边界检查最后以add_matrixdimGrad, dimBlock(a, b, c, N)启动。这个对比直观展示了从 SISD 串行到 SIMT 并行的编程范式转变。从编程模型到硬件执行Kernel 到 Warp 的映射从更宏观的角度看CUDA 编程模型与 NVIDIA 硬件的关系可以总结为一条链路程序员编写 SPMD 风格的 Kernel → 编译器将线程逻辑编译为标量指令流 → 硬件把执行相同指令的线程动态组织成 Warp每 Warp 32 线程→ Warp 作为 SM 调度单位交给 SIMD Pipeline 执行。正如 03SPMT.md 所归纳的SIMD是指令的执行方式和对应映射的硬件体系结构SIMT是以 SIMD 指令为主、具有 Warp Scheduler 等硬件模块、支持 SPMD 编程模型的硬件架构SPMD是一种具体的并行编程模型类似于 CUDA 所提供的编程模式DSA是特殊的硬件架构NPU/TPU 等专门针对 AI 的特殊硬件架构应用于大规模数据处理、分布式存储等场景。值得注意的是NVIDIA 在 GPU 架构设计中加入了专门用于神经网络矩阵计算、同时支持混合精度计算的 Tensor Core因此 NVIDIA GPU 也变成了SIMT DSA的模式。此外Warp 的调度特性使得 SIMT 相比 SIMD 在可编程性上有根本性优势硬件层面解决了大部分流水编排问题开发者无需像 SIMD 那样费时费力地把数据凑成合适的矢量长度再传入硬件线程还可以独立执行、允许每个线程有不同的分支——这正是 SIMT 的核心。编程模型 vs 硬件执行模型在理解 SIMD 与 SIMT 之后还需要区分两个经常被混淆的概念编程模型Programming Model是程序员用来编写程序的抽象概念定义了程序员如何组织和控制计算机程序的方式。它提供一种简化的视图使程序员能够专注于程序的逻辑结构而不必考虑底层硬件细节通常包括编程语言、数据结构、算法和并发机制等方面用于描述程序的行为和交互。硬件执行模型Hardware Execution Model描述计算机硬件如何执行程序包括硬件结构、指令集架构、寄存器、内存层次结构、缓存、并行执行方式等方面。它决定了程序在计算机硬件上的实际执行方式——指令的执行顺序、数据的传输方式、并发执行的策略等硬件执行 SIMD 和 SIMT。二者的区别和联系如下区别编程模型从程序员的角度描述程序的组织和行为硬件执行模型从计算机硬件的角度描述程序的执行方式。编程模型关注程序的逻辑结构和抽象行为硬件执行模型关注程序在实际硬件上的执行细节。联系编程模型定义了程序的行为和交互方式硬件执行模型决定了程序如何在硬件上执行。程序员编写的程序最终会被映射到硬件执行模型上执行。理解二者关系可以帮助程序员优化程序性能、充分利用硬件资源。编程模型最终会通过编译器转换为硬件执行模型因此二者在概念层面有明显的差异。这也正是 AISystem 知识体系中AI 编译器环节的切入点编译器要做的就是把 CUDA/SPMD 这类编程模型翻译成目标硬件GPU/NPU能够高效执行的 SIMD/SIMT 指令流。小结与思考SIMD 与 SIMT 均基于单指令多数据执行但 SIMD 要求数据严格对齐且同步执行而 SIMT 允许线程独立寻址且可异步执行提高了灵活性NVIDIA CUDA 通过 SIMT 架构实现高效的并行计算利用线程块Block和网格Grid结构通过 CUDA Core 进行调度优化了 GPU 的性能编程模型与硬件执行模型相互关联前者为程序员提供抽象概念以组织程序后者描述程序在硬件上的实际执行方式理解二者关系有助于程序性能优化从实现层面看SIMT 的本质是SIMT 前端做线程控制、SIMD 后端做数据执行而 CUDA 的blockIdx/blockDim/threadIdx索引体系将线程逻辑精确映射到硬件执行单元这是理解 AI 芯片算子实现如矩阵乘、卷积与 AI 编译器后端优化的关键基础。如果想继续深入可以接着阅读本系列课程的 03SPMT.mdSIMD/SIMT/SPMD/DSA 的完整关系与 Warp 调度细节和 04NVSIMT.mdCUDA 与 NVIDIA 硬件架构、AMD ROCm 对比并结合 code 目录中的 CUDA 示例动手实践。赞分享文档教程人工智能【免费下载链接】AISystemAISystem 主要是指AI系统包括AI芯片、AI编译器、AI推理和训练框架等AI全栈底层技术项目地址https://gitcode.com/GitHub_Trending/ai/AISystem点击查看免费下载相关推荐Apache RocketMQ Controller 模式部署与升级实战指南Apache RocketMQ Controller 模式部署与升级实战指南 本篇指南以 RocketMQ 自动主从切换Controller 模式为主题完文档教程人工智能PyPTO Pro 编程范式深度解析SIMD 与 SIMT 并行执行模型及 AI Core 硬件基础PyPTO Pro 编程范式深度解析SIMD 与 SIMT 并行执行模型及 AI Core 硬件基础 PyPTO ProParallel Tensor/Ti人工智能编译器模型编译深度学习高性能计算CANNAscendPyPTO Pro 编程范式详解SIMD 与 SIMT 并行模型及 AI Core 硬件抽象PyPTO Pro 编程范式详解SIMD 与 SIMT 并行模型及 AI Core 硬件抽象 PyPTO Pro 是 CANN/PyPTO 项目中面向 NPU人工智能编译器模型编译深度学习高性能计算CANNAscend上一篇MinecraftByExample网络通信教程让你的模组实现跨客户端交互下一篇VisiData 列操作完全指南类型、格式化、拆分与派生列实战创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 AISystem 硬件基石：深入解析 SIMD 与 SIMT 计算本质及 NVIDIA CUDA 硬件实现
文档教程人工智能【免费下载链接】AISystemAISystem 主要是指AI系统包括AI芯片、AI编译器、AI推理和训练框架等AI全栈底层技术项目地址https://gitcode.com/GitHub_Trending/ai/AISystem点击查看免费下载本篇文章聚焦于 AI 芯片中最核心的两种并行计算模型——SIMD单指令多数据与 SIMT单指令多线程从计算本质、硬件结构、流水线调度到 NVIDIA CUDA 的实际实现层层展开。读者学完本篇后将能准确区分 SIMD 与 SIMT 的执行差异理解 Warp、Thread Block、Grid 与 CUDA Core、SM 之间的软硬件对应关系并看懂基于blockIdx/blockDim编写 CUDA Kernel如矩阵乘的底层原理为后续理解 AI 编译器的向量化、算子优化与 NPU/GPU 架构设计打下基础。本篇内容以 02SIMTSIMD.md 为主体骨架并结合 07Thought 系列课程中 03SPMT.mdSPMD/CUDA 关系与 04NVSIMT.mdCUDA 硬件关系以及code/目录下的 CUDA 示例进行纵深扩充。SIMD 计算本质一条指令并行处理多个数据SIMDSingle Instruction Multiple Data是指对多个执行同样操作的处理元素在同一时刻执行完全相同的计算操作。它利用的是数据级并行Data-Level Parallelism而不是并发性即有多个计算但只有一个进程在运行。一条单一命令可以同时对多个数据值进行操作因此 SIMD 是提升 CPU 计算能力、实现数据并行最经典的手段其硬件代价是更宽位数的计算单元 ALU 与较小的控制逻辑。需要特别强调的是SIMD 依然是单线程操作硬件上只需要一个计算核心只不过一次操作多个数据。这与 GPU 的多线程并行有本质区别——SIMD 的计算本质是在多个数据上并行执行相同操作的硬件部分。以两个向量相乘为例假设向量 A 与向量 B 各有 4 个元素逐元素相乘得到结果 C$$C[0: 3] A[0: 3] × B[0: 3]$$为了使一次乘法能够同时完成多个元素的计算硬件上需要增加 ALU 单元的数量形成多个处理单元Process Unit同时增加功能单元的数据通路数量由控制单元Control Unit将数据传送给处理单元从而在单一时钟周期内整体提升硬件的计算吞吐量。在图片中可以看到SIMD 硬件由控制单元内含 PC 程序计数器与 IR 指令寄存器、多个处理单元每个处理单元包含 ALU 与寄存器文件 Reg File以及存储器构成指令由控制单元统一译码数据通过多条数据通路并行送入各处理单元执行这正是单指令、多数据通路的硬件体现。SIMD 的优缺点在实际计算过程中SIMD 有其鲜明的优缺点缺点SIMD 使用独立线程该线程能同时进行多个数据元素的计算但由于 ALU 宽度的限制计算时要求数据类型、格式、大小必须严格对齐无法对任意散乱的数据进行灵活操作。优点在一定程度上可以提升计算性能充分利用内存数据总线带宽多个数据可以同时从内存读出和写入。例如 $C[0: 3] A[0: 3] × B[0: 3]$ 操作在使用 SIMD 之后代码量为原来的 1/4执行周期也相应降为原来的 1/4。从非 SIMD 到 SIMD指令数降为 1/4以 $C[0: 3] A[0: 3] × B[0: 3]$ 计算为例计算机在没有使用 SIMD 时实际执行的指令如下总共有 4 组读-加-存序列对 4 个元素进行逐元素相加或相乘t1 LD B, i t2 LD C, i t3 t1 t2 ST A, i, t3 t1 LD B, i1 t2 LD C, i1 t3 t1 t2 ST A, i1, t3 t1 LD B, i2 t2 LD C, i2 r3 t1 t2 ST A, i2, t3 t1 LD B, i3 t2 LD C, i3 r3 t1 t2 ST A, i3, t3SIMD 本身是对指令的控制在使用 SIMD 之后只需要一组指令即可每条指令后面的4表示该指令执行时同时对 4 个元素进行操作。编译器会将下面的代码编译成硬件能够识别的 SIMD 指令v1 LD B, i, 4 v2 LD C, i, 4 v3 v1 v2, 4 ST A, i, 4, v3对比可见指令代码量变为原来的 1/4执行周期也相应降为原来的 1/4执行效率得到显著提升。指令集落地Intel MMX/SSE/AVX 与 ARM NEONSIMD 能力的落地依赖指令集扩展。Intel 从MMXMultiMedia eXtensions开始支持 SIMD它是 Intel 于 1996 年推出的一种 SIMD 指令集扩展用于对多个数据元素同时执行相同操作包含数据移动指令、整数运算指令、逻辑运算指令等可加速多媒体处理、图像处理等应用。随着技术发展Intel 后续推出 SSEStreaming SIMD Extensions、AVXAdvanced Vector Extensions等指令集扩展进一步提高了处理器对 SIMD 计算的支持和性能。ARM 则通过NEON将 SIMD 扩展引入 ARM-Cortex 架构。ARM NEON 是 2004 年推出的 SIMD 扩展技术其指令集包括数据加载/存储指令、整数运算指令、浮点运算指令等同样可以对多个数据元素同时执行相同操作为移动端多媒体、图像处理应用提供更高的计算性能。NEON SIMD 单元位宽为128-bit包含 16 个 128-bit 寄存器能够被用作 32 个 64-bit 寄存器。这些寄存器可被当作同等数据类型的 vector 使用此时数据是对齐的、元素格式也是相同的因此可以用一个进程对多个数据进行计算一个寄存器位宽 128 bit可存放 4 个 32-bit 元素。相关代码如下//对四个数据同时进行乘法操作 C[0:3] A[0:3]*B[0:3] //一个寄存器 128bit可以存放 4x32bits15 寄存器存放向量 B vldmia.32 r0!, {s15} //通过 s14 寄存器存放向量 A vldmia.32 r1!, {s14} // s15 s15*s14 vmul.f32 s15, s15, s14 //保存 s15 的计算结果 vstmia.32 r2!, {s15}综上SIMD 最重要且最本质的改变是增加了硬件计算单元的数量与数据读取通路的数量同时对上层提供更多指令集。在实际编程中程序员很少会直接操作 SIMD 指令通常由编译器自动向量化完成。SIMT 计算本质单指令多线程的 GPU 并行模型SIMTSingle Instruction Multiple Threads单指令多线程是英伟达NVIDIA提出的、基于 GPU 的新概念。与 SIMD 相比二者都通过将同样的指令广播给多个执行单元来实现数据并行和计算主要的不同在于SIMD 要求所有向量元素在统一的同步组里一个线程内同步执行而 SIMT 允许多个线程在一个 Warp 中独立执行。SIMT 类似 CPU 上的多线程有多个计算核心每一个核心中有独立的寄存器文件Register FileRF和计算单元Arithmetic Logic UnitALU但没有独立的指令缓存Instruction Cache、解码器和程序计数器PC命令从统一的指令缓存广播给多个 SIMT 核心。因此 SIMT 的所有核心各自独立在不同的数据上执行相同的计算操作——即执行命令相同、多个线程各有各的处理单元而 SIMD 则是共用同一个 ALU。还是以数组相乘 $C[0: 3] A[0: 3] × B[0: 3]$ 为例SIMT 给每个元素分配一个线程一个线程只需完成一个元素的乘法所有线程并行执行完成后两个数组的相乘就完成了。SIMT 硬件结构SIMT Core Cluster具体到 SIMT 的硬件结构SIMT 提供一个多核系统SIMT Core ClusterCPU 负责将算子Kernel加载到 SIMT Core Cluster 中。每个 SIMT 核SIMT Core有独立的 RFRegister File、ALU、Data Cache但只有一个程序计数器Program Counter和一个指令译码器指令被同时广播给所有的 SIMT 核执行具体计算。GPU 由多个 SIMT Core Cluster 组成每个 SIMT Core Cluster 由多个 SIMT Core 构成SIMT Core 中运行着多个 Thread Block。从图中可以看到完整的 SIMT 硬件组成外部 CPU 通过 Kernel Launch 将任务加载进来SIMT Core Cluster 内部包含 SIMT Stacks线程栈、Shared Memory、多个 Thread Block、Register File、Constant Cache、Texture Cache、Data Cache 与 Memory Port集群通过互联网络Network与内存分区Memory Partition含原子操作单元、末级缓存组与片外 DRAM交互数据。SIMT 硬件核心流水三个调度循环GPU 的 SIMT 可以看作一个特殊的 SIMD 结构SIMT 硬件核心流水可被分为 SIMT 前端SIMT front-end和 SIMD 后端SIMD back-end。流水线中存在三个调度循环取指循环包含 Fetch取指、I-Cache指令缓存、Decode译码和 I-Buffer指令缓冲四个阶段指令发射循环包含 I-Buffer、Score Board记分板、Issue发射和 SIMT-StackSIMT 栈四个阶段寄存器访问循环包含 Operand Collector操作数收集器、ALU 和 Memory 三个阶段。流水线中的三个调度循环共同组成 SIMT 硬件核心流水其中取指是将具体指令放在堆栈SIMT Stack中堆栈在运行时就会把所有的线程分发到具体的 ALU 中具体执行时采用 SIMD 的方式而 SIMT 主要完成具体线程的前端控制。换句话说SIMT 负责线程的前端控制SIMD 负责数据的后端执行这正是SIMT 前端 SIMD 后端流水线设计的精髓。SIMD 与 SIMT 的区别和联系结合上述内容SIMD 和 SIMT 的主要区别和联系可归纳如下SIMT 与 SIMD 的基本原理相同都是采用单指令多数据的思想SIMT 形式上是多线程但本质上在硬件端执行的还是单线程使用多个核心来实现多线程并行SIMT 比 SIMD 更灵活允许一条指令对数据分开寻址可以实现每个线程独立寻址SIMD 必须连续取址要求数据在类型、格式和大小方面严格对齐。因此SIMT 是 SIMD 的一种推广在编程模式上更加灵活对开发者更友好。NVIDIA CUDA 实现从 Grid、Block 到 Warp 的线程分级回顾 GPU 的线程分级在图形图像处理中会将图像进行切分网格Grid表示要执行的任务大的网格会被分成多个小的网格每个网格中包含很多相同线程Threads数量的块Blocks。此时线程分层执行块中的线程独立执行对像素数据进行处理和计算可以共享数据、同步数据交换。CUDA 并行编程模型基于单程序多数据Single Program Multiple DataSPMD模式SPMD 与 SIMT 的联系和区别在 03SPMT.md 中重点讲解。在 CUDA 编程中grid是线程块block的阵列集合线程块映射到 SM 上进行计算处理一个线程块可包含多个线程束Warp线程块的大小影响 CUDA Kernel 程序的性能GPU 执行时的最小单位是线程thread一个 block 中的线程可存取同一块共享内存并且可以快速进行同步。与 SIMD 不同的是SIMT 允许程序员为独立、标量线程编写线程级的并行代码也允许为协同线程编写数据并行代码。为了确保正确性开发者甚至可以忽略 SIMT 行为很少需要维护一个 warp 块内的线程分支只需维护相关代码即可获得硬件并行带来的显著性能提升。在一个线程块Thread Block中所有线程执行同一段代码在英伟达 GPU 中这段代码称为Kernel。每一个线程有一个自己的线程索引threadIdx.x用于计算内存地址和执行控制决策每个线程在执行时被分配唯一标识符因此可以通过程序准确控制每一个线程。块索引与线程内索引将多个线程块组合在一起就组成一个 Grid 线程组因此线程块是 SM 的基本调度单元SM 对应具体的硬件单元而线程块是编程抽象出来的概念。因为有多个线程块组合同时存在硬件计算单元在横向和纵向两个维度的排布因此线程索引通常由**块索引Block Index和线程内索引Thread Index Within Block**组成块索引用于标识当前线程所在的块Block线程内索引用于标识当前线程在所属块中的位置。在 CUDA 中使用blockIdx.x和blockDim.x访问块索引和块维度中的 x 分量blockIdx.x当前线程所在块的 x 方向索引。块索引是三维向量包括 x、y、z 三个分量blockDim.x当前块的 x 方向维度大小即该方向上线程数量。块维度同样是三维向量包括 x、y、z 三个分量。通过blockIdx.x和blockDim.x可以方便地获取当前线程所在块的 x 方向索引和当前块在 x 方向上的线程数量从而进行相应的计算和操作。结合threadIdx.x线程的全局一维 ID 通常写作int tid blockIdx.x * blockDim.x threadIdx.x;这一经典公式在仓库示例 cuda1.cu 中也有体现——它将 CPU 上的 for 循环for (int i 0; i10000; i) C[i] A[i] B[i];改写为 GPU Kernelint tid blockDim.x * blockIdx.x threadIdx.x;后每个线程独立取a[tid]、b[tid]并相加写入c[tid]。软件线程与硬件单元的对应关系回顾英伟达 GPU 软件和硬件之间的对应关系线程Thread对应 CUDA Core线程以线程块为单位被分配到 SM 上线程块Block对应 SMSM 维护线程块和线程 IDSM 管理和调度线程执行每个线程块又按照每个 Warp 中共32 个线程执行Warp 是 SM 的调度单位Warp 里的线程执行 SIMDBlock 线程块只在一个 SM 上通过 Warp 进行调度一旦在 SM 上调用了 Block 线程块就会一直保留到执行完 KernelSM 可以同时保存多个 Block 线程块块间并行执行。从图中可以清晰看到软硬件层级映射软件侧 Thread → Thread Block → Grid 分别对应硬件侧 CUDA Core → SM → Device完整 GPU 设备。这正是线程对应 CUDA Core、线程块对应 SM、网格对应 GPU的直观表达。需要补充的是Warp 之所以成为调度单位是因为SIMT 架构通过细粒度多线程Fine-Grained Multi-ThreadingFGMT调度将执行流水线细分使不同线程的指令交错执行从而以访存和计算并行来隐藏延迟详见 04NVSIMT.mdWarp 支持乱序执行以隐藏访存延迟线程寄存器值都保留在 RF 中英伟达通过 Warp Scheduler 硬件调度使 Warp 先访存完毕后交给 SIMD Pipeline 执行尽可能多的指令同时隐藏其它 Warp 的访存时间。实战CUDA 矩阵乘 Kernel在 AI 框架的开发流程方面首先会按照编程思想定义神经网络然后根据 AI 框架编写对应的程序AI 框架会自动构建计算正向图并根据自动微分原理构建反向图。其中在神经网络中比较重要的算子是矩阵乘。以 CUDA 代码为例实现 $C A × B$使用blockIdx.x和blockDim.x来访问块索引和块维度#include stdio.h #define N 4 // 矩阵大小 // 矩阵乘法的 CUDA 核函数 __global__ void matrixMultiplication(int *a, int *b, int *c) { // 使用 blockIdx.x 和 blockDim.x 来访问块索引和块维度 int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; int sum 0; for (int k 0; k N; k) { sum a[row * N k] * b[k * N col]; } c[row * N col] sum; } int main() { int a[N][N], b[N][N], c[N][N]; int *dev_a, *dev_b, *dev_c; // 分配内存 cudaMalloc((void**)dev_a, N * N * sizeof(int)); cudaMalloc((void**)dev_b, N * N * sizeof(int)); cudaMalloc((void**)dev_c, N * N * sizeof(int)); // 初始化矩阵 a 和 b for (int i 0; i N; i) { for (int j 0; j N; j) { a[i][j] i * N j; b[i][j] j * N i; } } // 将矩阵 a 和 b 传输到设备 cudaMemcpy(dev_a, a, N * N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * N * sizeof(int), cudaMemcpyHostToDevice); // 定义块大小和网格大小 dim3 blockSize(2, 2); dim3 gridSize(N / blockSize.x, N / blockSize.y); // 调用核函数 matrixMultiplicationgridSize, blockSize(dev_a, dev_b, dev_c); // 将结果传回主机 cudaMemcpy(c, dev_c, N * N * sizeof(int), cudaMemcpyDeviceToHost); // 打印结果 for (int i 0; i N; i) { for (int j 0; j N; j) { printf(%d , c[i][j]); } printf(\n); } // 释放内存 cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; }这段代码的要点在于线程索引计算row blockIdx.y * blockDim.y threadIdx.ycol blockIdx.x * blockDim.x threadIdx.x通过块索引、块维度与线程内索引的组合精确映射到输出矩阵的每个元素位置二维网格配置dim3 blockSize(2, 2)定义每个块为 2×2 共 4 个线程dim3 gridSize(N / blockSize.x, N / blockSize.y)定义网格为 2×2 共 4 个块恰好覆盖 4×4 矩阵的全部 16 个元素数据搬运流程cudaMalloc分配设备内存 →cudaMemcpyHostToDevice上传输入 → 启动 Kernel →cudaMemcpyDeviceToHost取回结果 →cudaFree释放这是 CUDA 编程的标准生命周期。仓库中的 cuda2.cu 与 cuda3.cu 展示了矩阵加法的CPU 串行版 → CUDA 并行版演进CPU 版本用双重 for 循环按index i j * N逐元素相加CUDA 版本则通过int i blockIdx.x * blockDim.x threadIdx.x;、int j blockIdx.y * blockDim.y threadIdx.y;让每个线程独立计算一个(i, j)位置的元素并用if (i N j N)做边界检查最后以add_matrixdimGrad, dimBlock(a, b, c, N)启动。这个对比直观展示了从 SISD 串行到 SIMT 并行的编程范式转变。从编程模型到硬件执行Kernel 到 Warp 的映射从更宏观的角度看CUDA 编程模型与 NVIDIA 硬件的关系可以总结为一条链路程序员编写 SPMD 风格的 Kernel → 编译器将线程逻辑编译为标量指令流 → 硬件把执行相同指令的线程动态组织成 Warp每 Warp 32 线程→ Warp 作为 SM 调度单位交给 SIMD Pipeline 执行。正如 03SPMT.md 所归纳的SIMD是指令的执行方式和对应映射的硬件体系结构SIMT是以 SIMD 指令为主、具有 Warp Scheduler 等硬件模块、支持 SPMD 编程模型的硬件架构SPMD是一种具体的并行编程模型类似于 CUDA 所提供的编程模式DSA是特殊的硬件架构NPU/TPU 等专门针对 AI 的特殊硬件架构应用于大规模数据处理、分布式存储等场景。值得注意的是NVIDIA 在 GPU 架构设计中加入了专门用于神经网络矩阵计算、同时支持混合精度计算的 Tensor Core因此 NVIDIA GPU 也变成了SIMT DSA的模式。此外Warp 的调度特性使得 SIMT 相比 SIMD 在可编程性上有根本性优势硬件层面解决了大部分流水编排问题开发者无需像 SIMD 那样费时费力地把数据凑成合适的矢量长度再传入硬件线程还可以独立执行、允许每个线程有不同的分支——这正是 SIMT 的核心。编程模型 vs 硬件执行模型在理解 SIMD 与 SIMT 之后还需要区分两个经常被混淆的概念编程模型Programming Model是程序员用来编写程序的抽象概念定义了程序员如何组织和控制计算机程序的方式。它提供一种简化的视图使程序员能够专注于程序的逻辑结构而不必考虑底层硬件细节通常包括编程语言、数据结构、算法和并发机制等方面用于描述程序的行为和交互。硬件执行模型Hardware Execution Model描述计算机硬件如何执行程序包括硬件结构、指令集架构、寄存器、内存层次结构、缓存、并行执行方式等方面。它决定了程序在计算机硬件上的实际执行方式——指令的执行顺序、数据的传输方式、并发执行的策略等硬件执行 SIMD 和 SIMT。二者的区别和联系如下区别编程模型从程序员的角度描述程序的组织和行为硬件执行模型从计算机硬件的角度描述程序的执行方式。编程模型关注程序的逻辑结构和抽象行为硬件执行模型关注程序在实际硬件上的执行细节。联系编程模型定义了程序的行为和交互方式硬件执行模型决定了程序如何在硬件上执行。程序员编写的程序最终会被映射到硬件执行模型上执行。理解二者关系可以帮助程序员优化程序性能、充分利用硬件资源。编程模型最终会通过编译器转换为硬件执行模型因此二者在概念层面有明显的差异。这也正是 AISystem 知识体系中AI 编译器环节的切入点编译器要做的就是把 CUDA/SPMD 这类编程模型翻译成目标硬件GPU/NPU能够高效执行的 SIMD/SIMT 指令流。小结与思考SIMD 与 SIMT 均基于单指令多数据执行但 SIMD 要求数据严格对齐且同步执行而 SIMT 允许线程独立寻址且可异步执行提高了灵活性NVIDIA CUDA 通过 SIMT 架构实现高效的并行计算利用线程块Block和网格Grid结构通过 CUDA Core 进行调度优化了 GPU 的性能编程模型与硬件执行模型相互关联前者为程序员提供抽象概念以组织程序后者描述程序在硬件上的实际执行方式理解二者关系有助于程序性能优化从实现层面看SIMT 的本质是SIMT 前端做线程控制、SIMD 后端做数据执行而 CUDA 的blockIdx/blockDim/threadIdx索引体系将线程逻辑精确映射到硬件执行单元这是理解 AI 芯片算子实现如矩阵乘、卷积与 AI 编译器后端优化的关键基础。如果想继续深入可以接着阅读本系列课程的 03SPMT.mdSIMD/SIMT/SPMD/DSA 的完整关系与 Warp 调度细节和 04NVSIMT.mdCUDA 与 NVIDIA 硬件架构、AMD ROCm 对比并结合 code 目录中的 CUDA 示例动手实践。赞分享文档教程人工智能【免费下载链接】AISystemAISystem 主要是指AI系统包括AI芯片、AI编译器、AI推理和训练框架等AI全栈底层技术项目地址https://gitcode.com/GitHub_Trending/ai/AISystem点击查看免费下载相关推荐Apache RocketMQ Controller 模式部署与升级实战指南Apache RocketMQ Controller 模式部署与升级实战指南 本篇指南以 RocketMQ 自动主从切换Controller 模式为主题完文档教程人工智能PyPTO Pro 编程范式深度解析SIMD 与 SIMT 并行执行模型及 AI Core 硬件基础PyPTO Pro 编程范式深度解析SIMD 与 SIMT 并行执行模型及 AI Core 硬件基础 PyPTO ProParallel Tensor/Ti人工智能编译器模型编译深度学习高性能计算CANNAscendPyPTO Pro 编程范式详解SIMD 与 SIMT 并行模型及 AI Core 硬件抽象PyPTO Pro 编程范式详解SIMD 与 SIMT 并行模型及 AI Core 硬件抽象 PyPTO Pro 是 CANN/PyPTO 项目中面向 NPU人工智能编译器模型编译深度学习高性能计算CANNAscend上一篇MinecraftByExample网络通信教程让你的模组实现跨客户端交互下一篇VisiData 列操作完全指南类型、格式化、拆分与派生列实战创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考