2025 秋季 国科大《GPU架构与编程》大作业一

基于 LeNet 结构的脉冲卷积神经网络（Spiking CNN, SCNN）推理，CUDA 实现 + PTX 优化。

---

1. 项目概述

• 任务：在 Fashion-MNIST 上进行 SCNN 推理分类，模型骨架参考 LeNet（Conv-Pool-Conv-Pool-FC-FC）。
• 特色：
  • CUDA 实现卷积、池化、膜电位更新与发放逻辑。
  • PTX 内核对卷积/FC 做访存与指令级优化。
• 最终性能：单次推理耗时 0.045 s,准确率42%。

2. 代码结构

• src/
  • 主程序，含数据预处理、推理入口、CUDA/PTX 内核调用。
  • ptx/*.ptx：手写或生成的 PTX 内核（若有）。
  • include/ops.h：卷积/池化/FC/膜电位更新等算子声明。
  • utils/dataset.cpp：Fashion-MNIST 读取与批次组织。
• docs/
  • README.md（本文件），实验截图/曲线等。

3. 实验方法

• 模型与编码
  • 结构：LeNet 风格 Conv→Pool→Conv→Pool→FC→FC，输出用发放计数判定类别。
  • 脉冲编码：Rate Coding，将输入像素归一化为发放率；时间步 T 内多次前向累积膜电位。
  • 神经元：IF/LIF 离散更新，阈值触发后重置；可配置泄露系数 λ。
• 前向流程
  1. 输入批次预处理与脉冲序列生成。
  2. GPU 卷积/池化/FC 并行计算，输出写入膜电位缓冲。
  3. 膜电位阈值比较并生成发放脉冲，计数累计。
  4. 发放计数最大者为预测类别。
• CUDA 实现要点
  • 卷积/FC：线程块按输出 tile 划分，shared memory 缓权重/输入，coalesced 访存，必要时 float4 向量化加载。
  • 池化：max-pool 用 shared 减少全局访存。
  • 膜电位更新：分层 kernel，阈值比较用掩码避免分支发散。
• PTX 优化思路（加分项）
  • 手写 PTX 控制寄存器与加载指令，减少溢出，必要时使用 ldmatrix/mma（Tensor Core）或 FMAD 展开。
  • 显式缓存策略（如 ld.global.ca/ld.global.cs）与 pragma unroll 控制展开。
  • 将膜电位写回与卷积输出 fuse，减少 kernel launch 开销。

4. 复现步骤

• 环境
  • 硬件：NVIDIA GPU（建议 Compute Capability ≥ 7.0），显存 ≥ 4 GB。
  • 软件：CUDA 11.x/12.x，nvcc，C++17；可选 cuDNN 作为基线对照。
• 数据
  • 运行 scripts/download_fmnist.sh 下载并解压至 data/。
• 编译
  • 单文件示例：nvcc -O3 --use_fast_math homework0.07.cu -o scnn
  • 多文件示例：nvcc -O3 --use_fast_math -Iinclude src/*.cu -o bin/scnn
• 推理
  • 示例：./scnn --data ./data --batch 256 --timesteps 8 --use-ptx
  • 关闭 PTX：去掉 --use-ptx 走纯 CUDA 基线。

5. 实验结果

• 速度：0.045 s/推理（V100；batch size 与 T 请补充）。
• 精度：Top-1 准确率 0.42（Fashion-MNIST）。

6. 性能优化总结

• 访存：输入/权重按对齐格式（如 NHWC/NCHW16）提升 coalescing；shared tiling 减少全局带宽。

• 计算：关键环路手动展开，融合膜电位写回，减少 kernel 数量。

• 指令级：PTX 控制寄存器分配、避免 bank conflict，调度 FMAD/FMA 顺序。

• 并行度：选择合适 blockDim（如 128/256）平衡 occupancy 与寄存器压力；--use_fast_math 提升吞吐。

• 声明：仅供 2025 年秋季 国科大《GPU架构与编程》课程学习与交流使用。