英伟达Blackwell张量核心真相：矩阵乘法竟变内存泄漏重灾区
内容提要
英伟达Blackwell架构将张量核心升级为独立协处理器，配备专属内存，矩阵乘法可由单线程异步发射，结果存入专用空间，无需线程集体等待。但需手动管理内存，易致泄漏，且读取结果需128线程。此设计彻底改变CUDA编程范式，从集体协作转向流水线分工，凸显张量核心从指令到独立单元的转变。
延伸解读
编程范式之变：从集体协作到流水线分工
Blackwell将张量核心升级为独立协处理器，矩阵乘法发射从128线程缩至1线程，结果读取却需128线程。这种不对称设计迫使开发者重构内核：六个线程束各司其职，通过硬件邮箱协作，彻底告别传统SIMT模式。理解这一转变，是优化Blackwell性能的关键。
内存管理新挑战：矩阵乘法也会泄漏
张量核心专属内存需手动分配和释放，忘记free会导致后续线程块卡死。这种内存泄漏隐蔽且致命，普通工具难以检测。NVIDIA虽提供sanitizer，但开发者仍需养成严格的内存管理习惯，否则在长期运行或复杂任务中极易触发问题。
寄存器瓶颈催生架构革新
寄存器文件容量七年未变，而张量核心算力增长十八倍，导致分块数据占满寄存器。Blackwell通过为张量核心配备独立内存，绕开这一瓶颈，但也带来新的编程复杂度。这一设计权衡反映了GPU架构在算力与存储间的艰难平衡。
Q&A
英伟达Blackwell架构中，张量核心的独立内存大小是多少？
在Blackwell架构中，每个流式多处理器（SM）上为张量核心划出了256KB的独立内存，具有独立的地址空间，只有张量核心能读写。
在Blackwell上，发射一个矩阵乘法需要多少个线程？
在Blackwell上，发射一个矩阵乘法只需要一个线程，因为结果直接存入张量核心的专属内存，无需线程集体等待。
为什么Blackwell上读取矩阵乘法的结果需要128个线程？
因为张量核心内存被切成了四等份，每个线程束（warp）只能看到自己的一份，一个线程束有32个线程，四份就需要128个线程（一个warpgroup）来读取结果。
Blackwell上矩阵乘法如何导致内存泄漏？
在Blackwell上，矩阵乘法的结果存储在张量核心的专属内存中，需要手动分配和释放。如果线程算完矩阵乘法后忘记free这块内存，就会导致内存泄漏，下一个被调度到同一SM上的线程块会卡死。
NVIDIA为Blackwell提供了什么工具来检测矩阵乘法的内存泄漏？
NVIDIA提供了一个sanitizer（检测器），打开开关后，运行时一旦发现有内存未释放，会立刻trap住，并精确指示泄漏的位置。
Blackwell上高速核函数（kernel）的线程束是如何分工的？
在Blackwell上，一个高速核函数通常由六个线程束组成流水线：第一个线程束负责把数据搬进张量核心内存，第二个线程束专门发射tcgen05.mma指令，剩下四个线程束负责把结果从张量核心内存中drain出来。它们通过硬件邮箱通信。
Blackwell架构中张量核心的角色发生了什么根本变化？
张量核心从线程束中的一条指令变成了一个独立的协处理器，拥有自己的内存空间、指令集和执行节奏。程序员需要手动分配内存、派发任务和清理现场，它不再是被集体调用的功能，而是独立运转的单元。
