261 lines
7.5 KiB
Markdown
261 lines
7.5 KiB
Markdown
# 实验 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. **数据映射**:
|
||
```cuda
|
||
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. **共享内存使用**:
|
||
```cuda
|
||
__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 内存层次结构的优化
|
||
- 学会了使用性能分析工具评估并行程序
|
||
|
||
---
|
||
|
||
## 附录
|
||
|
||
### 运行命令
|
||
```bash
|
||
# 编译所有程序
|
||
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/`: 生成的所有图表
|