WMMA 片段 API
加载、mma_sync 与存储片段
WMMA 片段 API 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 3 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
WMMA API
要直接编程使用张量核心,您需要使用WMMA,即 nvcuda::wmma 命名空间中的 warp 矩阵乘加 API。🧩
整个 warp 协同工作
WMMA 以 warp 为单位工作:一个warp中的全部 32 个线程共同处理一个分块。您需要以分块而不是单个线程为思维单位。
认识片段
片段是 WMMA 的数据类型,用于保存一个 warp 所负责的矩阵分块切片。每个线程分别拥有其中的几个元素。
三种片段角色
您可以声明标记为 matrix_a、matrix_b 或累加器的片段。该标记会告诉 WMMA 如何将分块用于乘加运算。
分块形状是固定的
片段使用固定的分块大小,例如 16 乘 16 乘 16。您需要选择受支持的形状;硬件只接受这些组合。
第 1 步:加载
首先调用 load_matrix_sync,将一个分块从内存加载到片段中。这次加载会替您将数据分散到整个 warp。
wmma::load_matrix_sync(a_frag, ptr, ldm);第 2 步:清零累加器
在进行加法前,使用 fill_fragment 将结果分块清零。干净的累加器意味着求和从已知值开始。
wmma::fill_fragment(c_frag, 0.0f);第 3 步:矩阵乘加
核心调用是 mma_sync。它直接使用张量核心对已加载的片段执行 D = A 乘以 B 再加上 C。
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);第 4 步:存储
最后,store_matrix_sync 会将累加器片段写回内存。这次存储会将每个线程负责的部分汇集为一个分块。
wmma::store_matrix_sync(out, c_frag, ldm, layout);同步意味着 warp 同步
_sync 后缀是一个提醒:每个 WMMA 调用都是warp 同步的。全部 32 个通道必须一起到达,否则行为未定义。
布局与前导维度
加载和存储需要前导维度(即行间步长),还需要行主序或列主序布局,以便正确读取内存。
快速检查
哪个 WMMA 调用会真正让张量核心执行矩阵乘加?
回顾
您已经完成 WMMA 流程:声明片段、加载分块、清零累加器、调用 mma_sync,然后存储。每一步都是 warp 同步的。🙌
用 AI 导师学习 C++ — 免费
在浏览器中编写并运行真实代码,获得全天候 AI 导师的即时帮助,并在网页或应用中继续学习。
- 课程
- 30
- 课程
- 120
常见问题解答
「WMMA 片段 API」课时是免费的吗?
是的 — 「WMMA 片段 API」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「WMMA 片段 API」这节课中我会学到什么?
加载、mma_sync 与存储片段 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗?
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 3 节课,共 4 节。
「WMMA 片段 API」课时需要多长时间?
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗?
能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- 张量核心计算什么
- 混合精度:FP16、BF16、TF32
- WMMA 片段 API
- 数值稳定性的权衡