[ GRX930 · G100 · GRXCP ]

Open Silicon.
No NDAs.
No Licensing Fees.

Gridworkz is building the Linux of chip design: a quad-core RISC-V SoC with dual NPUs, an open RISC-V GPU, and a photonic tensor accelerator in design, all on GitHub. Built for university labs and research teams.

  • Apache 2.0
  • RISC-V
  • SystemVerilog
  • Icarus Verilog
  • Verilator
~/grx930 — bash
# 1. Install the open-source simulator
$ sudo apt install iverilog
# 2. Clone the GRX930 RISC-V SoC
$ git clone https://github.com/kubeworkz/grx930.git
$ cd grx930/c930
# 3. Run the NPU GEMM regression (~1 min)
$ make npu
Illustrative waveformWaveform: clock, valid and data bus signalsclkvaliddata
Zero-cost simulation

From git clone to a passing
NPU test in minutes.

No NDAs and no EDA licenses to simulate. The test suites run on open-source simulators, on your laptop.

Simulate

52 testbenches, ready to run

GRX930 ships with Icarus Verilog testbenches, equivalence gates and formal checks: make npu, make soc, make quad. The G100 runs on its SimX and RTL simulators.

Program

CUDA-style kernels

GRXCP gives you __global__ kernels, <<<>>> launches and a CUDA-shaped runtime for the G100, with a header that maps cudaX to grxX.

Synthesize

Timing closed on FPGA

A Vivado flow and pin constraints for the Nexys Video (Artix-7 200T). The full quad-core SoC closes timing at ~121 MHz with zero failing paths.

Architecture

Every block, exposed.

Select a block to see the source behind it, its docs and how far it has been proven.

GRIDWORKZ PLATFORM

Select a block to view its source, docs and status →

The open stack

Layered like an operating system.

Layer 0 · Base

Unhindered bare-metal compute.

GRX930: four RV64IMAC cores with caches, coherent L2 and an AXI4 fabric. No microcode to hide behind. Change the pipeline, add instructions, rerun the test suite.

Layer 1 · Acceleration

Accessible tensor workloads.

Two systolic NPUs (INT4 to BF16) and the G100 GPU with tensor cores (fp8 to fp32), both with every register documented and every line of RTL readable.

Layer 2 · Frontier

Photonic matrix multiply, in design.

A 128 × 64 photonic GEMM tile, planned as a chiplet beside the G100. It's being designed against an error model first, so its accuracy is known before anything is fabricated.

Layer 3 · Software

Programmable hardware. Zero assembly required.

Write parallel C++ with GRXCP, built on the VOLT LLVM toolchain. If you can write CUDA, you can write kernels for the G100.

__global__ void axpy(const float* x, float* y, unsigned int n, float a) {
  const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i] + y[i];
}
GRXCP docs →
Documentation

Read the registers before you write the grant.

RTL architecture, register maps, pinouts, toolchain setup and errata, all public and versioned with the code.

GRX930 · NPU0 CSR0x4000_0000
OffsetRegisterAccess
0x00CTRL · STARTW
0x04STATUS · BUSY DONE ERRORR
0x08–0x10DIM_M / DIM_N / DIM_KRW
0x14–0x1CA_BASE / B_BASE / C_BASERW
0x20PREC · INT8…INT4RW
For labs

Request an evaluation board for your lab.

Ready to move from simulation to hardware? GRX930 runs on the Nexys Video (Artix-7 200T). Tell us about your research and we'll help you get set up.

Allocation details and grant datasheet →

Stop buying black boxes.
Start building the future of hardware.

Fork the code, run the test suites in open-source simulators, or build GRX930 for an FPGA board and test your ideas on real hardware.