HiCAIN HiGPU Design¶

1. Overview¶

HiGPU 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 HiCAIN 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	HiCAIN
NVIDIA GPU	HiGPU
NVLink	HiLink
NVSwitch	HiLink Switch
CUDA	HiCAIN Compute (HCC)
PTX	HiIR
nvcc	hcc
nvidia-smi	hi-smi
dcgmi / dcgm	hi-dcgmi / hi-dcgm
NCCL	HiCCL
NVSHMEM	HiSHMEM
Nsight Compute	hi-prof

2. Identifiers¶

Field	Value
PCI Vendor ID	`0x1ED5` (HiCAIN)
PCI Device ID	`0xCA20` (HiGPU)
QEMU device name	`higpu`
Char device	`/dev/higpu0`, `/dev/higpu1`, ...
sysfs class	`/sys/class/higpu/`

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 hi-smi / hi-dcgm
HiLink controller	Inter-GPU traffic to HiLink 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 HiGPU 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 HiIR Instruction Set¶

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

Existing LLVM RISC-V Vector backend handles vector ops
Tensor ops are added as HiCAIN custom instructions
hcc (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)	HiCAIN	`tmma`, `tload`, `tstore`, `tquant`, `tdequant`
Sync (custom)	HiCAIN	`bar.sync`, `bar.warp`, `membar.gl`
Cross-lane (custom)	HiCAIN	`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¶

hi_error hi_malloc(void **ptr, size_t bytes);
hi_error hi_free(void *ptr);
hi_error hi_memcpy(void *dst, const void *src, size_t bytes,
                   hi_memcpy_kind kind);

kind is one of HI_MEMCPY_HOST_TO_DEVICE, HI_MEMCPY_DEVICE_TO_HOST, HI_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
hi_malloc allocates inside that region (linear allocator + free list)
hi_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:

Host driver writes command record (kernel launch, DMA, etc) to ring
Host bumps tail register
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_HILINK_RX`	Frame arrived on HiLink port
`IRQ_ERROR`	Page fault, illegal instruction, etc
`IRQ_PERF_OVERFLOW`	Counter overflow (for hi-prof)

7. Software Stack¶

7.1 Layered View¶

Applications (matrix multiply, training, inference)
     |
HiBLAS, HiDNN, HiCCL, HiSHMEM, HiFFT, HiSPARSE, HiSOLVER
     |
HiCAIN Runtime API (libhicart.so)
     |
HiCAIN Driver API (libhi.so)
     |
ioctl(/dev/higpuN)
     |
higpu.ko (kernel driver)
     |
QEMU higpu device (PCIe)

7.2 Kernel Driver (`higpu.ko`)¶

Responsibility	Detail
PCIe probe	Match vendor/device IDs
MMIO mapping	Map BAR0, BAR2
Char device	`/dev/higpuN` 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 hi-smi

7.3 Runtime / Driver Libraries¶

Library	API surface
`libhi.so` (Driver API)	hiInit, hiCtxCreate, hiModuleLoad, hiLaunchKernel, hiMemAlloc, hiMemcpy
`libhicart.so` (Runtime API)	hicaMalloc, hicaMemcpy, hicaLaunchKernel, hicaDeviceSynchronize, kernel dispatch via `<<<>>>` syntax
`libhccrt`	hcc-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 (hcc)¶

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

7.5 Numerical Libraries¶

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

All implemented as HiIR kernels callable via Runtime API.

7.6 Communication Libraries¶

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

7.7 Tools¶

Tool	Modeled on	Function
`hi-smi`	nvidia-smi	List GPUs, util %, memory, temp (synthetic), power (synthetic), processes
`hi-dcgmi`	dcgmi	Health checks, field groups, policy management
`hi-dcgm`	dcgm daemon	Continuous telemetry, REST/gRPC API
`hi-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¶

hicain-container-toolkit mirrors nvidia-container-toolkit. Wraps Docker runtime to inject /dev/higpuN, libhi.so, runtime libs into containers on --gpus flag.

8. Multi-GPU and HiLink¶

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

Summary as it relates to HiGPU:

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

9. Telemetry¶

Each HiGPU 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
hilink_rx_bytes / hilink_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 hi-smi shows realistic-looking output.

10. Phasing¶

Phase	Scope	Outcome
1	QEMU higpu PCIe device, basic MMIO, BAR1 device memory	`lspci` shows device
2	`higpu.ko` kernel driver, char device, ioctl plumbing	`/dev/higpu0` appears
3	`libhi.so` + `libhicart.so` minimal: alloc, memcpy, launch	`hicaMalloc` + memcpy works
4	hcc compiler, HiIR codegen, vector add kernel	Vector add passes
5	Tensor MMA emulation, `tmma` opcode	GEMM kernel passes
6	HiBLAS subset (sgemm, hgemm, axpy)	Throughput numbers
7	hi-smi, hi-dcgmi tools	Telemetry visible
8	HiDNN subset (conv2d, pool, relu, softmax)	Run small CNN
9	HiLink endpoint + HiLink Switch + HiCCL	AllReduce across N GPUs
10	HiSHMEM	One-sided put/get works
11	hi-prof, container toolkit	Profile a kernel inside container

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

11. Directory Structure (planned)¶

src/
  higpu/
    qemu/                      # QEMU device source (lives in qemu/hw/misc/ via submodule fork)
    driver/                    # higpu.ko Linux module
    runtime/                   # libhicart.so, libhi.so
    compiler/                  # hcc wrapper, HiIR backend (LLVM)
    libs/
      hiblas/
      hidnn/
      hifft/
      hiccl/
      hishmem/
    tools/
      hi-smi/
      hi-dcgmi/
      hi-dcgm/
      hi-prof/
    container-toolkit/
docs/
  HiCAIN_HiGPU_Design.md       # this file
  HiCAIN_HiLink_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_GPL access), or dual MIT/GPL like our existing modules?
HiIR stability: freeze ABI at v1, or version it and allow per-card capability bits?