在 GPU 架构与底层算子开发如 CUDA、TileLang、CuTe、Triton中Global全局内存、Shared共享内存和Fragment寄存器片段代表了 GPU 金字塔式存储层级中的三个核心层级。它们在容量、访问延迟以及作用域上逐级递减共同决定了计算与访存的流水线设计。存储层级物理位置容量规模访问延迟作用域 (Scope)典型用途Global Memory板载显存 (HBM / GDDR)数十 GB ~ 上百 GB数百个时钟周期整个 GPU (Grid 内所有线程均可访问)存储模型的原始输入、权重以及输出张量Shared Memory片上 SRAM (SM 内部)几十 KB ~ 几百 KB / SM几十个时钟周期单个 Thread Block (块内线程共享)存放 Tiling 切片数据实现块内线程协同与数据复用Fragment片上寄存器 (Register File)每个线程几十 KB 寄存器1~几个时钟周期单个 Thread / Warp 私有存放 Tensor Core (MMA) 计算所需的矩阵小块与累加器1. Global Memory全局内存特性位于 SM 外部的板载 DRAM如 HBM3e、GDDR6容量极大但访问延迟高。算子中的角色数据的起点与终点。算子启动时从 Global 读取输入张量计算完成后再写回 Global。优化重点由于访问开销大必须通过分块Tiling尽量减少对 Global 的重复读取并保证合并访存Coalesced Access或使用硬件异步拷贝指令如 CUDAcp.async/ TMA。2. Shared Memory共享内存特性位于 SM 内部的高速 SRAM由同一个 Thread Block 内的所有线程共同共享。算子中的角色作为数据中转站Staging Buffer。在 GEMM 或 Attention 算子中程序会把 Global 中的大矩阵切成一小块一小块Tile存入 Shared 中供 Block 内的 Warp 反复读取计算将 Global 的访存次数降低数十倍。优化重点必须避免 Bank ConflictBank 冲突现代算子通常结合Swizzle或 Padding 技术来打散物理存储地址。3. Fragment寄存器片段特性分布在线程私有的寄存器堆Register File中是距离计算单元Tensor Core / ALU最近、速度最快的存储区域。算子中的角色Tensor Core 执行 MMA 指令时的直接操作对象。一条硬件矩阵乘法指令如16×8×1616 \times 8 \times 1616×8×16的 GEMM通常需要一个 Warp 中的 32 个线程协同完成每个线程持有一小段寄存器数据即 Fragment共同拼接出一个完整的数据块。典型的数据流转流水线Data Pipeline以 TileLang / CUDA 中的 GEMM 计算为例数据在三者之间的流动过程如下Global Memory→异步搬运 (cp.async / TMA)Shared Memory→寄存器加载 (ldmatrix)Fragment→Tensor Core 计算 (MMA)Fragment (累加器)\text{Global Memory} \xrightarrow{\text{异步搬运 (cp.async / TMA)}} \text{Shared Memory} \xrightarrow{\text{寄存器加载 (ldmatrix)}} \text{Fragment} \xrightarrow{\text{Tensor Core 计算 (MMA)}} \text{Fragment (累加器)}Global Memory异步搬运(cp.async / TMA)Shared Memory寄存器加载(ldmatrix)FragmentTensor Core计算(MMA)Fragment (累加器) global/shared/fragment
在 GPU 架构与底层算子开发如 CUDA、TileLang、CuTe、Triton中Global全局内存、Shared共享内存和Fragment寄存器片段代表了 GPU 金字塔式存储层级中的三个核心层级。它们在容量、访问延迟以及作用域上逐级递减共同决定了计算与访存的流水线设计。存储层级物理位置容量规模访问延迟作用域 (Scope)典型用途Global Memory板载显存 (HBM / GDDR)数十 GB ~ 上百 GB数百个时钟周期整个 GPU (Grid 内所有线程均可访问)存储模型的原始输入、权重以及输出张量Shared Memory片上 SRAM (SM 内部)几十 KB ~ 几百 KB / SM几十个时钟周期单个 Thread Block (块内线程共享)存放 Tiling 切片数据实现块内线程协同与数据复用Fragment片上寄存器 (Register File)每个线程几十 KB 寄存器1~几个时钟周期单个 Thread / Warp 私有存放 Tensor Core (MMA) 计算所需的矩阵小块与累加器1. Global Memory全局内存特性位于 SM 外部的板载 DRAM如 HBM3e、GDDR6容量极大但访问延迟高。算子中的角色数据的起点与终点。算子启动时从 Global 读取输入张量计算完成后再写回 Global。优化重点由于访问开销大必须通过分块Tiling尽量减少对 Global 的重复读取并保证合并访存Coalesced Access或使用硬件异步拷贝指令如 CUDAcp.async/ TMA。2. Shared Memory共享内存特性位于 SM 内部的高速 SRAM由同一个 Thread Block 内的所有线程共同共享。算子中的角色作为数据中转站Staging Buffer。在 GEMM 或 Attention 算子中程序会把 Global 中的大矩阵切成一小块一小块Tile存入 Shared 中供 Block 内的 Warp 反复读取计算将 Global 的访存次数降低数十倍。优化重点必须避免 Bank ConflictBank 冲突现代算子通常结合Swizzle或 Padding 技术来打散物理存储地址。3. Fragment寄存器片段特性分布在线程私有的寄存器堆Register File中是距离计算单元Tensor Core / ALU最近、速度最快的存储区域。算子中的角色Tensor Core 执行 MMA 指令时的直接操作对象。一条硬件矩阵乘法指令如16×8×1616 \times 8 \times 1616×8×16的 GEMM通常需要一个 Warp 中的 32 个线程协同完成每个线程持有一小段寄存器数据即 Fragment共同拼接出一个完整的数据块。典型的数据流转流水线Data Pipeline以 TileLang / CUDA 中的 GEMM 计算为例数据在三者之间的流动过程如下Global Memory→异步搬运 (cp.async / TMA)Shared Memory→寄存器加载 (ldmatrix)Fragment→Tensor Core 计算 (MMA)Fragment (累加器)\text{Global Memory} \xrightarrow{\text{异步搬运 (cp.async / TMA)}} \text{Shared Memory} \xrightarrow{\text{寄存器加载 (ldmatrix)}} \text{Fragment} \xrightarrow{\text{Tensor Core 计算 (MMA)}} \text{Fragment (累加器)}Global Memory异步搬运(cp.async / TMA)Shared Memory寄存器加载(ldmatrix)FragmentTensor Core计算(MMA)Fragment (累加器)