6.1 概述

计算两个维度分别为 (wA, hA) 和 (wB, wA) 的矩阵 A 和 B 的乘积 C 的任务，
以下列方式分为多个线程：

• 每个线程负责计算 C 的一个平方子矩阵 C_sub
• 块内的每个线程负责计算 C_sub 的一个元素

选择 C_sub 的维度 block_size 等于 16，以便：

• 每块的线程数是 warp 大小的倍数（参见 5.2）
• 保持低于每块的最大线程数（参见附录 A）

[图片 fig6-1-a.png / fig6-1-b.png]
图 6-1: 矩阵乘法

C_sub 等于两个矩形矩阵的乘积：

• 维度为 (wA, block_size) 的子矩阵 A，与 C_sub 具有相同的行索引
• 维度为 (block_size, wA) 的 B 的子矩阵，与 C_sub 具有相同的列索引

为了适应设备的资源，这两个矩形矩阵可根据需要划分为许多维度为 block_size
的平方矩阵，并且 C_sub 计算为这些平方矩阵的乘积之和。

其中每个乘积的执行过程：

1. 使用一个加载每个矩阵一个元素的线程，将两个相应的平方矩阵从全局内存
   加载到共享内存
2. 让每个线程计算乘积的一个元素
3. 每一线程将其中每个乘积的结果累计到寄存器中
4. 执行完毕后，将结果写入全局内存

通过以这种方式分块计算：

• 有效利用快速的共享内存
• 节省许多全局内存带宽（A 和 B 仅从全局内存读取 (wA / block_size) 次）

注意：编写此示例是为了清楚地说明各种 CUDA 编程原则，并非为一般的矩阵
乘法提供高性能的内核，所以不应如此构造。

6.2 源码清单

[图片 fig6-code2.png] 源码清单（第1页）

[图片 fig6-code3.png] 源码清单（第2页）

[图片 fig6-code4.png] 源码清单（第3页）

6.3 源码攻略

源码包含下列两个函数：

(1) Mul() — 作为 Muld() 的包装器的主机函数
(2) Muld() — 在设备上执行矩阵乘法的内核

6.3.1 Mul()（主机端包装函数）

Mul() 接受下列输入：

• 指向 A 和 B 的元素的主机内存的两个指针
• A 的高度和宽度，B 的宽度
• 指向应该写入 C 的主机内存的指针

Mul() 执行下列操作：

1. 使用 cudaMalloc() 将足够的全局内存分配到 A、B 和 C 中
2. 使用 cudaMemcpy() 将 A 和 B 从主机内存复制到全局内存
3. 调用 Muld() 在设备上计算 C
4. 使用 cudaMemcpy() 将 C 从全局内存复制到主机内存
5. 使用 cudaFree() 释放为 A、B 和 C 分配的全局内存

伪代码流程：

主机分配设备内存
↓
主机→设备: 复制矩阵 A, B
↓
设备: Muld() 内核计算 C = A × B
↓
设备→主机: 复制结果矩阵 C
↓
释放设备内存

6.3.2 Muld()（设备端内核函数）

除了指针指向设备内存而非主机内存之外，Muld() 与 Mul() 具有相同的输入。

对于每个块，Muld() 迭代处理所有需要计算 C_sub 的 A 和 B 的子矩阵。

在每次迭代中，此函数执行：

(1) 将 A 的一个子集和 B 的一个子集从全局内存加载到共享内存中

(2) __syncthreads()
同步以确保两个子矩阵都由块内的所有线程完全加载

(3) 计算两个子集的乘积并将其加到上一次迭代期间获得的乘积中

(4) __syncthreads()
再次同步以确保在开始下一次迭代之前两个子集的乘积已经完成

按照 5.1.2.1 和 5.1.2.4 所述，编写 Muld() 是为了最大化内存性能：

┌───────────────────────────────────────────────────┐
│ 全局内存合并 │
│ │
│ 假设 wA 和 wB 是 16 的倍数（如 5.1.2.1 所建议）， │
│ 则确保了全局内存合并，因为 a、b 和 c 都是 │
│ BLOCK_SIZE 的倍数，BLOCK_SIZE 等于 16。 │
└───────────────────────────────────────────────────┘

┌───────────────────────────────────────────────────┐
│ 共享内存库冲突避免 │
│ │
│ 对于每个半 warp，没有任何共享内存库冲突： │
│ - 所有线程的 ty 和 k 都是相同的 │
│ - tx 在 0 到 15 之间变化 │
│ │
│ 访问模式分析： │
│ As[ty][tx] → 每个线程访问不同的库 ✓ │
│ Bs[ty][tx] → 每个线程访问不同的库 ✓ │
│ Bs[k][tx] → 每个线程访问不同的库 ✓ │
│ As[ty][k] → 每个线程访问同一个库（使用广播）✓ │
└───────────────────────────────────────────────────┘

================================================================================
本章体现的核心 CUDA 编程模式

1. 分块策略 (Tiling)

   • 将大矩阵分解为线程块大小的小块
   • 使用共享内存作为每块的本地缓存

2. 内存层次利用
   全局内存 ──加载──→ 共享内存 ──计算──→ 寄存器 ──写回──→ 全局内存
   (慢, 大) (快, 小) (最快) (合并写)

3. 同步机制

   • __syncthreads() 确保块内所有线程的数据一致性
   • 加载同步 → 计算 → 累加同步 → 下一次迭代

4. 合并访问优化

   • 数据布局使线程访问满足合并条件
   • 块大小选择 warp 大小的倍数

5. 库冲突避免

   • 通过正确设计数据访问模式消除共享内存库冲突