寄存器与本地内存
了解每个线程最快的存储方式及其溢出
寄存器与本地内存 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 1 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
您拥有的最快内存
每个线程都有自己的私有寄存器，这是芯片上速度最快的存储空间。它们紧邻计算单元，因此读取它们几乎不需要成本。
普通变量会变成寄存器
在内核中编写普通的局部变量时，编译器通常会将其保存在寄存器中。完全不需要特殊语法。🙂
__global__ void k() {
int x = 5; // x lives in a register
float y = x * 2.0f;
}属于单个线程
一个寄存器值只属于一个线程。线程 7 无法查看线程 8 的寄存器，因此每个线程的工作能够保持清晰分离。
寄存器是共享资源池
每个流式多处理器都有固定大小的寄存器文件。所有常驻线程都会分用它，因此每个线程使用的寄存器越少，就能同时运行越多线程。
寄存器用尽时
如果内核需要的寄存器超过可用数量，多出的值就会被移到其他位置。这个过程称为寄存器溢出。
溢出数据去哪里：本地内存
溢出的值会进入本地内存，尽管名称如此，它并不在芯片上，而是实际存放在速度较慢的片外 DRAM 中。
本地内存仍然是线程私有的
本地内存和寄存器一样，对每个线程都是私有的。区别仅在于速度，因为本地内存的访问速度要慢得多。
大型本地数组会导致溢出
声明大型的线程私有数组通常会迫使它进入本地内存，因为寄存器根本不足以容纳每个元素。
__global__ void k() {
float buf[64]; // likely lives in local memory
}溢出会损害性能
每次溢出都会将免费的寄存器访问变成一次缓慢的 DRAM 访问。降低寄存器压力，是您能够进行的最简单有效的优化之一。
查看寄存器数量
您可以要求编译器报告每个线程使用的寄存器数量。向 nvcc 传入 --ptxas-options=-v，它就会打印每个内核的使用情况。
nvcc --ptxas-options=-v kernel.cu有意限制寄存器数量
__launch_bounds__限定符会提示编译器限制寄存器数量，以少量降低每个线程的速度为代价，换取更多线程同时运行。
__global__ void __launch_bounds__(256)
myKernel() { /* ... */ }快速检查
您声明了一个较大的线程私有数组，性能却下降了。可能发生了什么？
回顾：快速且私有
您已经学会：寄存器是线程私有的最快存储空间，但资源池有限，溢出后会进入速度较慢的本地内存。请保持较低的寄存器压力。🎯
用 AI 导师学习 C++ — 免费
在浏览器中编写并运行真实代码，获得全天候 AI 导师的即时帮助，并在网页或应用中继续学习。
- 课程
- 30
- 课程
- 120
常见问题解答
「寄存器与本地内存」课时是免费的吗？
是的 — 「寄存器与本地内存」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「寄存器与本地内存」这节课中我会学到什么？
了解每个线程最快的存储方式及其溢出 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 1 节课，共 4 节。
「寄存器与本地内存」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- 寄存器与本地内存
- 全局内存的权衡
- 常量内存及其缓存
- 建立层次结构的心智模型