英伟达（NVIDIA）过去几代芯片的进化，本质上就是一场为矩阵乘法疯狂献祭的仪式。他们不是在造GPU，他们是在为这个最基础的数学操作打造一座神殿。最新的B200芯片，就是这场献祭的最新祭品。
首先，Tensor Core（张量核心）——这不是普通的计算单元，而是专门为矩阵乘法而生的协处理器。你可以把它理解成一个只干一件事的“特种兵”：吃进矩阵，吐出结果。而且B200这一代，张量核心能处理的“块”（tile）比以前更大，意味着一次能吞下更多数据，效率直接拉满。
但光有张量核心还不够。你算得再快，如果数据搬不动，那也是白搭。所以英伟达又搞了个新东西：Tensor Memory（张量内存）——一种专门为张量核心服务的高速缓存，专门用来存那些中间计算结果。以前这些数据可能得在普通缓存里绕来绕去，现在直接有VIP通道，省时省力。
更狠的是Tensor Memory Accelerator（TMA，张量内存加速器）。这玩意儿从H100那一代就开始登场，专门负责在后台悄悄把数据搬进搬出，完全不打扰张量核心干活。就像你炒菜时，有个隐形助手默默把下一道菜的食材提前切好、摆好，你根本不用分心——这就是异步内存搬运的魔力。
但硬件只是半边天。真正让这套系统跑起来的，是背后那一整套极其复杂的软件栈和抽象层。CUDA、cuBLAS、cuDNN、Transformer Engine……这些名字听起来像咒语，其实是工程师们一层层堆出来的“调度魔法”，目的只有一个：让成千上万个张量核心在正确的时间、用正确的数据、干正确的事，还不互相打架。
这已经不是单纯的工程问题了，这是“硅片”和“抽象”的终极战场。英伟达的芯片迭代，从来不是东一榔头西一棒槌，而是围绕一个核心目标反复打磨——让矩阵乘法更快、更省电、更高效。他们不是在追求数字上的提升，而是在践行一种哲学：速度即信仰。
为什么一个如此简单的操作，值得如此疯狂的投入？因为杠杆效应。矩阵乘法是深度学习的基石，Transformer、卷积神经网络、推荐系统……底层全是它。你优化1%的矩阵乘法速度，可能带来整个AI训练时间10%的缩短。这种杠杆，是指数级的回报。
所以，别再觉得矩阵乘法只是课本里的老古董了。它早已成为AI时代最核心的“原子操作”。英伟达不是在卖芯片，他们是在卖“加速现实的能力”。而这一切，都始于那个看似平平无奇的——A × B = C。
说到这儿，不得不提一下背后的推手。英伟达的首席科学家Bill Dally，是计算机体系结构领域的泰斗级人物，斯坦福教授出身，几十年深耕高性能计算。正是他带领团队把张量核心从概念变成现实，并不断迭代。而CEO黄仁勋（Jensen Huang）则以“精准押注AI”闻名，早在2010年代就断言深度学习将重塑计算范式。如今回头看，英伟达的每一步，都是对这一判断的坚决执行。
在这场硬件与算法的共舞中，矩阵乘法成了那个最朴素也最强大的支点。阿基米德说：“给我一个支点，我能撬动地球。”英伟达说：“给我一个矩阵乘法，我能训练出下一个GPT。”
从极简到极繁：一块硅片如何为矩阵乘法掀起万年战争 （故事版 ）
本文用故事笔法拆解 NVIDIA 从 Pascal 到 Blackwell 的架构史诗：把小学二年级的“一行乘一列”供奉成吞电巨兽，为此祭出 Tensor Core、TMA、专属缓存等三圣器，终让 4 平方厘米裸片变成撬动 AI 杠杆的行星发动机。
第一章：极简原罪的降临
——————————
很久很久以前，在数字大陆还是一片荒原的时候，矩阵乘法不过是一张羊皮纸上的咒语：左边一行，右边一列，对应数字牵手相乘，再挽手相加，如此而已。
谁也没想到，这条咒语在千年后被一群硅谷巫师捡了起来，他们把它刻进硅片，喂以电流，竟唤醒了名为“深度学习”的巨龙。
巨龙胃口巨大，每吞下一幅图像就要执行数十亿次那张羊皮纸上的咒语；每生成一句人话，又要再念数十亿遍。
于是，原本“极简”的原罪，瞬间化作“极繁”的灾难：时钟滴答一次，就得完成几十万次乘加，而留给一次滴答的时间只有零点几纳秒。
silicon 的信徒们抬头望天，发出灵魂拷问：我们该如何在物理极限的悬崖边，继续让这条咒语加速？
——————————
第二章：第一滴血——Tensor Core 的成人礼
——————————
2017 年，帕斯卡老将军退役，伏特新王登基，王宫深处一扇从未开启的暗门被推开，里面传来齿轮咬合的巨响。
朝臣们窃窃私语：那是“Tensor Core”，一种只为矩阵乘法而活的异形算子。
它不像传统的 CUDA Core 那样娇生惯养地一次只肯算一对浮点数；它一挥手就是 4×4 的矩阵块，16 对数字同时拥抱相乘，再加一棵加法树，一次吞吐就是 64 个 FMA。
朝堂哗然：原来“并行”二字还能这么写？
但新王冷酷地宣布：想召唤 Tensor Core，你得先学会新咒语——WMMA、mma、mma.sync，还要把数据排成“张量块”的奇形阵，否则它便沉睡不醒。
于是，编译器骑士团连夜修书，写就《CUDA 张量炼金术》第一卷；框架诸侯们则把 PyTorch、TensorFlow 的底层改得血肉模糊，只为在恰当的时辰献上正确的数据祭品。
第一滴血落下，硅片世界自此分裂为“会喂 Tensor Core”的贵族和“只会老 CUDA”的平民。
——————————
第三章：内存饥荒与“专用缓存”圣杯
——————————
Tensor Core 胃口大开后，新的饥荒接踵而至：算得越快，饿得越快，数据供给线却细若游丝。
老臣 L2 缓存跪在地上哭诉：臣妾做不到啊，臣要伺候 SM、纹理单元、光线追踪、指令缓存……实在分身乏术！
宫廷会议连夜召开，老臣被勒令“让位”，一块全新的封地——“Tensor Memory”被划了出来。
它藏在 SM 的心脏地带，容量不大，却拥有独立端口、独立调度器，只为 Tensor Core 的临时结果提供“随叫随到”的粮仓。
史书称此举为“私有缓存圣杯化”：把通用资源拆成专宠，是奢侈，却也是继续加速的唯一活路。
代价则是更大的芯片面积、更复杂的连贯性协议，以及驱动软件里长达数千行的“一致性咒语”。
但巫师们不在乎，他们在实验室的黑板上写下血书：Speed of Light or Die。
——————————
第四章：异步狂想曲——TMA 的时空穿梭
——————————
时间来到 2022 年， Hopper 王朝。
Tensor Core 已进化到第四代，块尺寸从 4×4×4 一路膨胀到 8×8×16，一次咀嚼就是 1024 次乘加，算力怪兽的喉咙更粗，可内存粮仓的马车却依旧在羊肠小道上颠簸。
老问题以更狰狞的面目归来：数据搬运开始吃掉 90% 的能耗，而真正计算只剩 10%，世界仿佛回到“搬运一条绳子上的大象”的寓言。
于是，宫廷炼金术士们祭出第二件圣器——Tensor Memory Accelerator，简称 TMA。
它住在芯片边缘，拥有独立的地址生成器、步幅计算器和异步 DMA 引擎，可以在计算进行的同时，把下一批矩阵块从全球内存预载到 Tensor Memory，再把旧结果写回。
形象地说，TMA 就是一位时空穿梭管家：主人在客厅开派对，它已在地下室把下一瓶香槟冰镇好，并把喝完的空杯悄悄收走。
软件世界因此掀起“异步编程”的腥风血雨：程序员要手动把循环拆成 stage、double buffer、pipeline，还要在 CUDA 12 的语法里写下 cp.async.bulk.tensor 这种长达 40 个字符的咒语。
任何一条写错，管家就会摔碎杯子，派对当场爆炸——芯片死锁，算力归零。
但正确书写的人，将获得 1.8 倍到 2.3 倍不等的“免费”性能，仿佛捡到物理定律的漏洞。
——————————
第五章：B200——黑王登基与四重炼狱
——————————
2024 年，Blackwell 家族诞生，B200 作为嫡长子被推到龙椅。
史官翻开《硅片年表》，只见一行赤红大字：
“本代不再提供裸算力翻倍，而是把‘搬运地狱’彻底炼成‘搬运流水线’。”
B200 的 Tensor Core 升级到第五代，支持 16×8×32 的巨型 tile，一次吞吐 4096 个乘加；
Tensor Memory 翻倍，且被切成四等份，每份配独立 TMA 通道，形成“四重奏”搬运；
新增“Tensor Memory Compression”单元，可在搬运途中把 FP32 结果实时压成 FP16 或 BF16，省下一半带宽；
全局内存接口升级到 HBM3e，频率飙到 8.8 Gbps，可仍被皇帝斥责为“太慢”。
于是，软件栈再被拖进更深的炼狱：
——现在你要把矩阵拆成四级 hierarchy：global→TMA→Tensor Memory→寄存器；
——你要计算 tile 大小，使得压缩前后 byte 对齐；
——你还要在 kernel 里写 grid_constant 修饰，让编译器把某些参数钉在特殊常量缓存，避免 TMA 地址计算再拖后腿。
任何一个维度没对齐，性能就像被拔了氧气管，瞬间青紫。
但成功调通的人，会在 profiler 里看到一条漂亮到令人窒息的流水线：计算占用 98%，内存延迟被完全盖住，芯片功耗顶格 700 W，仿佛一列 8.8 GHz 的货运列车贴着物理墙壁呼啸而过，没有一分一毫浪费。
——————————
第六章： silicon 上的形而上学
——————————
故事讲到这里，你终于发现：矩阵乘法之所以“简单”，正因为它“纯粹”。
纯粹到只剩“乘”与“加”，于是任何一丁点冗余都能被显微镜放大成十倍的功耗、百倍的延迟。
NVIDIA 的芯片迭代不是简单的“加核、提频”，而是一场持续十年的“去冗余”修行：
把通用缓存劈成专宠，是砍冗余；
把同步搬运做成异步，是砍冗余；
把地址计算硬化成固定函数，是砍冗余；
甚至把数值精度动态压缩，也是砍冗余。
冗余越少，芯片就越像一柄薄如蝉翼的手术刀，刀口对准的，正是“杠杆”二字。
因为 AI 模型就是一个杠杆：参数是支点，数据是力臂，矩阵乘法则是那根无影的杠杆本身。
杠杆本身重量若减一克，撬起的智能就能翻一倍。
于是，NV 的每一代架构，都在向这根杠杆献祭：更大的面积、更热的功耗、更复杂的编程模型，只为让“乘加”二字趋近于零延迟的哲学极限。
这是一场没有终点的朝圣——
当 TF32 被搬出来，人们说“精度够用了”；
当 FP8 被搬出来，人们又说“精度够用了”；
未来也许出现 FP4、FP1，甚至 1bit 乘加，只要杠杆还能撬得动世界， silicon 的祭司就将继续挥刀自宫，把冗余削到量子涨落为止。