hpc-lab-code/lab4/实验报告模板.md

实验 4: CUDA 程序设计与优化

实验 4.2: CUDA程序的编译和运行

实验目的

1. 掌握 CUDA 程序的基本结构和编译方法
2. 理解向量加法的并行实现
3. 分析数据规模对程序性能的影响

实验结果

数据规模与执行时间关系

数据规模 N	执行时间 (ms)	吞吐量 (elements/s)
128
256
512
1024
2048

性能分析

图表: 见 experiment_data/figures/vectoradd_performance.png

分析:

• 随着数据规模增加，执行时间的变化趋势是：
• 时间复杂度分析：
• GPU 并行效率分析：

---

实验 4.3: 基于CUDA优化矩阵乘法

思考题解答

思考一: matMultCUDAKernel1 对于矩阵的数据划分策略是什么？

答案:

matMultCUDAKernel1 采用的是 二维线程块和网格 的数据划分策略：

1. 线程组织:

   • 每个线程块 (Block) 的大小为 16×16 = 256 个线程
   • 每个线程负责计算结果矩阵 C 中的一个元素

2. 数据映射:

   int row = blockIdx.y * blockDim.y + threadIdx.y;
   int col = blockIdx.x * blockDim.x + threadIdx.x;
   

   • threadIdx.x 和 threadIdx.y: 线程在线程块内的局部坐标
   • blockIdx.x 和 blockIdx.y: 线程块在网格中的全局坐标
   • row 和 col: 直接映射到结果矩阵 C 的行列索引

3. 计算过程:

   • 每个线程计算 C[row][col] = Σ(A[row][k] × B[k][col])
   • 需要访问 A 的第 row 行和 B 的第 col 列

4. 优缺点:

   • ✓ 优点: 实现简单，每个线程独立计算
   • ✗ 缺点: 每个线程需要多次访问全局内存，带宽利用率低

---

思考二: matMultCUDAKernel2 对于矩阵运算的优化策略是什么，线程同步是否是必要的，为什么？

答案:

matMultCUDAKernel2 采用的是 共享内存分块 (Tiling) 优化策略：

1. 优化策略:

   • 将矩阵 A 和 B 分成小块 (Tile)，大小为 TILE_WIDTH × TILE_WIDTH
   • 每个线程块协作加载一个 Tile 到共享内存
   • 所有线程从共享内存读取数据进行计算，减少全局内存访问

2. 共享内存使用:

   __shared__ float shared_A[TILE_WIDTH][TILE_WIDTH];
   __shared__ float shared_B[TILE_WIDTH][TILE_WIDTH];
   

3. 线程同步的必要性:

   • 第一次 __syncthreads(): 确保所有线程完成数据加载到共享内存
   • 第二次 __syncthreads(): 确保所有线程完成当前 Tile 的计算，才能加载下一个 Tile

   为什么必要？

   • 共享内存是线程块级别的共享资源
   • 如果不同步，部分线程可能在其他线程完成数据加载前就开始计算
   • 会导致读取未初始化的数据，产生错误结果

4. 性能提升:

   • 共享内存带宽 ~ 1.5 TB/s，远高于全局内存 ~ 50 GB/s
   • 每个元素被重复使用 TILE_WIDTH 次，但只需加载一次到共享内存

---

思考三: matMultCUDAKernel2 还有没有可以继续优化的空间？

答案:

是的，还有多个优化方向：

1. 寄存器分块 (Register Tiling):

   • 将部分计算结果暂存在寄存器中
   • 进一步减少共享内存访问次数
   • 预期性能提升: 1.2-1.5x

2. 循环展开 (Loop Unrolling):

   • 展开内层计算循环，减少循环开销
   • 编译器可以更好地优化指令级并行

3. 内存合并访问优化:

   • 确保全局内存访问是合并的 (Coalesced)
   • 调整数据布局或访问模式

4. Warp 级别优化:

   • 使用 Warp Shuffle 指令在线程间直接交换数据
   • 减少共享内存使用

5. 流式多处理器 (SM) 优化:

   • 调整 BLOCK_SIZE 以最大化占用率 (Occupancy)
   • 平衡每个 SM 的线程块数量

6. 使用 Tensor Core (现代 GPU):

   • 利用 Volta/Turing 架构的 Tensor Core 进行矩阵乘法
   • 可达数倍性能提升

---

实验一: CPU vs GPU 性能对比

测试环境

• GPU: (从 gpu_info.txt 填写)
• CPU: (填写 CPU 型号)
• 编译器: nvcc, gcc
• 优化级别: -O3

性能数据

CPU (OpenMP) 不同线程数性能:

矩阵规模	线程数	时间 (ms)	GFLOPS	加速比
512×512	1			1.00
512×512	8
512×512	64
512×512	256
1024×1024	1			1.00
...	...

CUDA Kernel1 (基础版本):

矩阵规模	时间 (ms)	GFLOPS	相对CPU加速比
512×512
1024×1024
2048×2048
4096×4096

CUDA Kernel2 (共享内存优化):

矩阵规模	时间 (ms)	GFLOPS	相对CPU加速比	相对Kernel1提升
512×512
1024×1024
2048×2048
4096×4096

性能分析

图表: 见 experiment_data/figures/cpu_vs_gpu_comparison.png

关键发现:

1. CPU 多线程扩展性:
2. GPU 相对 CPU 的优势:
3. Kernel2 相对 Kernel1 的优化效果:
4. 不同矩阵规模下的性能趋势:

---

实验二: BLOCK_SIZE 对性能的影响

性能数据

矩阵规模	BLOCK_SIZE	时间 (ms)	GFLOPS
256×256	4
256×256	8
256×256	16
256×256	32
512×512	4
...	...

性能分析

图表: 见 experiment_data/figures/blocksize_analysis.png

最优 BLOCK_SIZE 分析:

1. 小矩阵 (256×256):

   • 最优 BLOCK_SIZE:
   • 原因:

2. 中等矩阵 (512×512 - 1024×1024):

   • 最优 BLOCK_SIZE:
   • 原因:

3. 大矩阵 (2048×2048):

   • 最优 BLOCK_SIZE:
   • 原因:

BLOCK_SIZE 影响因素:

• 共享内存大小限制 (每个 SM 有限)
• 线程束 (Warp) 的执行效率
• 占用率 (Occupancy) 的平衡
• 内存访问模式的优化

---

实验总结

主要发现

1. CUDA 并行计算相比 CPU 的性能优势:
2. 共享内存优化的重要性:
3. BLOCK_SIZE 对性能的影响规律:

性能优化建议

1. 对于小规模矩阵:
2. 对于大规模矩阵:
3. 通用优化策略:

实验收获

• 掌握了 CUDA 编程的基本方法
• 理解了 GPU 内存层次结构的优化
• 学会了使用性能分析工具评估并行程序

---

附录

运行命令

# 编译所有程序
cd lab4
xmake

# 运行实验并收集数据
./lab4.sh

# 生成图表 (需要安装 matplotlib)
./plot_results.py

数据文件

• experiment_data/gpu_info.txt: GPU 硬件信息
• experiment_data/vectoradd_results.txt: 向量加法测试数据
• experiment_data/matrixmul_comparison.txt: CPU vs GPU 对比数据
• experiment_data/blocksize_analysis.txt: BLOCK_SIZE 分析数据
• experiment_data/figures/: 生成的所有图表