营造司的掌班人

一、三千工匠，乱成一锅粥

京城要建天下粮仓。十万石储量，工期三个月。工部一道急令，从直隶三十六县调了三千工匠，黑压压地涌进了城东那片空地。

头七天过去，督造官卢大人来巡工。他在工地上站了不到半刻钟，脸就绿了。

东头有二十几个瓦匠在砌墙，泥还没干，西头又有一队人推着车过来，把刚砌好的墙角撞塌了一块。南边堆着三千块青砖没人搬，北边的木匠在等砖，干坐了一上午。还有七八十个工匠散在各处，晒太阳的、抽旱烟的、靠墙打盹的——不是不想干，是根本没告诉他们该干什么。

"三千人，七天，只砌了三面半墙。"卢大人把营造司的掌班鲁端木叫到跟前，"照这个进度，别说三个月，三年也封不了顶。"

鲁端木没吭声。他在工地上蹲了三天，把三千人的动线全看了一遍。他发现问题不是工匠懒，也不是监工不尽力——是这三千人根本没有编组。谁跟谁搭伴、谁取料谁砌墙、哪个队管哪块工区，全凭工匠自己商量。三千人乱成了一锅粥。

"给我三天。"鲁端木站起来，拍了拍膝上的灰，"三天之后，工地上多一个人都不要，但每天砌的墙要比现在多五倍。"

卢大人盯着他看了好一阵："你拿什么保证？"

"我不保证。"鲁端木已经转身朝工棚走去了，"我做给你看。"

二、九十三个队，每队三十二人

鲁端木做的第一件事，是画了一张巨大的编组图。

他把三千人编成了九十三个工队，每队正好三十二人。多出来的二十几个工匠——年纪太大、腿脚不便——编成了补给队，只管烧水送饭，不上一线。

"为什么是三十二？"鲁端木的小徒弟阿岑问他。

"因为一面铜锣只能让三十二个人同时听见。"鲁端木头也不抬，"再多一个，锣声就有人在半拍之后才反应过来。半拍的延迟——一块砖就砌歪了。"

他把编组图贴在工棚外面的木墙上。三千个名字，九十三个队，每个队分到了一块自己的工区——不是整座粮仓，只是粮仓的一小段。第一队砌东墙左起第一段，第二队砌左起第二段，以此类推。九十三个队像九十三个拼图块，各自只管自己面前那一块。

每个工区正中间，鲁端木让人搭了一个矮木台。

"这叫工具台。"他把阿岑带到一台前，"瓦刀、抹子、线锤、水平尺——全队最常用的家伙全搁在这台上。一个人伸手就够得着。谁要用什么，走两步就到，不用去总库房领。"

"那总库房呢？"

"总库房在大门外面。"鲁端木往远处指了一下，阿岑眯着眼找了半天才看见——一个灰扑扑的大棚子，孤零零地蹲在工地最角落里，"青砖、石灰、木料、铁钉全在那儿。但你不能每缺一块砖就跑一趟总库房——来回小半个时辰，一天跑六趟，天就黑了。"

他转过身，拍了拍阿岑的肩膀。

"记住：能伸手的，别走两步。能走两步的，别跑半里地。工具就搁在脚边。材料才去总库房取。"

三、三十二人只听一声锣

鲁端木做的第三件事，是给每个工队配了一面铜锣。

锣是定制的——不大，铜面的，一敲声音又脆又短，刚好传到三十二个人的耳朵里，不会串到隔壁工区。

"一面锣，四个令。"鲁端木把锣递给一个工队长，竖起手指：

"一响——所有人砌砖。二响——所有人抹灰。三响——所有人拉线对齐。四响——所有人放下工具，歇一刻。"

"能不能有第五个令？"工队长问。

"不能。"

"为什么？"

"因为一面锣就那么几种敲法。"鲁端木拿起锣槌，当当地敲了两下，声音清亮利落，"一声、两声、三声、四声——多了就分不清了。分不清就乱。乱了就慢。"

他让这个工队长试了两天，然后做出了一个更细致的安排：同一个工队里，三十二块工位尽量挨着、干的活尽量一样。砌标准砖的都砌标准砖，抹灰的都抹灰，拉线的都拉线。

阿岑不解："为什么非要干的活一样？"

"因为三十二个人在听同一面锣。"鲁端木指着正在砌墙的工队，"锣响一声，三十二个人必须同时砌砖。如果有一个人的位置需要斜切砖——他做不到一响锣就切完。其他三十一个人就得停下来等他。这三十一个人的手是空着的，瓦刀是提着的，时间是一样在走的。"

阿岑盯着那个工队看了好一会儿。看得出来——全队的速度，是被最慢的那个人拖住的。

"那碰到需要斜切的情况呢？"

"不放在这个队里。"鲁端木已经朝下一个工区走去了，"不一样的事，交给不一样的队。一个队只干一种活。活一样，锣就齐。锣齐了，三十二个人才能同时在同一个节拍上动手。"

四、取料取对路，一趟就当十趟

编组和铜锣解决之后，新问题冒出来了。

阿岑发现工匠们还是花太多时间在路上——不是走去总库房那半里地的问题，是到了总库房以后的问题。一个人进去翻半天货架，找不到自己要的那批青砖，回头又叫三个人来帮忙，四个人堵在过道里，后面的全进不去。

"总库房的货架是按品类排的。"阿岑皱着眉头跟鲁端木汇报，"青砖在一号架，石灰在五号架，铁钉在十二号架。一个工队要同时取青砖、石灰和铁钉，就得跑三个货架——"

"那就不要同时取。"鲁端木说。

"可是——"

"我说的不是分时取。"鲁端木从怀里掏出一本小册子，翻到中间一页，上面画着总库房的货架布局，"我说的是——取料的人编号。一号取一号架，二号取二号架，三号取三号架。"

阿辰愣住了。

"不是让一个人跑三个货架。是让取料的三个人，编号相邻，去相邻的货架。一号取完一号架直接递给二号，二号接了一号和自己的料一起递给三号——三个人一趟车拉回去。"

"所以取料的速度不取决于总库房有多大——"

"取决于取料的人是不是相邻编号、取的是不是相邻货架。相邻的人取相邻的料——这是最快的方法。随便取——比这慢十倍。"

阿岑飞快地在纸上算了一下。三千人，每队三十二人，如果每队派三个人取料，取料的人编号挨着，取的货架也挨着——一趟取料能省下四分之三的时间。

他把笔一扔，叫了一声好。

五、手上的家伙够不够使

编组定好了，铜锣挂上了，取料也理顺了。但鲁端木还是不满意。

阿岑发现师傅盯着编组图看了整整一个下午，笔在纸上画了又划，划了又画。

"您在算什么？"

"算每个人手上能拿几样家伙。"鲁端木把纸推过来。纸上列着每个工队的人数、每人随身带的工具数、工具台上放的公用工具数。密密麻麻好几页。

"每个工匠只有两只手。"鲁端木说，"瓦刀占一只手，抹子占另一只，水平尺就只能插在腰带上——不一定顺手。这跟人数多少没关系，是每个人手上的位置就这么多。"

"那就每个人少带几样？"

"少带了，就得频繁跑去工具台上取。工具台虽然近，不如就在手里。多带了随身的东西，占着手，慢。"

"那怎么办？"

"找到一个数。"鲁端木在纸上圈了几个数字，"每个人带三到五件最常用的——瓦刀、抹子、线锤、水平尺——其余的放工具台。三件太少，八件太多。在这个数之间，手上够用，工具台也不挤。

这还没完。他又拉了一组数——每个工队的人数，和三十二比起来，能不能再调？

"三十二个人用一个工具台，刚好。但如果队里的人每个人手上带的工具太多，工具台的格子就不够全队用。那就得减人——三十个、二十八个——让剩下的人手上的家伙和台子上的格子刚好匹配。"

阿岑听出了门道："所以不是固定的队越大越好——"

"对。是手上的家伙和台子上的格子，合在一起刚刚好——才好。人多了，家伙不够分，干等着。人少了，台子空着，浪费。刚刚好——才是最快的。"

鲁端木最后敲定：大部分工队保持三十二人，每个人随身四件工具。有三个人数特别密集的工区，每队减到二十八人。有两个活特别复杂的工区，每队只放二十四个人——因为每个人需要随身带六件工具。

编组图最终定稿那天，阿岑把九十三个队的数据誊写了一遍。他忽然发现，鲁端木从头到尾只做了三件事：让同一个队里的人干一样的活，让取料的人按相邻编号取相邻货架上的料，让每个人手上的工具和台子上的格子刚好对得上。但他记得第一天工地上那片乱象。三千人，七天，三面半墙。

"师父，明天卢大人来巡工。"

鲁端木嗯了一声，把编组图从墙上揭下来，叠好放进了工具箱里。

六、二十七天，粮仓封顶

卢大人如约来了。

但这回他不是来看进度的——他是来通知鲁端木，工期提前了。

"北边边关吃紧，军粮要提前调拨。"卢大人站在工棚里，脸色比上一次还难看，"原定三个月——现在最多给你三十天。从今天算起。"

阿岑站在旁边，差点把手里的茶壶摔了。三十天。上次三千人、七天、三面半墙，卢大人当场差点把鲁端木的乌纱帽摘了。现在同样的三千人，整个粮仓，三十天。

"够不够？"鲁端木忽然问了一句。

卢大人愣住了："你说什么？"

"我问你——三十天够不够。"

"你是掌班，你问我？"

"对，我问你。"鲁端木直视着卢大人，"你定三十天，一定有你的道理。我要知道这个道理是什么。"

卢大人沉默了一会儿："三十天之后，第一批军粮必须入库。过了三十天，前面的仗怎么打，就不是我能管的事了。"

"明白了。"鲁端木站起来，拎起工具箱就往外走，"三十天。一分不多。"

接下来那些日子，阿岑后来回想起来都觉得像一场梦。

九十三个工队的铜锣从日出敲到日落。每一声锣响，三十二双手同时起落——砌砖、抹灰、拉线。三十二声锣在九十三个工区之间此起彼伏，错落有致。取料的人按照编号排成三列，拿着号牌去总库房，相邻货架的料一趟拉回来。谁手里的工具不顺了，走两步到工区中央的木台，放下旧的拿起新的，锣声没断过。

总有意外。第三工区砌到一半，发现地基下面的土是松的——底下得先灌一层石灰浆。换在以前，整队人就得停下来等。但鲁端木的处理很简单：把这个队拆成两个半个队。半个队去灌浆，半个队先砌相邻的墙。进度没耽误。

第十一工区有十一个人同时发现手里的瓦刀不好使——刀刃崩了。鲁端木立刻从补给队调了四把新瓦刀过来，让崩瓦刀的人去工具台拿备用的，其余的人接着干。全队多等了不到半刻钟。

第十二天的时候，阿岑在编组图上画了一笔——工程量过半。

第十九天，西墙、南墙、东墙全部封顶。还剩北墙和房顶。

第二十四天，北墙砌完。

第二十七天，粮仓封顶。

卢大人站在刚封顶的粮仓前面，仰着头看了很久。十万石容量的粮仓——攒尖顶、厚砖墙、十二个通风口齐齐整整——像一座小山蹲在京城东面。三千个工匠已经散了，工地上只剩下风吹青砖的声音。

"上次你说，每天砌的墙要比原来多五倍。"卢大人没回头，"多了几倍？"

"十一倍。"阿岑抢着答了，随即又缩了回去。

"十一倍。"卢大人重复了一遍这三个字，像在咀嚼一块味道很复杂的糖。

鲁端木什么也没说。他蹲在地上，用一根树枝在泥里画了几个圈。

"卢大人，"他抬起头，"你知道为什么十一倍吗？"

"怎么讲？"

"不是因为工匠比以前勤快了——工匠还是那批工匠，力气还是那么多力气。是因为我把三千人的力气，编到了一起。"

他用树枝指着地上的圈：

"每个人手上不超过四件工具。每队人干的活一模一样。取料的人按相邻编号取相邻货架。工队之间互不串扰——一个队慢了不会拖累旁边的队。"

他把树枝放下，站起来。

"快不是多使劲。快是不让力气白花在路上，白花在等人上，白花在找工具上。快是把三千个人的每一分力气，都花在它应该花的地方。"

卢大人低头看了看地上那幅用树枝画出来的编组图。风吹过来，浮土盖住了几个圈，但圈和圈之间的位置关系反而看得更清楚了。

"有道理。"他转身走了几步，又回过头来，"鲁掌班，北边还有三座粮仓要建。工期一样紧张。你愿意去吗？"

鲁端木终于笑了一下。

"我编过的工队不用重编。人去就行。铜锣是随身带的。"

三千人的力气，不在多，在编。编好了，铜锣一响，千手同起——没有比这更快的了。

技术解读

CUDA（Compute Unified Device Architecture）是 NVIDIA 于 2006 年推出的并行计算平台和编程模型，它允许开发者用 C++ 扩展语法直接控制 GPU 上的数千个线程同时执行计算任务。GPU 之所以能成为深度学习革命的硬件基石——从卷积神经网络的像素级并行到 Transformer 大语言模型的矩阵乘法——根本原因就在于其精密的线程层级体系：线程（Thread）、线程束（Warp）、线程块（Block）和网格（Grid），四个层级协同工作，将数万亿次计算拆成可并行的最小单元，同时通过共享内存（Shared Memory）和合并访存（Coalesced Access）大幅压缩数据传输延迟。CUDA 编程的核心不是"怎么写"，而是"怎么编"——编组合理，一个消费级 GPU 就能跑出超算的效率；编组不当，再多的 CUDA Core 也在空转。

核心概念回顾

故事中的隐喻对照

为什么这个故事对应 CUDA/GPU 并行计算？

1. 线程层级是 GPU 并行的核心设计：三千工匠→九十三个队→每队三十二人→每人听一声锣——Grid→Block→Warp→Thread 的四级层级精确映射了 GPU 的硬件调度模型。理解这一层级是写出高效 CUDA 程序的前提：数据怎么分配给线程、线程怎么组织成块——一切性能优化都从这里出发。

2. Warp 是调度的原子单位，SIMT 是执行的根本模式：每队三十二人同时听一声锣做同一件事，这正是 warp 的 SIMT 本质。一个 warp 的 32 个线程在每个时钟周期执行同一指令（PC 共享），只是数据不同（寄存器不同）。GPU 的高吞吐量不是靠每个线程跑得多快，而是靠这种"三十二人同一步调"的规模化并行。

3. Warp 发散是性能杀手：故事中"有人要斜切砖，其他三十一人干等"是对 warp divergence 最直白的解释。当 warp 内线程遇到 if-else 分支而条件不同时，GPU 必须执行完一条路径再执行另一条——执行 if 时走 else 的线程闲置，反之亦然。这就是为什么好的 CUDA 代码尽量让同一 warp 内的线程走相同分支。

4. 内存层级决定了实际性能：GPU 有三层存储——寄存器（私人工具带，0 延迟）、共享内存（工区工具台，20 周期）、全局内存（总库房，~600 周期）。数据放在哪里，直接决定程序快慢。优秀的 CUDA 程序尽量把频繁访问的数据放在共享内存，正如老掌班鲁端木把常用工具放在工区中央的木台上。

5. 合并访存是用好全局内存的前提：相邻线程访问相邻地址——这是 GPU 全局内存的硬件优化核心。一次合并访存可以把 32×4 字节的请求压缩成一次 128 字节宽位传输，速度提升接近 32 倍。故事中"编号相邻的人取相邻货架"正是这一规则的生动呈现。

6. 占用率是资源和速度的博弈：每个 SM 能驻留的 warp 数量受制于寄存器、共享内存等物理资源。线程用的寄存器越多、Block 用的共享内存越大，则 SM 能同时运行的 warp 越少——但隐藏内存延迟的能力也越弱。鲁端木反复计算"手上的工具数和台子的格子数要匹配"便是占用率优化的核心问题。

7. C++ 是 CUDA 的原生语言：CUDA 的核函数用 __global__ 声明、<<<grid, block>>> 启动——语法是 C++ 的扩展，编译工具链（nvcc）基于 LLVM/Clang。现代 C++ 特性（模板、constexpr、lambda）可以直接在设备端使用。正是 C++ 的零开销抽象让 GPU 编程在保持极致性能的同时具备了表达力。

后记：GPU 的真正能力不在几千个 CUDA Core 的纸面数字——在你怎么编。编不好，三千人各干各的，三天砌三面墙。编好了，铜锣一响，千手同起，二十七天封了一座粮仓。鲁端木没发明任何新工具——铜锣是路边买的，工具台是木板钉的，编组图是一张纸。他只做了一件事：让每一个人的力气都花在应该花的地方。快不是多使劲，快是不白费劲。 GPU 如此，世上的许多事也如此。