Skip to content

CnuasGPU Design

1. Overview

CnuasGPU is a virtual GPU plugged into QEMU as a PCIe device. It mirrors the NVIDIA software ecosystem (CUDA Runtime, NCCL, NVSHMEM, nvidia-smi, dcgm, Triton) using clean Cnuas naming, so students learn the real GPU programming model on emulated hardware and the same code patterns transfer to NVIDIA, AMD, or Intel GPUs.

Real World Cnuas
NVIDIA GPU CnuasGPU
NVLink CnuasLink
NVSwitch CnuasLink Switch
CUDA Cnuas Compute
PTX CnuasIR
nvcc cnuascc
nvidia-smi cnuas-smi
dcgmi / dcgm cnuas-dcgmi / cnuas-dcgm
NCCL CnuasCCL
NVSHMEM CnuasSHMEM
Nsight Compute cnuas-prof

2. Identifiers

Field Value
PCI Vendor ID 0x1AF4 (Red Hat/Qumranet, emulated; experimental)
PCI Device ID 0x10F1 (CnuasGPU)
QEMU device name cnuasgpu
Char device /dev/cnuasgpu0, /dev/cnuasgpu1, ...
sysfs class /sys/class/cnuasgpu/

3. Hardware Blocks

3.1 Compute

Block Quantity (default) Function
Streaming Multiprocessor (SM) 16 Top-level compute tile
Vector ALU lanes per SM 32 SIMD execution backing the SIMT model
Tensor MAC unit per SM 1 Matrix tile op (FP16/BF16/INT8/FP8)
Warp scheduler per SM 1 Schedules 32-thread warps
Special function unit (SFU) per SM 1 sin/cos/sqrt/rcp

Configurable via QEMU device properties: sm_count, lanes_per_sm, tensor_size.

3.2 Memory Hierarchy

Level Size (default) Latency model
Register file per SM 64 KB 1 cycle
L1 / shared memory per SM 128 KB ~30 cycles
L2 cache (chip-wide) 8 MB ~200 cycles
Device memory (HBM equiv) 8 GB ~500 cycles

Latency is emulated by QEMU clock dilation, not real wall time, so emulation stays fast.

3.3 Engines

Engine Function
Command processor Reads work submission queues from host memory
Copy engines (2) DMA H2D, D2H, D2D
MMU Per-context page tables (discrete addressing in v1, UVM in v2)
Performance counters Per-SM and chip-wide for cnuas-smi / cnuas-dcgm
CnuasLink controller Inter-GPU traffic to CnuasLink Switch

4. Programming Model

4.1 SIMT (audience-facing)

Same hierarchy as CUDA:

Grid → Block → Warp (32 threads) → Thread

Programs launched as kernels with <<<grid, block>>> style or via the runtime API call.

4.2 SIMD (emulation backend)

QEMU emulates each warp as a SIMD vector on the host CPU. Divergent control flow uses an active-mask register (predicated execution), exactly as real GPUs do. Host SIMD used:

Host arch Vector width used Backend
x86-64 256-bit AVX2 (or 512-bit AVX-512 if available) QEMU TCG vector ops
ARMv8 128-bit NEON (or SVE if available) QEMU TCG vector ops
RISC-V RVV 1.0 QEMU TCG vector ops

This means CnuasGPU emulation runs at host SIMD speed, and porting to FPGA later keeps the same SIMT-on-SIMD architecture (one DSP slice per lane).

4.3 CnuasIR Instruction Set

CnuasIR is a virtual ISA inspired by PTX but defined in terms of RISC-V Vector extension (RVV 1.0) plus Cnuas-specific tensor opcodes. This means:

  • Existing LLVM RISC-V Vector backend handles vector ops
  • Tensor ops are added as Cnuas custom instructions
  • cnuascc (the compiler) is an LLVM target with a small extension
Class Source Examples
Scalar integer RV64I add, sub, xor, lw, sw
Floating point RV F/D/Q fadd, fmul, fsqrt
Vector RVV 1.0 vadd.vv, vmul.vv, vfmacc.vv, vle32.v
Tensor (custom) Cnuas tmma, tload, tstore, tquant, tdequant
Sync (custom) Cnuas bar.sync, bar.warp, membar.gl
Cross-lane (custom) Cnuas shfl, vote.any, vote.all

Tensor ops operate on tile registers (e.g. 16x16 FP16). One tmma performs a full tile MMA.

4.4 Warps and Divergence

Warp size: 32 threads (matches CUDA). Divergence handled via per-lane active mask. No warp-stealing or hardware MIMD; standard SIMT divergence with reconvergence stack.

5. Memory Model (v1, Discrete)

5.1 Allocation

cnuasdev_error cnuasdev_malloc(void **ptr, size_t bytes);
cnuasdev_error cnuasdev_free(void *ptr);
cnuasdev_error cnuasdev_memcpy(void *dst, const void *src, size_t bytes,
                   cnuasdev_memcpy_kind kind);

kind is one of CNUASDEV_MEMCPY_HOST_TO_DEVICE, CNUASDEV_MEMCPY_DEVICE_TO_HOST, CNUASDEV_MEMCPY_DEVICE_TO_DEVICE.

5.2 Implementation

  • Device memory backed by a contiguous mmap'd region in the host (per QEMU process), exposed to guest via PCIe BAR1
  • cnuasdev_malloc allocates inside that region (linear allocator + free list)
  • cnuasdev_memcpy is just memcpy plus latency simulation
  • No paging, no migration, no fault handling

5.3 v2 Path (UVM)

UVM adds page tables, IOMMU integration, and copy-on-fault migration. Not in v1. Driver design leaves room for it (separate VMA class).

6. Host ↔ GPU Communication

6.1 PCIe BARs

BAR Size Purpose
BAR0 64 KB MMIO control registers (doorbell, command queue head/tail, IRQ)
BAR1 up to 16 GB Device memory window (configurable)
BAR2 4 KB Per-context doorbell pages

6.2 Command Submission

Standard ring buffer in host pinned memory:

  1. Host driver writes command record (kernel launch, DMA, etc) to ring
  2. Host bumps tail register
  3. QEMU emulator picks up command, executes, advances head
  4. On completion, MSI raised

6.3 IRQ Sources

IRQ Trigger
IRQ_DMA_DONE Copy engine completed
IRQ_KERNEL_DONE Kernel finished
IRQ_CNUASLINK_RX Frame arrived on CnuasLink port
IRQ_ERROR Page fault, illegal instruction, etc
IRQ_PERF_OVERFLOW Counter overflow (for cnuas-prof)

7. Software Stack

7.1 Layered View

Applications (matrix multiply, training, inference)
     |
CnuasBLAS, CnuasDNN, CnuasCCL, CnuasSHMEM, CnuasFFT, CnuasSPARSE, CnuasSOLVER
     |
Cnuas Runtime API (libcnuasrt.so)
     |
Cnuas Driver API (libcnuasdev.so)
     |
ioctl(/dev/cnuasgpuN)
     |
cnuasgpu.ko (kernel driver)
     |
QEMU cnuasgpu device (PCIe)

7.2 Kernel Driver (cnuasgpu.ko)

Responsibility Detail
PCIe probe Match vendor/device IDs
MMIO mapping Map BAR0, BAR2
Char device /dev/cnuasgpuN for userspace ioctl
Memory mgmt Pin user memory, set up DMA mappings
Context create/destroy One context per process
Command queue Per-context ring submission
IRQ handling MSI vector → completion event
sysfs Expose telemetry for cnuas-smi

7.3 Runtime / Driver Libraries

This section describes the target design. For the API that exists in the tree today, see the Cnuas Software Stack Datasheet. The shipped runtime offers a fixed set of kernels (cnuasSgemm, cnuasVectorAddF32, cnuasVectorDotF32, cnuasVectorScaleF32), plus cnuasModuleLoad and cnuasLaunchKernel, which run a CnuasIR object on a backend that can execute one. There is no context API and no triple angle bracket launch syntax, and no compiler emits CnuasIR yet, so a kernel has to be written as an instruction stream by hand.

Library Target API surface
libcnuasdev.so (Driver API) cnuasInit, cnuasCtxCreate, cnuasModuleLoad, cnuasLaunchKernel, cnuasMemAlloc, cnuasMemcpy
libcnuasrt.so (Runtime API) cnuasMalloc, cnuasMemcpy, cnuasLaunchKernel, cnuasDeviceSynchronize, kernel dispatch via <<<>>> syntax
libcnuasccrt cnuascc-emitted device runtime (printf, math, sync)

The Runtime API is the everyday one (mirrors CUDA Runtime). Driver API is the lower-level one for tools and complex use cases.

7.4 Compiler (cnuascc)

Component Function
Frontend Clang with --cnuascc-host and --cnuascc-device flags
Splitter Separates host and device code, emits two object files
Device backend LLVM RISC-V Vector + Cnuas tensor opcodes → CnuasIR object
Host backend Standard host LLVM pipeline, links in Cnuas runtime stubs
Linker Embeds CnuasIR fatbin into host ELF, runtime extracts at load

That table is the version 1 target. Version 0.1, which is what the tree builds, takes a shorter route: a self contained lexer, parser and type checker in tools/cnuascc/, reading a subset of C cut down to what CnuasIR v0.1 can express, with no Clang and no host and device splitting. The subset and the reasons for it are in src/cnuasgpu/docs/CnuasCC_Language.md.

7.5 Numerical Libraries

Library Modeled on Scope (v1)
CnuasBLAS cuBLAS GEMM, GEMV, AXPY, batched GEMM
CnuasDNN cuDNN Conv2D, MaxPool, ReLU, Softmax, BatchNorm
CnuasFFT cuFFT 1D/2D forward and inverse FFT
CnuasSPARSE cuSPARSE SpMV, SpMM, CSR/CSC formats
CnuasSOLVER cuSOLVER LU, Cholesky, SVD on dense matrices

All implemented as CnuasIR kernels callable via Runtime API.

7.6 Communication Libraries

Library Modeled on Transport
CnuasCCL NCCL Collectives (AllReduce, Broadcast, AllGather, ReduceScatter, AllToAll) over CnuasLink
CnuasSHMEM NVSHMEM One-sided put/get, atomics, signals over CnuasLink

7.7 Tools

Tool Modeled on Function
cnuas-smi nvidia-smi List GPUs, util %, memory, temp (synthetic), power (synthetic), processes
cnuas-dcgmi dcgmi Health checks, field groups, policy management
cnuas-dcgm dcgm daemon Continuous telemetry, REST/gRPC API
cnuas-prof Nsight Compute Per-kernel metrics, occupancy analysis
hi-cuda-gdb style tool cuda-gdb Device-side debugging via PCIe register tap

7.8 Container Toolkit

cnuas-container-toolkit mirrors nvidia-container-toolkit. Wraps Docker runtime to inject /dev/cnuasgpuN, libcnuasdev.so, runtime libs into containers on --gpus flag.

The CnuasLink fabric is described in detail in CnuasLink Switch Design.

Summary as it relates to CnuasGPU:

  • Each CnuasGPU has a CnuasLink endpoint (4 lanes by default)
  • Lanes appear as MMIO mailboxes in BAR0
  • CnuasCCL collectives use CnuasLink as primary transport
  • If two CnuasGPUs are on the same host, ivshmem provides direct GPU↔GPU mapping
  • If on different hosts (or for >1 hop), traffic goes through cnuasgpu-link-switchd

9. Telemetry

Each CnuasGPU exposes the following counters via sysfs and OTel:

Counter Unit
sm_active_cycles per-SM cycles
sm_busy_pct percent
memory_used bytes
memory_bandwidth bytes/sec
tensor_ops count
vector_ops count
pcie_rx_bytes / pcie_tx_bytes bytes
cnuaslink_rx_bytes / cnuaslink_tx_bytes bytes
temperature_c (synthetic) celsius
power_w (synthetic) watts

Synthetic values driven by load (e.g. temp = 30 + 50 * sm_busy_pct) so cnuas-smi shows realistic-looking output.

10. Phasing

Phase Scope Outcome
1 QEMU cnuasgpu PCIe device, basic MMIO, BAR1 device memory lspci shows device
2 cnuasgpu.ko kernel driver, char device, ioctl plumbing /dev/cnuasgpu0 appears
3 libcnuasdev.so + libcnuasrt.so minimal: alloc, memcpy, launch cnuasMalloc + memcpy works
4 cnuascc compiler, CnuasIR codegen, vector add kernel Vector add passes
5 Tensor MMA emulation, tmma opcode GEMM kernel passes
6 CnuasBLAS subset (sgemm, hgemm, axpy) Throughput numbers
7 cnuas-smi, cnuas-dcgmi tools Telemetry visible
8 CnuasDNN subset (conv2d, pool, relu, softmax) Run small CNN
9 CnuasLink endpoint + CnuasLink Switch + CnuasCCL AllReduce across N GPUs
10 CnuasSHMEM One-sided put/get works
11 cnuas-prof, container toolkit Profile a kernel inside container

UVM, RT cores, video codecs, MIG, confidential compute deferred to v2.

11. Directory Structure (planned)

src/
  cnuasgpu/
    qemu/                      # QEMU device source (lives in qemu/hw/misc/ via submodule fork)
    driver/                    # cnuasgpu.ko Linux module
    runtime/                   # libcnuasrt.so, libcnuasdev.so
    compiler/                  # cnuascc wrapper, CnuasIR backend (LLVM)
    libs/
      cnuasblas/
      cnuasdnn/
      cnuasfft/
      cnuasccl/
      cnuasshmem/
    tools/
      cnuas-smi/
      cnuas-dcgmi/
      cnuas-dcgm/
      cnuas-prof/
    container-toolkit/
docs/
  CnuasGPU_Design.md       # this file
  CnuasLink_Switch_Design.md

12. Open Questions

  1. Container toolkit scope: full OCI runtime hook, or thin wrapper?
  2. Profiler depth: just counters, or also instruction-level sampling (which needs QEMU TCG plugin work)?
  3. Driver licensing: GPL-only kernel driver (for EXPORT_SYMBOL_GPL access), or dual MIT/GPL like our existing modules?
  4. CnuasIR stability: freeze ABI at v1, or version it and allow per-card capability bits?