# 实验 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/`: 生成的所有图表