yly/hpc-lab-code
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
hpc-lab-code/lab4/experiment_data/实验分析报告.md
yly 45d06345cb save results
2026-01-22 04:31:52 +08:00

11 KiB
Raw Permalink Blame History

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硬件信息