一句话定义：TMA（Tensor Memory Accelerator，张量内存加速器）是 NVIDIA 从 Hopper 架构开始在 GPU 里集成的专用硬件单元，负责把一整块多维张量数据在显存（HBM/global memory）和片上共享内存（shared memory）之间异步搬来搬去，全程不需要线程逐个元素参与。
以前是怎么搬的
GPU 的计算核心（SM）不能直接对显存里的数据做矩阵乘法，必须先把数据搬到离计算单元更近的共享内存。传统做法是：一个线程块里几百个线程，每人算好自己的地址，用一条 load 指令搬一个或几个元素进寄存器，再写进共享内存。
问题很明显：地址计算要算、指令要发射、数据要占寄存器。数据还没开始算，寄存器已经被"过路的货"挤满了，这就是常说的寄存器压力（register pressure）。而且每次只能搬一维的一小段，遇到二维矩阵的边界还得写一堆判断代码。
TMA 的做法
打个比方：以前是几百个工人排队，每人抱着一个箱子从仓库跑到车间。TMA 相当于请了一支专业搬运队，你只需要填一张"提货单"（在 CUDA 里叫 tensor map，描述数据在哪、一共几维、每维多大步长、要搬多大一块、怎么排布），交给搬运队，它自己就把整块货送到了共享内存，送完发个信号通知你。
于是变化是：一条指令搬一整块多维数据；搬运过程不经过寄存器；搬运期间线程可以继续干别的（异步，配合 mbarrier 之类的完成机制等待）；多出来的边界情况由硬件自动处理，越界部分按规则补零，省掉大量 if 判断；数据落到共享内存时还能顺便做 swizzle 重排，避免 bank 冲突。Blackwell 及后续架构沿用了这套思路，并把它和新的张量核心数据通路配合得更紧（具体硬件细节以官方文档为准）。
和相邻概念的区别
| 谁在搬 | 粒度 | 占寄存器吗 | |
|---|---|---|---|
| 普通 load/store | 每个线程 | 元素/少量元素 | 占 |
cp.async 一类的异步拷贝 | 每个线程发指令 | 每线程一小段 | 不占寄存器，但仍需逐线程算地址 |
| TMA | 专用硬件单元 | 整个多维 tile | 不占，一条指令搞定 |
对你意味着什么
如果你写 PyTorch，基本感知不到它——TMA 是被 cuBLAS、AI 词典：FlashAttention">FlashAttention 这类底层库和编译器悄悄用掉的，你享受的是更高吞吐、更省显存带宽的效果。如果你在写 CUDA kernel，TMA 已经是高性能矩阵乘、注意力机制里的标配工具，代价是编程模型变复杂：要学 descriptor 的构造、对齐要求、异步同步语义。可以先从"把原来手写的分块拷贝换成 TMA"这一步入手，收益通常最直观。