CnuasGPU Design¶
1. Overview¶
CnuasGPU is a virtual AI and ML accelerator that runs in either of two forms: as a Soft-GPU in an ordinary host process, with no hypervisor, guest or kernel module, and as a PCIe Gen5 x16 device plugged into QEMU for guest-side and rack deployments. Both forms execute the same kernels through the same compute backends. See CnuasGPU as a Soft-GPU for the standalone form.
It mirrors the NVIDIA software ecosystem (CUDA Runtime, NCCL, NVSHMEM, nvidia-smi, dcgm, Triton) using clean Cnuas naming, so the real GPU programming model is learned on emulated hardware and the same code patterns transfer to NVIDIA, AMD, or Intel GPUs.
1.1 Research and implementation boundary¶
CnuasGPU is a vertically open co-design target: researchers can modify the device model, CnuasIR, compiler, compute backends, kernel interface, runtime, numerical and collective libraries, and CnuasLink integration rather than being limited to the software surface of a fixed vendor accelerator.
The FPGA path is an experimental implementation in progress. Current evidence covers functional RTL simulation of an eight-lane FP32 descriptor engine for vector addition, scaling, SAXPY and matrix multiplication. It is not a timing-accurate FPGA model, and timing closure, resource utilization, released bitstreams and board operation have not been established. Architectural blocks described below must therefore be read alongside their implementation status; their presence in the target architecture does not itself establish executable hardware.
| 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 |
|---|---|
| Host bus | PCIe Gen5 x16, 32 GT/s per lane |
| 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 |
| TVP (Tile Vector Processor) lanes per SM | 32 | SIMD execution backing the SIMT model, activations and elementwise work, and the array every primitive below runs on |
| TTU (Tile Tensor Unit) per SM | 1 | Matrix by matrix multiply with direct accumulation, 16 by 16 tile edge (FP16/BF16/INT8/FP8) |
| TCP (Tile Control Processor) per SM | 1 | Pointer arithmetic, address generation and loop control |
| LSU (Load Store Unit) per SM | 1 | Operand movement between device memory and the tile |
| CCP (Cluster Control Processor) per device | 1 | Reads work submission queues and hands descriptors to tiles |
| 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.
The math the vector array carries is one level 1 group (add, scale, dot), a
level 2 general matrix vector product, two level 3 matrix multiplies (AB and
AB'), a reduction, five elementwise activations and a row-wise softmax. Each
is a cnuas_compute_ops entry point implemented once per backend, and a
backend that does not implement one does not advertise its capability bit, so
a caller is refused rather than given a wrong answer. FP32 is the only format
any primitive accepts; the other formats above are architectural. The tensor
MAC unit is reached only from inside the matrix multiply. See the
CnuasGPU datasheet section 4.
3.1.1 Tile blocks¶
An SM is a tile, and the device is the cluster of tiles. The four tile blocks carry the names the wider accelerator field uses, because the same four appear in most inference parts and one vocabulary then covers the emulated device, the fabric implementation and the parts a reader already knows.
| Block | Status | Where it is today |
|---|---|---|
| TVP | Implemented | The 32 vector lanes. Every level 1 primitive, every activation and the softmax run here, and so does the arithmetic inside the matrix multiply |
| CCP | Implemented | The command processor, reached through the BAR2 doorbell pages |
| LSU | Implemented | Inside the primitives, not separately issued |
| TTU | Architectural | Accumulation into a destination tile happens inside sgemm, and the AVX-512 backend blocks with a 16 by 16 C tile, but no caller can address a tile register file and tmma has no implementation |
| TCP | Partial | Scalar registers, address generation and branches exist in the CnuasIR ISA and the interpreter executes them. There is no block a caller addresses, and no separate instruction stream per tile |
Naming the blocks is not cosmetic. A block with a name and an interface can be given RTL, a capability bit and a test; a block that exists only as a sentence in a specification cannot. The two blocks that are not implemented are exactly the two that have no caller-visible interface, and closing that is the same piece of work as section 3.1.2.
TTU operand formats¶
TF32 and BF16 are specified as the TTU operand formats for training work, with
FP32 accumulation in both cases, and 2:4 structured sparsity is specified as a
property of the weight operand. CNUAS_CCAP_GEMM_TF32, CNUAS_CCAP_GEMM_BF16
and CNUAS_CCAP_SPARSE_2_4 are reserved in the backend ABI and advertised by
no backend.
These sit behind the tile primitive rather than beside it. A narrowed operand
needs a unit to be routed to, and a sparse operand needs an entry point that
takes index metadata alongside the values, so neither can be reached while the
TTU is only implicit inside sgemm. The order of work is therefore the tile
primitive first, then TF32 on the emulated backends, then BF16 with its own
entry point and conversion, then the 2:4 metadata layout.
3.1.2 CnuasGPU FPGA IP¶
The compute blocks above, as synthesisable RTL on real fabric, with a RISC-V host. The reference target is the Microchip PolarFire SoC on the Icicle Kit, chosen because it carries PCIe hard IP and four Linux-capable RISC-V cores on one die, so the board is a complete system rather than a fabric that needs a host beside it.
| Resource | Icicle Kit value | What it carries |
|---|---|---|
| Fabric | MPFS250T, 254 K logic elements | The tile blocks |
| Math blocks | 784 by 18 by 18 MACC | The TVP lanes, about 64 of them for the v0.1 8-lane vector, and the TTU once it has a caller-visible interface |
| LSRAM | 17.6 Mb | Vector register file and the device memory window |
| Hard processors | 4 by SiFive U54 RV64GC, 1 by E51 monitor | The CCP and the TCP, as a firmware daemon under Linux |
| Host bus | PCIe Gen2 x4 hard IP | The endpoint, in place of the emulated Gen5 x16 |
Nothing above libcnuasdev changes, because the board is another device
backend. The emulated device becomes the reference model the RTL is checked
against, which is the practical reason to have built it first. Fabric
parameters, the RTL module list, the descriptor format and the bring-up phases
are in the CnuasGPU FPGA design note in the component repository.
3.1.3 Planned attachments¶
UALink is planned as an open accelerator to accelerator interconnect alongside PCIe, a role CnuasLink fills today.
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) | Configurable, 256 MiB default | ~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:
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.
This table is the CnuasIR specification, not a list of what executes today. No
backend implements the tensor opcodes: matrix multiply reaches the accelerator
through sgemm, which accumulates but exposes no tile register file. See the
CnuasGPU datasheet section 3.
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_mallocallocates inside that region (linear allocator + free list)cnuasdev_memcpyis justmemcpyplus 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 | 16 MiB to 1 TiB | Device memory window, 64-bit prefetchable, configurable |
| BAR2 | 4 KB | Per-context doorbell pages |
6.2 Command Submission¶
Standard ring buffer in host pinned memory:
- Host driver writes command record (kernel launch, DMA, etc) to ring
- Host bumps
tailregister - QEMU emulator picks up command, executes, advances
head - 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. cnuascc compiles a device source file into such an
object, so a kernel no longer has to be written as an instruction stream by
hand. There is still no context API and no triple angle bracket launch syntax,
and code generation covers the scalar subset only, so a kernel that needs the
vector unit is written 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.
8. Multi-GPU and CnuasLink¶
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¶
- Container toolkit scope: full OCI runtime hook, or thin wrapper?
- Profiler depth: just counters, or also instruction-level sampling (which needs QEMU TCG plugin work)?
- Driver licensing: GPL-only kernel driver (for
EXPORT_SYMBOL_GPLaccess), or dual MIT/GPL like our existing modules? - CnuasIR stability: freeze ABI at v1, or version it and allow per-card capability bits?