为内积分块
每个阶段加载 A 和 B 的子数据块
为内积分块 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 2 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
分块思想
分块技术会将矩阵拆分成适合高速片上内存的小型方形分块。线程协作加载一个分块,然后多次重复使用它。
为什么使用共享内存
分块存放在 __shared__ 内存中,块内的每个线程都可以访问。读取它的速度远快于一次又一次访问全局内存。⚡
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];一个块,一个输出分块
每个 block 负责输出 C 中一个 TILE×TILE 的区域。该块中的线程协作完成整个区域的计算。
分块大小与块匹配
您可以将 TILE 设置为块的宽度,通常为 16 或 32。这样,每个线程就能为每个分块加载恰好一个元素。
#define TILE 16
dim3 threads(TILE, TILE);将线程映射到分块槽位
在分块内部,线程的槽位就是它的 threadIdx。它的全局行和列仍然来自块索引和线程索引。
int ty = threadIdx.y, tx = threadIdx.x;
int row = blockIdx.y*TILE + ty;
int col = blockIdx.x*TILE + tx;加载 A 的一个分块
每个线程将 A 当前分块中的一个元素复制到共享内存。整个块协作完成对 A 的一个完整 TILE×TILE 区域的加载。
As[ty][tx] = A[row*N + (phase*TILE + tx)];加载 B 的一个分块
与此同时,每个线程加载 B 分块中的一个元素。现在两个分块都位于片上内存中,可以快速重复读取。
Bs[ty][tx] = B[(phase*TILE + ty)*N + col];计算前进行同步
请调用 __syncthreads(),确保所有线程都完成加载后再读取分块。跳过这一步会得到无效结果。
__syncthreads();在分块上计算
现在,每个线程都会遍历这个分块,只读取共享内存。这些读取的代价远低于全局内存读取。
for (int k = 0; k < TILE; ++k)
sum += As[ty][k] * Bs[k][tx];重复使用带来的收益
每个已加载的值会被 TILE 个线程使用,而不是只被一个线程使用。这种数据重用正是分块矩阵乘法能够高速运行的原因。
一个分块还不够
单个分块只能覆盖点积的一部分。您需要在多个阶段中重复执行加载、同步和计算步骤,下一课将对此进行介绍。
快速检查
请回忆处理共享分块时各步骤的顺序。
回顾
您将 A 和 B 的分块暂存到共享内存中,进行同步,然后利用低成本的片上读取完成计算。收益来自重复使用。接下来将介绍各个阶段。🧱
用 AI 导师学习 C++ — 免费
在浏览器中编写并运行真实代码,获得全天候 AI 导师的即时帮助,并在网页或应用中继续学习。
- 课程
- 30
- 课程
- 120
常见问题解答
「为内积分块」课时是免费的吗?
是的 — 「为内积分块」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「为内积分块」这节课中我会学到什么?
每个阶段加载 A 和 B 的子数据块 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗?
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 2 节课,共 4 节。
「为内积分块」课时需要多长时间?
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗?
能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。