# GPU-Homework1

**Repository Path**: madi-peak/gpu-homework1

## Basic Information

- **Project Name**: GPU-Homework1
- **Description**: 2025年秋季国科大《GPU架构与编程》大作业一
- **Primary Language**: Unknown
- **License**: Not specified
- **Default Branch**: master
- **Homepage**: None
- **GVP Project**: No

## Statistics

- **Stars**: 0
- **Forks**: 1
- **Created**: 2026-01-12
- **Last Updated**: 2026-01-25

## Categories & Tags

**Categories**: Uncategorized

**Tags**: None

## README

# 2025年秋季国科大《GPU架构与编程》

本目录包含：
- FashionMNIST 数据下载脚本
- SCNN 训练脚本（导出权重）
- CUDA 推理程序（`infer_optimized.cu`）
- 已训练好的权重（`weights/`）

下面给出**完整使用顺序**和**infer_optimized.cu 的优化思路**。

## 目录结构与用途

```
./
├── cmd.sh                    # 简单编译脚本（sm_70）
├── data/                     # FashionMNIST 数据目录（raw 文件）
├── download_fmnist.py        # 下载 FashionMNIST 测试集到 ./data
├── program1_scnn_train.py    # 训练并导出权重到 ./weights
├── infer_optimized.cu        # CUDA 推理实现（已优化）
├── weights/                  # 预训练权重（txt）
├── requirement.txt           # Python 依赖
└── readme.md                 # 说明文档
```

## 环境依赖

### Python 依赖（用于下载数据/训练）

```bash
pip install -r requirement.txt
```

### CUDA 编译环境（用于推理）
- NVIDIA GPU + CUDA Toolkit
- 推荐架构：sm_70 / sm_80 / sm_90（代码含 WMMA 路径）

## 使用顺序（从零开始）

### 1. 安装依赖（只用于下载/训练）
```bash
pip install -r requirement.txt
```

### 2. 下载 FashionMNIST 数据
```bash
python download_fmnist.py
```
会生成 `./data/FashionMNIST/raw` 目录，推理程序会自动查找该路径。

### 3. （可选）训练并导出权重
> 如果直接使用 `weights/` 里的预训练权重，可跳过此步。

```bash
python program1_scnn_train.py
```
- 训练完成后会将权重导出到 `./weights/*.txt`
- 训练脚本会自动下载数据集（若 data 目录缺失）

### 4. 编译 CUDA 推理程序

**方式 A：使用脚本（sm_70）**
```bash
bash cmd.sh
```
生成可执行文件 `./inference_prog`

**方式 B：手动编译（可扩展架构）**
```bash
nvcc -O3 -std=c++14 infer_optimized.cu -o infer_opt \
    -gencode arch=compute_70,code=sm_70 \
    -gencode arch=compute_80,code=sm_80 \
    -gencode arch=compute_90,code=sm_90
```

### 5. 运行推理
```bash
./inference_prog ./weights
# 或
./infer_opt ./weights
```
程序会自动在以下路径中查找 FashionMNIST raw 数据：
- `./data/FashionMNIST/raw`
- `./weights/../data/FashionMNIST/raw`
- `./weights/../../..//data/FashionMNIST/raw`

输出格式：
```
<seconds>:<accuracy>
```
例如：`0.1234:0.8978`

## infer_optimized.cu 优化思路

### 1. 算子融合与并行策略
- **Conv1 + IF + MaxPool 融合**：一次 kernel 完成卷积、脉冲更新与 2x2 池化，减少中间结果回写。
- **Conv2 + IF 融合**：stride=3 的卷积层与 IF 更新合并。
- **每线程负责一个池化输出块**：Conv1 中一次计算 2x2 的 4 个位置并直接池化。

### 2. PTX 级指令优化
- 自定义 `ptx_fma` 与 `ptx_max`，减少指令开销。
- 关键卷积/全连接内核使用 FMA 以提升吞吐。

### 3. 内存访问优化
- 使用 `__ldg()` 加速只读权重与输入访问。
- 连续内存布局配合展开循环提升访存合并度。
- 预先分配并复用 GPU/CPU 资源，避免重复分配开销。

### 4. 资源复用与流水
- 使用静态 `InferenceResources`：
  - Pinned host memory（`cudaMallocHost`）
  - 单一 stream + 异步 H2D/D2H
- 批次内多步时间循环 (`T_STEPS=2`) 复用已分配 buffer。
- warmup 预热，规避首次 kernel 启动开销。

### 5. 精度与性能权衡
- 默认使用 **FP32 版本的全连接层**，保证准确率（≈0.8978）。
- 代码内保留 **WMMA TensorCore 路径**：
  - `linear_if_wmma_kernel`
  - `linear_accum_wmma_kernel`
- 若追求更高吞吐，可切换为 WMMA 版本，但精度可能略降。

### 6. 可进一步尝试的优化（可选）
- Conv2 采用 shared memory tile + vectorized load
- FC 层进行 block-level reduction + cooperative groups
- 针对固定 batch size 做 kernel 特化

## 权重说明
推理程序要求权重文件存在并匹配尺寸：
- `conv1.weight.txt` (250)
- `conv1.bias.txt` (10)
- `conv2.weight.txt` (1440)
- `conv2.bias.txt` (16)
- `fc1.weight.txt` (32768)
- `fc1.bias.txt` (128)
- `fc2.weight.txt` (4096)
- `fc2.bias.txt` (32)
- `fc3.weight.txt` (320)
- `fc3.bias.txt` (10)

若权重缺失或维度错误，程序会直接退出并提示。