NOTEBOOK

GPU Emulator

公众号链接

在深度学习、科学计算、高性能计算等领域飞速发展的今天，GPU编程已成为实现高效并行计算的关键。开发者通过CUDA、PyTorch等工具编写kernel，但商业GPU的闭源特性，让kernel的实际执行过程难以观测，架构创新也面临诸多限制。GPU模拟器的出现，为这些问题提供了解决方案。

1 为什么需要GPU模拟器

商业GPU的微架构设计与内部优化策略不公开，给开发者和研究者带来两大核心痛点：

而GPU模拟器作为CPU上运行的软件工具，恰好弥补了这些短板：

2 GPU模拟器是什么

GPU模拟器的核心是通过软件建模，还原GPU的微架构和执行行为。以主流的GPGPU-Sim为例，其完整工作流程如下：

3 怎么本地部署(以GPGPU-SIM为例)

GPGPU-Sim仅支持Linux平台（32/64位均可），部署过程分为环境依赖配置、编译构建、运行验证三步，操作简洁高效。

3-1 环境依赖

需提前配置CUDA Toolkit及其他辅助依赖包，具体命令如下：

# 安装CUDA Toolkit 11.8
wget https://developer.download.nvidia.com/compute/cuda/11.8.0/local_installers/cuda_11.8.0_520.61.05_linux.run
sudo sh cuda_11.8.0_520.61.05_linux.run

# 配置CUDA环境变量
export CUDA_INSTALL_PATH=/usr/local/cuda-11.8
export PATH=$CUDA_INSTALL_PATH/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda-11.8/lib64:$LD_LIBRARY_PATH

# 安装其他依赖包
sudo apt-get install build-essential xutils-dev bison zlib1g-dev flex libglu1-mesa-dev doxygen graphviz python-pmw python-ply python-numpy libpng12-dev python-matplotlib libxi-dev libxmu-dev libglut3-dev

3-2 构建

通过Git克隆源码后，执行简单命令即可完成编译，支持调试和Release两种模式：

# 克隆源码仓库
git clone https://github.com/gpgpu-sim/gpgpu-sim_distribution.git
cd gpgpu-sim_distribution

# 配置环境（默认release模式，加debug参数可启用调试模式）
source setup_environment  # 或 source setup_environment debug

# 编译构建
make

注：Release模式模拟速度更快；调试模式需配合gdb，用于修改模拟器源码后的测试。

3-3 运行验证（以向量加法CUDA程序为例）

3-3-1 编写CUDA测试程序(vector_add.cu)

/* file: vector_add.cu */
#include<stdio.h>#define CHECK(call) \{ \
    const cudaError_t error = call; \if (error != cudaSuccess) { \printf("Error: %s:%d, ", __FILE__, __LINE__); \printf("code: %d, reason: %s\n", error, cudaGetErrorString(error)); \exit(1); \} \}

__global__ void vector_add(const float *a, const float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < n)
        c[idx] = a[idx] + b[idx];}

int main() {
    int n = 16;
    size_t bytes = n * sizeof(float);

    float h_a[16], h_b[16], h_c[16];for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;}

    float *d_a, *d_b, *d_c;

    CHECK(cudaMalloc(&d_a, bytes));
    CHECK(cudaMalloc(&d_b, bytes));
    CHECK(cudaMalloc(&d_c, bytes));

    CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

    vector_add<<<1, 16>>>(d_a, d_b, d_c, n);

    CHECK(cudaPeekAtLastError());
    CHECK(cudaDeviceSynchronize());

    CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));printf("Result: ");for (int i = 0; i < n; i++) {printf("%.1f ", h_c[i]);}printf("\n");

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);return0;}

3-3-2 编译CUDA程序

确保链接GPGPU-Sim的cudart库，编译命令如下：

nvcc --cudart shared -o vector_add vector_add.cu

3-3-3 配置模拟架构并运行

从GPGPU-Sim的预定义配置中选择目标架构（以SM86_RTX3070为例），复制配置文件后运行程序：

# 复制架构配置文件到当前目录
cp /path/to/gpgpu-sim_distribution/configs/tested-cfgs/SM86_RTX3070/* . -r

# 运行模拟程序
./vector_add

执行后会输出模拟耗时、指令执行速率、周期速率等基础信息，同时打印计算结果（如 Result: 0.0 3.0 6.0 ... 45.0）。此外，还会生成详细日志，包含模拟架构参数（SM数量、Warp Scheduler配置、Tensor Core数量等）和kernel执行细节（寄存器/共享内存使用、warp调度、流水线阻塞原因等）。

4 GPU模拟器能做哪些分析

GPU模拟器的价值不仅在于“运行”GPU程序，更在于“看透”运行过程、“验证”优化方案，主要体现在两大维度：

4-1 解析kernel的GPU执行过程

GPU采用SIMT（单指令多线程）模型，成百上千个线程通过warp协同执行，受限于共享内存、寄存器和缓存，执行逻辑远复杂于CPU的顺序执行。模拟器可解答以下关键问题：

例如，运行GEMM（矩阵乘法）kernel时，可通过模拟器观察每个warp的矩阵计算进度、共享内存/寄存器的缓存数据，快速定位访存瓶颈，为优化提供精准依据。

4-2 支撑GPU架构优化与创新

模拟器为架构研究者提供了安全可控的实验环境，无需依赖真实硬件即可验证创新设计：

典型案例：Accel-Sim框架通过模拟Volta架构GPU，发现L1缓存已非性能瓶颈，而内存调度策略对吞吐量影响显著，为后续架构优化指明了方向。

5 总结

GPU 模拟器为现代并行计算研究提供了不可或缺的工具。通过前面几节内容，我们可以总结出 GPU 模拟器的核心价值与应用场景：