356 lines
11 KiB
Markdown
356 lines
11 KiB
Markdown
# CUDA矩阵乘法性能实验分析报告
|
||
|
||
## 实验环境
|
||
- GPU: NVIDIA GeForce RTX 3090 (详见gpu_info.txt)
|
||
- CUDA版本: 根据代码推断为CUDA 11.x或更高版本
|
||
- CPU: 多核处理器(支持OpenMP)
|
||
|
||
---
|
||
|
||
## 实验一：CPU、CUDA Kernel1、CUDA Kernel2性能对比
|
||
|
||
### 1.1 实验数据汇总表
|
||
|
||
#### 表1-1：不同实现方式的执行时间对比（单位：ms）
|
||
|
||
| 矩阵规模 | CPU(8线程) | CPU(64线程) | CPU(256线程) | CUDA Kernel1 | CUDA Kernel2 |
|
||
|---------|-----------|------------|-------------|--------------|--------------|
|
||
| 512×512 | 747.483 | 743.606 | 748.649 | 0.316 | 0.827 |
|
||
| 1024×1024| 6033.205 | 6049.318 | 6051.757 | 2.374 | 6.484 |
|
||
| 2048×2048| 51065.609 | 50995.406 | 51083.363 | 19.190 | 53.599 |
|
||
| 4096×4096| - | - | - | 152.897 | 433.242 |
|
||
|
||
#### 表1-2：不同实现方式的性能对比（GFLOPS）
|
||
|
||
| 矩阵规模 | CPU(8线程) | CPU(64线程) | CPU(256线程) | CUDA Kernel1 | CUDA Kernel2 |
|
||
|---------|-----------|------------|-------------|--------------|--------------|
|
||
| 512×512 | 0.36 | 0.36 | 0.36 | 849.49 | 324.65 |
|
||
| 1024×1024| 0.36 | 0.35 | 0.35 | 904.75 | 331.22 |
|
||
| 2048×2048| 0.34 | 0.34 | 0.34 | 895.23 | 320.52 |
|
||
| 4096×4096| - | - | - | 898.90 | 317.23 |
|
||
|
||
#### 表1-3：GPU相对于CPU(8线程)的加速比
|
||
|
||
| 矩阵规模 | CUDA Kernel1加速比 | CUDA Kernel2加速比 |
|
||
|---------|------------------|------------------|
|
||
| 512×512 | 2365.45倍 | 903.85倍 |
|
||
| 1024×1024| 2541.37倍 | 930.48倍 |
|
||
| 2048×2048| 2661.05倍 | 952.73倍 |
|
||
|
||
### 1.2 详细分析
|
||
|
||
#### 1.2.1 CPU性能分析
|
||
|
||
**关键发现：**
|
||
1. **小矩阵规模(256×256)的可扩展性**
|
||
- 8线程: 86.012ms, 0.39 GFLOPS
|
||
- 64线程: 78.420ms, 0.43 GFLOPS (加速比1.14)
|
||
- 256线程: 76.496ms, 0.44 GFLOPS (加速比1.28)
|
||
- **结论**: 小矩阵可以放入CPU缓存，多线程扩展性较好
|
||
|
||
2. **中大矩阵规模的性能瓶颈**
|
||
- 从512×512开始，增加线程数几乎无性能提升
|
||
- 所有线程配置的性能都在0.34-0.36 GFLOPS
|
||
- **原因**: 受限于内存带宽，而非计算能力
|
||
|
||
3. **性能天花板**
|
||
- CPU最高性能仅0.44 GFLOPS
|
||
- 远低于GPU的300-900 GFLOPS
|
||
- **根本原因**: CPU的并行度有限，内存带宽远低于GPU
|
||
|
||
#### 1.2.2 CUDA Kernel1性能分析
|
||
|
||
**关键特点：**
|
||
1. **稳定的性能表现**
|
||
- 所有矩阵规模下性能稳定在850-905 GFLOPS
|
||
- 不随矩阵规模变化而明显波动
|
||
- **原因**: 简单的线程映射，良好的内存合并访问
|
||
|
||
2. **巨大的性能优势**
|
||
- 相比CPU(8线程)实现2000-2700倍加速比
|
||
- 相比CPU(256线程)实现2000-2700倍加速比
|
||
- **核心优势**: GPU的大规模并行计算能力
|
||
|
||
3. **设计优势**
|
||
- 每个线程计算一个结果元素，逻辑简单
|
||
- 全局内存访问模式良好，支持合并访问
|
||
- 无同步开销，执行效率高
|
||
|
||
4. **设计劣势**
|
||
- 每个线程需要重复访问全局内存
|
||
- 没有数据重用，内存带宽利用率低
|
||
- **优化空间**: 可以通过共享内存提升性能
|
||
|
||
#### 1.2.3 CUDA Kernel2性能分析
|
||
|
||
**意外发现：**
|
||
1. **性能反而下降**
|
||
- 性能稳定在317-331 GFLOPS
|
||
- 相比Kernel1性能下降约2.7-2.8倍
|
||
- **教训**: 盲目优化可能适得其反
|
||
|
||
2. **性能下降的根本原因**
|
||
|
||
**a) TILE_WIDTH=4太小**
|
||
- 共享内存的开销大于收益
|
||
- 每个tile只有16个元素，数据重用率低
|
||
- 频繁的tile加载增加了全局内存访问
|
||
|
||
**b) 同步开销**
|
||
- 每个tile需要两次`__syncthreads()`
|
||
- 对于小矩阵，同步开销占比很高
|
||
- 线程块内同步会阻塞所有线程
|
||
|
||
**c) 共享内存利用率低**
|
||
- 4×4的tile太小，无法充分利用共享内存带宽
|
||
- 现代GPU的共享内存设计用于更大的数据块
|
||
- Bank conflicts可能进一步降低性能
|
||
|
||
3. **设计问题**
|
||
- 过早优化：在没有充分理解硬件特性的情况下使用共享内存
|
||
- Tile size选择不当：4×4对于现代GPU来说太小
|
||
- 忽略了同步开销：小tile导致同步频率过高
|
||
|
||
#### 1.2.4 综合对比分析
|
||
|
||
**性能排名（从高到低）：**
|
||
1. CUDA Kernel1: ~900 GFLOPS
|
||
2. CUDA Kernel2: ~325 GFLOPS
|
||
3. CPU (任何线程数): ~0.36 GFLOPS
|
||
|
||
**关键结论：**
|
||
1. **GPU的绝对优势**: 即使是最简单的GPU实现，也比CPU快2000-2700倍
|
||
2. **优化需谨慎**: 设计不当的"优化"反而会降低性能
|
||
3. **简单往往更好**: Kernel1的简单设计优于Kernel2的复杂设计
|
||
4. **硬件理解很重要**: 必须根据GPU架构特性选择优化策略
|
||
|
||
---
|
||
|
||
## 实验二：BLOCK_SIZE对CUDA程序性能的影响
|
||
|
||
### 2.1 实验数据汇总表
|
||
|
||
#### 表2-1：不同BLOCK_SIZE下的执行时间（单位：ms）
|
||
|
||
| 矩阵规模 | 4×4 | 8×8 | 16×16 | 32×32 |
|
||
|---------|-----|-----|-------|-------|
|
||
| 256×256 | 0.116 | 0.040 | 0.029 | 0.026 |
|
||
| 512×512 | 0.831 | 0.265 | 0.189 | 0.178 |
|
||
| 1024×1024 | 6.539 | 2.022 | 1.397 | 1.364 |
|
||
| 2048×2048 | 54.023 | 16.080 | 11.454 | 11.019 |
|
||
|
||
#### 表2-2：不同BLOCK_SIZE下的性能（GFLOPS）
|
||
|
||
| 矩阵规模 | 4×4 | 8×8 | 16×16 | 32×32 |
|
||
|---------|-----|-----|-------|-------|
|
||
| 256×256 | 289.26 | 838.19 | 1170.29 | 1292.94 |
|
||
| 512×512 | 323.04 | 1014.10 | 1423.49 | 1506.57 |
|
||
| 1024×1024 | 328.40 | 1061.88 | 1536.94 | 1574.44 |
|
||
| 2048×2048 | 318.01 | 1068.38 | 1499.84 | 1559.16 |
|
||
|
||
#### 表2-3：相对于4×4的加速比
|
||
|
||
| 矩阵规模 | 8×8加速比 | 16×16加速比 | 32×32加速比 |
|
||
|---------|----------|------------|------------|
|
||
| 256×256 | 2.90倍 | 4.00倍 | 4.46倍 |
|
||
| 512×512 | 3.14倍 | 4.40倍 | 4.67倍 |
|
||
| 1024×1024 | 3.23倍 | 4.68倍 | 4.79倍 |
|
||
| 2048×2048 | 3.36倍 | 4.72倍 | 4.90倍 |
|
||
|
||
### 2.2 详细分析
|
||
|
||
#### 2.2.1 BLOCK_SIZE对性能的影响规律
|
||
|
||
**性能提升趋势：**
|
||
1. **4×4 → 8×8**: 性能提升约3倍（289→838 GFLOPS）
|
||
2. **8×8 → 16×16**: 性能提升约1.5倍（838→1423 GFLOPS）
|
||
3. **16×16 → 32×32**: 性能提升约1.05倍（1423→1574 GFLOPS）
|
||
|
||
**关键发现：**
|
||
- 性能提升幅度递减，呈现边际效应递减规律
|
||
- 32×32接近性能饱和点
|
||
- 不同矩阵规模下规律一致
|
||
|
||
#### 2.2.2 性能提升的深层原因分析
|
||
|
||
**1. 共享内存利用率提升**
|
||
|
||
**数据重用率分析：**
|
||
- 4×4 tile: 每个元素被重用4次
|
||
- 16×16 tile: 每个元素被重用16次
|
||
- 32×32 tile: 每个元素被重用32次
|
||
|
||
**全局内存访问减少：**
|
||
```
|
||
全局内存访问次数 ∝ 矩阵大小 / TILE_SIZE
|
||
```
|
||
- TILE_SIZE越大，全局内存访问次数越少
|
||
- 减少全局内存访问是性能提升的关键
|
||
|
||
**2. 线程级并行提升**
|
||
|
||
**线程块大小对比：**
|
||
- 4×4: 每个block只有16个线程
|
||
- 16×16: 每个block有256个线程
|
||
- 32×32: 每个block有1024个线程
|
||
|
||
**延迟隐藏效果：**
|
||
- 更多的线程可以更好地隐藏内存延迟
|
||
- GPU的warp scheduler有更多调度选择
|
||
- 提高了SM的利用率
|
||
|
||
**3. 计算与内存访问平衡**
|
||
|
||
**计算强度分析：**
|
||
- 小tile: 内存访问时间 > 计算时间（内存受限）
|
||
- 大tile: 计算时间 ≈ 内存访问时间（平衡）
|
||
- 最优tile: 计算与内存访问充分重叠
|
||
|
||
**指令级并行：**
|
||
- 大tile提供了更多的独立计算
|
||
- 编译器和硬件可以更好地优化指令调度
|
||
- 提高了流水线效率
|
||
|
||
#### 2.2.3 性能饱和现象分析
|
||
|
||
**从16×16到32×32性能提升有限的原因：**
|
||
|
||
**1. 共享内存容量限制**
|
||
- 每个SM的共享内存有限（如64KB）
|
||
- 32×32的tile已经占用较多共享内存
|
||
- 进一步增大tile会减少并发block数量
|
||
|
||
**2. 寄存器压力**
|
||
- 更大的tile需要更多寄存器存储累加器
|
||
- 寄存器使用过多可能导致spilling
|
||
- Spilling会将数据溢出到本地内存，严重降低性能
|
||
|
||
**3. 线程块调度效率**
|
||
- 过大的block会减少SM上驻留的block数量
|
||
- 降低了线程级并行度
|
||
- 可能导致SM资源利用率下降
|
||
|
||
**4. 内存带宽饱和**
|
||
- 当计算强度达到一定水平后
|
||
- 性能瓶颈转移到共享内存带宽
|
||
- 进一步增大tile无法提升性能
|
||
|
||
#### 2.2.4 最优BLOCK_SIZE选择策略
|
||
|
||
**针对当前GPU架构（RTX 3090）：**
|
||
- **最优选择**: 32×32
|
||
- **性能**: 1506-1574 GFLOPS
|
||
- **相比4×4提升**: 4.5-4.9倍
|
||
|
||
**通用选择原则：**
|
||
1. **考虑GPU架构**
|
||
- 不同架构有不同的最优值
|
||
- 需要查阅GPU架构文档
|
||
- 可以通过实验确定
|
||
|
||
2. **考虑问题规模**
|
||
- 小矩阵可能不适合大tile
|
||
- 需要平衡tile大小和矩阵规模
|
||
- 边界处理会增加复杂度
|
||
|
||
3. **资源平衡**
|
||
- 共享内存使用
|
||
- 寄存器使用
|
||
- 线程块调度
|
||
|
||
4. **性能调优方法**
|
||
- 使用CUDA性能分析工具（nvprof, Nsight）
|
||
- 监控共享内存使用率
|
||
- 监控寄存器使用情况
|
||
- 测试多个tile size选择最优
|
||
|
||
#### 2.2.5 与Kernel1的对比
|
||
|
||
**性能对比：**
|
||
- Kernel1 (无共享内存): ~900 GFLOPS
|
||
- Kernel2 (32×32共享内存): ~1574 GFLOPS
|
||
- **性能提升**: 1.75倍
|
||
|
||
**关键结论：**
|
||
1. **正确的共享内存优化非常有效**
|
||
- 从900提升到1574 GFLOPS
|
||
- 提升幅度达75%
|
||
|
||
2. **Tile size是关键**
|
||
- 4×4: 性能差（323 GFLOPS）
|
||
- 32×32: 性能优（1574 GFLOPS）
|
||
- 相差近5倍
|
||
|
||
3. **优化需要系统性思考**
|
||
- 不能盲目使用共享内存
|
||
- 必须选择合适的tile size
|
||
- 需要考虑硬件特性
|
||
|
||
---
|
||
|
||
## 总体结论与建议
|
||
|
||
### 3.1 主要发现
|
||
|
||
1. **GPU相比CPU有压倒性优势**
|
||
- 性能提升2000-2700倍
|
||
- 对于计算密集型任务，GPU是必然选择
|
||
|
||
2. **优化策略的重要性**
|
||
- 简单实现(Kernel1)已经很好
|
||
- 正确优化(Kernel2+32×32)可以再提升75%
|
||
- 错误优化(Kernel2+4×4)反而降低性能
|
||
|
||
3. **Tile size的关键作用**
|
||
- 4×4: 性能灾难
|
||
- 32×32: 性能最优
|
||
- 选择合适的tile size比使用共享内存本身更重要
|
||
|
||
### 3.2 实践建议
|
||
|
||
**对于CUDA矩阵乘法优化：**
|
||
|
||
1. **从简单实现开始**
|
||
- 先实现Kernel1这样的基础版本
|
||
- 确保正确性和基本性能
|
||
- 作为性能对比的基准
|
||
|
||
2. **谨慎使用共享内存**
|
||
- 理解共享内存的优势和代价
|
||
- 选择合适的tile size（至少16×16，推荐32×32）
|
||
- 避免过小的tile（如4×4）
|
||
|
||
3. **系统化性能调优**
|
||
- 使用性能分析工具
|
||
- 测试多个tile size
|
||
- 监控资源使用情况
|
||
|
||
4. **考虑更高级的优化**
|
||
- 寄存器分块
|
||
- 循环展开
|
||
- 使用Tensor Cores（现代GPU）
|
||
- 使用cuBLAS库
|
||
|
||
### 3.3 实验的价值
|
||
|
||
本实验很好地展示了：
|
||
1. 不同实现策略的巨大性能差异
|
||
2. 优化不当可能带来的负面影响
|
||
3. 系统化性能分析的重要性
|
||
4. 硬件特性对优化策略的影响
|
||
|
||
这些经验对于其他CUDA程序优化同样适用。
|
||
|
||
---
|
||
|
||
## 附录：图表说明
|
||
|
||
实验生成的图表：
|
||
1. `experiment1_analysis.png`: CPU、Kernel1、Kernel2性能对比
|
||
2. `experiment2_analysis.png`: 不同BLOCK_SIZE对性能的影响
|
||
|
||
原始数据文件：
|
||
1. `matrixmul_comparison.txt`: CPU、Kernel1、Kernel2的原始数据
|
||
2. `blocksize_analysis.txt`: 不同BLOCK_SIZE的原始数据
|
||
3. `gpu_info.txt`: GPU硬件信息
|