谈谈如何手搓GPU - 科技区角

谈谈如何手搓GPU图1

> 本文转载自「Disaggregated-AI」

笔者的 vibe-ai-infra 项目已经走完了Coding Agent、Claw 、RL Infra 和 Virt 全栈。这篇文章谈谈如何用 AI 手搓一个 GPU，一个理论上可以上板或者流片的硬件版本，项目取名 SpinalGPU。

地址：https://github.com/lastweek/SpinalGPU

动手前的问题

对接哪个硬件生态？

我最终选择了让 SpinalGPU 对接英伟达的 CUDA 生态，具体来说，对接 NVIDIA PTX 一层（支持部分核心关键指令集）。
原因比较简单，设计任何一款系统或者芯片，最重要的是起其北向 API。英伟达芯片的北向 API - PTX ISA指令集，有较为全面的文档，因此我可以参考这些文档去实现底层的功能，实现完之后也能去跑英伟达生态下的算子等，可加速学习的过程。适配国产AI 芯片是一个道理，笔者把这个留给 Future Work 或者有感兴趣的同学可以联系我一起搞。

用那个硬件开发语言？

我使用了 SpinalHDL 来开发，所以取了 SpinalGPU 这个项目名。
对于硬件开发描述不熟悉的同学，解释如下：传统硬件开发语言是大家可能听过的 Verilog、VHDL，它们是类似于 C 语言的定位，是非常底层描述数字系统硬件结构和行为的语言，学习和开发成本比较陡峭；在 2012 年~2014 年，涌现了更高抽象层次的硬件编程语言例如 Chisel 和 SpinalHDL，它们是基于Scala语言构建的硬件构造语言，用户可使用 Scala 高级语言开发硬件，编译器会把 Scala 编译成 Verilog 或 VHDL，大大加速了硬件开发迭代速度。
Chisel 和 SpinalHDL 两个语言很像，我最终选择 SpinalHDL 的原因是因为我比较喜欢 SpinalHDL 的语法和内置组件。比如可以用 a >> b 来串联两个端口。我在博士期间也用过 SpinalHDL 开发过 2 篇论文：Clio, ASPLOS'22 和 SuperNIC, FPGA'24. 目前基于 Chisel 比较有名的项目有中科院计算所包云岗老师的香山开源芯片项目（https://xiangshan.cc/zh/）。

怎么验收？

这个问题的答案有两个。第一个选择是功能仿真 only，也就是最后只用基于 CPU 的仿真来验收功能用例，不看最后硬件的时序等；第二个是选择上 FPGA 板子真实部署，也就是需要引入类似 Vivado 的软件、时序调试、上板部署等。
作为个人娱乐项目，我选择扣1。

SpinalGPU总览

我们先看看 SpinalGPU 在 Vibe 完后的架构总览。

简单来说，它运行 PTX 编译后的算子二进制文件，SpinalGPU 内部包含接受算子下发的组件，包含多个 SM、包含共享内存等。SpinalGPU 中所有的计算资源、内存资源都是可配置的。我在本地仿真基本就用 1~2 个 SM，每个 SM 用 128 个 CUDA Cores，每个 warp 32 threads ，每个 SM 用 1 个 Tensor Core、SFU。

这是我们 SpinalGPU 的顶层定义，包括时钟域、面向 Host 的控制模块（例如接受 Host 的算子下发）、计算单元GpuCluster（包含可配置个数的 SM）。这个 GPU Top 整体非常清晰简单，我们后面逐步展开如何填充核心的计算单元：

class GpuTop(val config: GpuConfig = GpuConfig.default) extends Component {  val io = GpuTopIo(config)
  val coreClock = in Bool()  val coreReset = in Bool()
  val coreClockDomain = ClockDomain(    clock = coreClock,    reset = coreReset,    config = ClockDomainConfig(resetKind = SYNC)  )
  val core = new ClockingArea(coreClockDomain) {    val hostControlBlock = new HostControlBlock(config)    val gpuCluster = new GpuCluster(config)
    hostControlBlock.io.axi <> io.hostControl    gpuCluster.io.command.command := hostControlBlock.io.command    gpuCluster.io.command.start := hostControlBlock.io.start    gpuCluster.io.command.clearDone := hostControlBlock.io.clearDone    hostControlBlock.io.executionStatus := gpuCluster.io.command.executionStatus    io.debugExecutionStatus := gpuCluster.io.command.executionStatus    io.memory <> gpuCluster.io.memory  }}

指令集 PTX ISA

我们第一步要解决和生态的对接问题。

如上所示，我们对接了英伟达芯片的指令集：PTX （Parallel Thread Execution）。PTX 是桥接上层编程语言 CUDA 和下层 GPU 芯片的中间层：CUDA 算子编译成 PTX 汇编语言，然后再由汇编器把 PTX 编译成针对不同 GPU 的二进制文件，这些二进制是最后执行在 GPU 上的算子。PTX ISA 最新版本是 9.2（as of Mar 27, 2026），大约有 135 个不同的指令，其中计算和数据相关指令占大头，如下图所示。

在 SpinalGPU 项目中，我最后大约实现了约 30~40 条 PTX指令。

PTX 指令是公开的，但最后运行在芯片中的 binary 形式不是（也许也是？笔者这里没有深究）。SpinalGPU 为了方便起见，有一套自己的指令 encoding 方式（比如一条指令 32 字节，前 4 个四节为 opcode 等）。为了能把 PTX 指令编译成可以在 SpinalGPU 上运行的二进制，我们引入了一个简单的 PTX Assembler for SpinalGPU，它的输入是 PTX 代码，输出是 .bin 二进制文件，可直接运行在 SpinalGPU 上。下图展示了我们的编译流是如何嵌到 CUDA 现有的编译流程中的：

PTX 不是什么黑科技，它只是任何芯片都有的 ISA。并且在大模型的帮助下，理解、生成这些代码都变得容易。好奇 PTX 代码长什么的同学，可以参考SpinalGPU中内置 PTX 测试Kernels，见：

https://github.com/lastweek/SpinalGPU/tree/main/kernels

Streaming Multiprocessor (SM)

解决了生态对接之后，下一步就是实现 GPU 最核心的计算单元：Streaming Multiprocessor（SM）。

一个 GPU 通常由多个 SM 组成。SM 是 NVIDIA GPU 的基本执行与资源管理单元，负责驻留并调度多个 warp，并在本地的执行流水线、寄存器文件、shared memory / L1 cache 等片上资源上推进指令执行。

从 CUDA 编程模型看，thread block 以 SM 为驻留边界：一个 block 中的所有线程都会在同一个 SM 上执行；而在 SM 内部，线程再按 32 个线程组织成一个 warp，warp 是硬件调度的基本粒度。

从微结构上看，SM 并不是若干 CUDA Cores 的简单堆叠，而是一个完整的片上执行节点。它内部通常包含多类执行与访存部件，例如负责通用标量算术的 CUDA Cores、负责矩阵乘加的 Tensor Cores、负责特殊函数计算的 SFU、负责访存的 LSU，以及配套的寄存器文件、片上缓存与调度逻辑。以 H100 为例，一个 SM 包含 128 个 FP32 CUDA Cores 和 4 个第四代 Tensor Cores。

进一步看，NVIDIA 近几代架构通常会把一个 SM 划分为多个内部处理分区。以 H100 这一代为例，可以把一个 SM 视为由 4 个 SM sub-partitions 组成；每个分区内部都有自己的 warp scheduler、register file 和一组执行单元，而一些缓存与共享资源则在整个 SM 范围内共享。也正因为如此，H100 的一个 SM 并不是“最多只能运行 4 个 warps”，而是“内部有 4 个主要处理分区，可在其上调度多个 ready warps”；从 occupancy 的角度看，一个 H100 SM 最多可驻留 64 个 warps，只是任一时刻真正被发射、向前推进的 warp 数会受到调度器和执行资源的限制。

在 SpinalGPU 中，我们首先设计支持 multi-SM。其次，在一个 SM 内部，我们设计包含一些公共资源和多个 SmPartition，一个 SmParition 会包含如下资源：

Warp Queue 任务队列
Warp Scheduler 调度器
执行 SIMT 指令的 CUDA Cores
执行矩阵计算加速的 Tensor Cores
执行特殊计算加速的 Special Function Unit (SFU)
执行内存读写的 Load/Store Unit (LSU)
数据和指令的 Cache
大量的 Registers

下面是生成的 SM 模块，包含可多个 SmPatition 和一些公共资源：

// SpinalGPU written in SpinalHDL Language.class StreamingMultiprocessor(val config: SmConfig = SmConfig.default) extends Component {  val io = StreamingMultiprocessorIo(config)
  private val subSmPartitions = Array.fill(config.subSmCount)(new SubSmPartition(config))  private val smAdmissionController = new SmAdmissionController(config)  private val warpStateTable = new WarpStateTable(config)  private val warpBinder = new WarpBinder(config)  private val l0InstructionCaches = Array.fill(config.subSmCount)(new L0InstructionCache(config))  private val l1InstructionCache = new L1InstructionCache(config)  private val l1DataSharedMemory = new L1DataSharedMemory(config)  private val sharedMemory = new SharedMemory(config)  private val externalMemoryArbiter = new ExternalMemoryArbiter(config)  private val externalMemoryAdapter = new ExternalMemoryAxiAdapter(config)  ...

SM 中的其他代码基本都是 plumbing code，比如下面是串联 SubSM 和各种 L0/L1 Cache：

// 在 SpinalHDL 语言中， A端口 <> B 端口， 意味着双向对接 A 和 B // A >> B， 意味着 A 单向到 Bfor (subSm <- 0 until config.subSmCount) {    l0InstructionCaches(subSm).io.request <> subSms(subSm).io.fetchMemReq    subSms(subSm).io.fetchMemRsp <> l0InstructionCaches(subSm).io.response    l1InstructionCache.io.subSmReq(subSm) <> l0InstructionCaches(subSm).io.l1Req    l0InstructionCaches(subSm).io.l1Rsp <> l1InstructionCache.io.subSmRsp(subSm)、  .....  }

SmParition 里则是真正的核心计算资源，比如 register、Cuda Core、Tensor Core，SFU 等：

// SpinalGPU written in SpinalHDL Language.class SubSmPartition(config: SmConfig) extends Component {  val io = new Bundle {    ..  }   // 定义每个 SubPM 拥有的核心资源，比如寄存器、CudaCore、TensorCore 等  private val registerFile = new WarpRegisterFile(config, slotCount = config.residentWarpsPerSubSm)  private val slotTable = new LocalWarpSlotTable(config)  private val scheduler = new LocalWarpScheduler(config)  private val fetchUnit = new InstructionFetchUnit(config)  private val decodeUnit = new DecodeUnit(config)  private val specialRegisterUnit = new SpecialRegisterReadUnit(config)  private val cudaCoreArray = new CudaCoreArray(config)  private val loadStoreUnit = new LoadStoreUnit(config)  private val specialFunctionUnit = new SpecialFunctionUnit(config)  private val tensorCoreBlock = new TensorCoreBlock(config)  ..

SmParition 的逻辑简单来说就是 warp调度+指令计算，具体来说：每个 SmPartition 都有一个 warp scheduler 和一个 warp 队列。在每个运行 cycle，warp scheduler 会检查队列有没有可以执行的 warp，如果有，则调度执行这个 warp 的指令。对于 warp 中的每一条执行指令，会有一个大的 dispatch 模块把指令发给对应的计算模块，如果计算模块需要长时间才能完成，则 warp scheduler 选选择调度其他的 ready warps：

 // SpinalGPU written in SpinalHDL Language. // decoding and dispatch a warp's instruction    ...    } elsewhen (decodeUnit.io.decoded.target === ExecutionUnitKind.CUDA) {      cudaCoreArray.io.issue.valid := True      when(cudaCoreArray.io.issue.ready) {        capturePendingOp(decodeUnit.io.decoded)        engineState := EngineState.WAIT_CUDA      }    } elsewhen (decodeUnit.io.decoded.target === ExecutionUnitKind.LSU) {      loadStoreUnit.io.issue.valid := True      when(loadStoreUnit.io.issue.ready) {        capturePendingOp(decodeUnit.io.decoded)        engineState := EngineState.WAIT_LSU      }    } elsewhen (decodeUnit.io.decoded.target === ExecutionUnitKind.SFU) {      specialFunctionUnit.io.issue.valid := True      when(specialFunctionUnit.io.issue.ready) {        capturePendingOp(decodeUnit.io.decoded)        engineState := EngineState.WAIT_SFU      }    } elsewhen (decodeUnit.io.decoded.target === ExecutionUnitKind.TENSOR) {      tensorCoreBlock.io.issue.valid := True      when(tensorCoreBlock.io.issue.ready) {        capturePendingOp(decodeUnit.io.decoded)        engineState := EngineState.WAIT_TENSOR      }

总计来说，AI 自动生成的 SM 顶层模块非常清晰，你可以很直观的看到各个核心组件以及串联逻辑。更有挑战的调度算法、调度效率、SIMT divergence 之后的执行效率，以及特别是数据读写效率：It's the memory, stupid! 这些笔者就按下不表了，感兴趣的同学可以自行深挖。

SIMT CUDA Cores

完成了 SM 的顶层设计之后，下一步就是实现 SM 内最核心的执行部件：CUDA Cores。

简单来说，CUDA Core 是 SM 内负责通用数值计算的基础执行单元。它承担整数和浮点算术、比较、逻辑与类型转换等最常见的线程级操作，是 GPU 通用计算能力的基本来源。

在执行时，一个 warp 中的 32 个线程会在 SIMT 模型下并行推进，同一条指令会分发到多个 CUDA Cores 上执行。因此，从软件视角看，CUDA Core 不只是在做“单个标量运算”，而是在 warp 级并行中支撑起整个通用计算数据通路。早期 GPU 上的矩阵计算，本质上也是将矩阵乘法拆成大量标量乘加，再由许多 CUDA Cores 协同完成；而在后续架构中，像高吞吐矩阵乘加这样的工作负载，则进一步演进到由专门的 Tensor Core 承担。

总结来说，从设计上看，单个 CUDA Core 要做的事情非常简单，就是各种精度的标量计算。SpinalGPU 中生成CUDA Core代码核心其实就是一个大 switch，根据指令 OP 做相应的数值计算：

  // SpinalGPU written in SpinalHDL Language.  private def opResult(opcode: Bits, operandA: Bits, operandB: Bits, operandC: Bits): Bits = {    ..    switch(opcode) {      ..
      is(B(Opcode.FADD, 8 bits)) {        result := Fp32Math.add(operandA, operandB).asUInt      }      is(B(Opcode.FMUL, 8 bits)) {        result := Fp32Math.mul(operandA, operandB).asUInt      }      ..      is(B(Opcode.HADD, 8 bits)) {        result := Fp16Math.add(operandA(15 downto 0), operandB(15 downto 0)).asUInt.resize(config.dataWidth)      }      is(B(Opcode.HMUL, 8 bits)) {        result := Fp16Math.mul(operandA(15 downto 0), operandB(15 downto 0)).asUInt.resize(config.dataWidth)      }      ..      is(B(Opcode.CVTF16X2E4M3X2, 8 bits)) {        result := Fp8Format.e4m3x2ToF16x2(operandA(15 downto 0)).asUInt.resize(config.dataWidth)      }      is(B(Opcode.CVTF16X2E5M2X2, 8 bits)) {        result := Fp8Format.e5m2x2ToF16x2(operandA(15 downto 0)).asUInt.resize(config.dataWidth)      }      ..    }    result.asBits  }

也正因为CUDA Core 做的事情简单，CUDA Core 的芯片资源占用并不高。笔者的demo里面 FP32/FP16/FP8 都用了不同的芯片资源。更合理的设计应该是共享计算资源，所以正常商用的 AI 芯片里低精度计算 FLOPs 是高精度计算 FLOPs 的 N 倍。

在有 PTX Spec 的前提下，我用 AI 完善了 SpinalGPU CUDA Core 可以运行的 PTX 指令，可以运行数据搬迁、Scalar 计算指令、Vector 计算指令、以及 FP32/FP16/FP8 三种精度。支持的 PTX 指令有：

Load / Store- ld: loads data from .param, .global, or .shared into registers, so it is the main read path for kernel inputs and intermediate data.- st: stores register data back to .global or .shared, so it is the main write path for results.- mov: copies values between registers or materializes special values, so it is the basic data-shuffling/setup instruction.- cvt: converts between supported numeric formats, mainly for f16, f32, and packed FP8 teaching paths.
Scalar Floating Point- add / sub: do basic floating-point arithmetic on scalar values.- mul: does scalar floating-point multiplication.- fma: does multiply-accumulate work and is the core math op in many matrix-style kernels.- neg / abs: apply simple unary FP transforms like sign flip and absolute value.- setp: compares floating-point values and produces predicates.- selp: picks between floating-point values without branching.- min / max: support clamp-style floating-point selection.
Vector- mov.v2/v4: moves tuple-style vector data between FP registers, mainly for float2/float4 teaching kernels.- ld.v2/v4 / st.v2/v4: load or store tuple-style vector data from global memory, but the frontend lowers them into repeated scalar ops.- Packed add / mul on f16x2: operate on two packed half values in one register, which is the closest thing to native packed vector arithmetic today.
Matrix- No dedicated matrix PTX instruction family exists yet, so matrix kernels are built from scalar mov, ld/st, add, mul, fma, setp, and bra.- No mma, wmma, or tensor PTX is exposed, so “matrix” here means matrix-shaped indexing and accumulation, not special tensor instructions.
Control- bra: jumps to a label, including predicated forms, so it provides the current subset’s control flow.- ret: exits the kernel, not a full PTX function-call return model.- trap: intentionally raises a runtime fault for negative tests and fault handling.

在没有做任何性能优化的情况下，当前一个 warp（32 threads）的 FLOPS 性能大致如下（横轴是芯片的 Frequency，纵轴是 GLOPS）。按 H100 的 144 个 SM 算，在 1.8GHz 左右，总共有约 2 TFLOPs @ FP16.

Tensor Cores

完成了 CUDA Core 后，下一步自然是 Tensor Core。

细心的读者可能发现上一章节 CUDA Core 支持的 PTX 指令中并没有 matrix 相关的指令。不是 SpinalGPU 没实现，而是 PTX Spec 就没有。在 Tensor Core 出现之前，matrix kernel 都是基于 scalar/vector PTX 指令构成的。大矩阵的计算效率自然不高。这也是为什么类似 Tensor Core 这样的加速核存在的原因。

我们可以看看英伟达生态下，矩阵乘代码用 cuda 和 tensor core 写的对比：

GEMM CUDA Core 版本:

__global__ void gemm_cuda_core(const float* A,                               const float* B,                               float* C,                               int M, int N, int K) {    int row = blockIdx.y * blockDim.y + threadIdx.y;    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < M && col < N) {        float acc = 0.0f;        for (int k = 0; k < K; ++k) {            acc += A[row * K + k] * B[k * N + col];        }        C[row * N + col] = acc;    }}

.. 其生成的 PTX 核心代码为（运行在 CUDA Core 上）：

LOOP:    ld.global.f32 %f1, [A...]       ld.global.f32 %f2, [B...]    fma.rn.f32    %f0, %f1, %f2, %f0   // 标量乘    bra LOOP

做为对比， GEMM Tensor Core版本基于 wmma API 构建：

__global__ void gemm_tensor_core(const half* A,                                 const half* B,                                 float* C,                                 int M, int N, int K) {    int tile_m = blockIdx.y * 16;    int tile_n = blockIdx.x * 16;    if (tile_m >= M || tile_n >= N) return;    wmma::fragment<wmma::matrix_a, 16, 16, 16, half,  wmma::row_major> a_frag;    wmma::fragment<wmma::matrix_b, 16, 16, 16, half,  wmma::row_major> b_frag;    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;    wmma::fill_fragment(c_frag, 0.0f);    for (int k = 0; k < K; k += 16) {        wmma::load_matrix_sync(a_frag, A + tile_m * K + k, K);        wmma::load_matrix_sync(b_frag, B + k * N + tile_n, N);        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);    }    wmma::store_matrix_sync(C + tile_m * N + tile_n, c_frag, N, wmma::mem_row_major);}

.. 生成的 PTX 核心代码为：

ldmatrix...  // 加载矩阵mma.sync.aligned....  // 计算stmatrix...  // 保存矩阵

总结来说，为了加速矩阵计算，英伟达的 Tensor Core 有一些新的矩阵计算 PTX 指令。做为加速 DNN 的核心组件，Tensor Core 在近几年有多个版本的 API 更新。简洁起见，我们尝试实现最基本的几个指令：同步的ldmatrix/mma-sync/stmatrix，以及异步 tcgen05 家族。具体代码可见：

https://github.com/lastweek/SpinalGPU/blob/main/src/main/scala/spinalgpu/TensorCoreBlock.scala

SFU

完成 CUDA Core 和 Tensor Core 之后，剩下的就是 SM 中的 Special Function Unit （SFU），负责一些特殊的数学计算（Tensor Core 和 SFU 都是 special 的..）。SFU 负责的一些常用数学计算有：

sqrt.approx.f32 — fast approximate square root for positive FP32 inputs.
rsqrt.approx.f32 — fast approximate reciprocal square root, common in normalization code.
sin.approx.f32 — fast approximate sine for FP32 inputs.
cos.approx.f32 — fast approximate cosine for FP32 inputs.
lg2.approx.f32 — fast approximate base-2 logarithm.
ex2.approx.f32 — fast approximate base-2 exponential, useful as a building block for exp-like math
...

这些都是用 CUDA Core 实现性能开销大，但用专有硬件实现代价小收益高的数值运算。

总结

我们使用 Codex vibe 了 SpinalGPU，一个模仿英伟达 GPU 的学习项目。SpinalGPU 使用高级编程语言 SpinalHDL 完成，硬件代码逻辑清晰、组件化、可快速仿真。

时间问题，我只过了计算部分，内存结构方面没有做过多设计。但是从任何体系架构师的视角看: It's the memory, stupid. 如何不要让计算等待，如何让计算互通，2 个永恒的问题。有兴趣的读者可以进一步自行探索。

这个项目只是走了硬件开发最基础也最简单的一步，我们没有碰例如时序优化、布局布线等挑战的工作。希望读者们也能从中复习AI加速器的相关知识和一些硬件开发的思路。

-- 完 --