feature: Hy3 support

.github/workflows/check.yml

@@ -16,6 +16,8 @@ jobs:
       - uses: actions/checkout@v4
       - name: make check
         run: make -C c check
+      - name: Build hy3 engine
+        run: make -C c hy3
 
   windows:
     # The job that would have caught #68/#137 pre-merge: native MinGW-w64
@@ -37,6 +39,9 @@ jobs:
       - name: make check
         run: make -C c check
 
+      - name: Build hy3 engine
+        run: make -C c hy3
+
       # The Windows engine as this branch builds it. Small, and it makes an
       # antivirus false-positive report (#527) verifiable on a PR instead of
       # only after a release is tagged and published.
@@ -55,3 +60,5 @@ jobs:
         run: brew install libomp
       - name: make check
         run: make -C c check
+      - name: Build hy3 engine
+        run: make -C c hy3

.github/workflows/ci.yml

@@ -12,8 +12,8 @@ jobs:
     runs-on: ubuntu-latest
     steps:
       - uses: actions/checkout@v4
-      - name: Build colibri + inkling
-        run: cd c && make colibri inkling
+      - name: Build engines
+        run: cd c && make colibri hy3 inkling
       - name: C test suite
         run: cd c && make test-c
 
@@ -59,7 +59,7 @@ jobs:
         run: |
           cd c
           rc=0
-          for t in colibri inkling kimi_k3 olmoe; do
+          for t in colibri inkling kimi_k3 olmoe hy3; do
             echo "::group::$t"
             make $t || { echo "FAILED: $t"; rc=1; }
             echo "::endgroup::"
@@ -122,6 +122,29 @@ jobs:
           python3 tools/make_tiny_inkling.py tiny_inkling
           SNAP=tiny_inkling ./inkling 8 0 tiny_inkling/ref_inkling.json
 
+  hy3-oracle:
+    name: Hy3 oracle (teacher-forcing)
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+      - uses: actions/setup-python@v5
+        with:
+          python-version: '3.12'
+          cache: pip
+          cache-dependency-path: c/tools/oracle-requirements.txt
+      - name: Install torch (CPU) + transformers
+        run: pip install -r c/tools/oracle-requirements.txt
+      - name: Build hy3
+        run: cd c && make hy3
+      - name: Tiny-model fixture + teacher-forcing oracle (fp32, 32/32 positions)
+        run: |
+          cd c
+          # Random-init HYV3ForCausalLM via HF transformers (see make_hy3_oracle.py).
+          # Weights are not committed (*.safetensors is gitignored); regenerate here
+          # like inkling-oracle does for tiny_inkling/.
+          python3 tools/make_hy3_oracle.py
+          SNAP=./hy3_tiny TF=1 ./hy3 64 16 16
+
   engine-hip-syntax:
     name: HIP syntax check
     runs-on: ubuntu-latest

.github/workflows/site.yml

@@ -22,6 +22,9 @@ concurrency:
 
 jobs:
   deploy:
+    # GitHub Pages is configured on the upstream repo only; forks lack a Pages site
+    # and configure-pages fails with HTTP 404 until one is set up manually.
+    if: github.repository == 'JustVugg/colibri'
     runs-on: ubuntu-latest
     environment:
       name: github-pages

.gitignore

@@ -18,6 +18,8 @@ c/olmoe
 c/inkling
 c/kimi_k3
 c/kimi_k3.exe
+c/hy3
+c/hy3.exe
 c/olmoe.exe
 c/iobench
 c/iobench.exe

README.md

@@ -20,7 +20,8 @@ hierarchy.
 
 Four families run today: **GLM-5.2** (744B), **Inkling** (975B), **Kimi K3**
 (2.8T) and **OLMoE** (7B) — one C file each, the same `coli chat` /
-`coli serve` / `coli web` front end. [Full roster ↓](#other-supported-models)
+`coli serve` / `coli web` front end. **Hy3** (Tencent, 295B) is also supported
+as a sibling engine. [Full roster ↓](#other-supported-models)
 
 > **Colibrì is an inference engine you can run today, and an open research
 > platform.** Its primary goal is to pursue inference-side performance across
@@ -368,6 +369,7 @@ the model's `config.json`):
 | **Inkling** (Thinking Machines) | 975B / 41B | [`nbeerbower/Inkling-colibri-int4`](https://huggingface.co/nbeerbower/Inkling-colibri-int4) (469 GB) | `make -C c inkling` | [inkling.md](docs/inkling.md) |
 | **Kimi K3** (Moonshot) | 2.8T / 104B | [`moonshotai/Kimi-K3`](https://huggingface.co/moonshotai/Kimi-K3) — original checkpoint, routed experts stay **native MXFP4** | `make -C c kimi_k3` | [kimi_k3.md](docs/kimi_k3.md) |
 | **OLMoE** (AI2) | 7B / 1B | converted with `c/tools/convert_olmoe_merged.py` | `make -C c olmoe` | — |
+| **Hy3** (Tencent) | 295B / 21B | [UnderstandLing/Hy3-colibri-int4](https://huggingface.co/UnderstandLing/Hy3-colibri-int4) (~142 GB) | `make -C c hy3` | [hy3.md](docs/hy3.md) |
 
 Kimi K3 needs no conversion: its QAT-trained MXFP4 experts are streamed straight from
 the original Hugging Face shards, and the bf16 dense set is quantized at load time.
@@ -403,10 +405,12 @@ changes between models**. Build the engine you want once, then just point
 make -C c glm                                     # GLM-5.2
 make -C c inkling                                 # Inkling
 make -C c kimi_k3                                 # Kimi K3
+make -C c hy3                                     # Hy3
 
 COLI_MODEL=/nvme/glm52_i4      ./coli chat        # TUI
 COLI_MODEL=/nvme/inkling_i4    ./coli chat
 COLI_MODEL=/nvme/kimi_k3       ./coli chat
+COLI_MODEL=/nvme/hy3_i4        ./coli chat
 
 ./coli web --model /nvme/inkling_i4               # API + dashboard, same port
 ./coli web --model /nvme/kimi_k3

c/Makefile

@@ -111,6 +111,25 @@ EXE     =
 endif
 endif
 
+# Optional io_uring async expert reads (PIPE=2). Set IOURING_PREFIX to a liburing
+# install prefix, or IOURING=1 to auto-probe /usr and /tmp/liburing-install.
+IOURING       ?= 0
+IOURING_PREFIX ?=
+IOURING_CFLAGS  =
+IOURING_LDFLAGS =
+ifeq ($(IOURING),1)
+ifneq ($(IOURING_PREFIX),)
+IOURING_CFLAGS  = -DCOLI_IOURING -I$(IOURING_PREFIX)/include
+IOURING_LDFLAGS = -L$(IOURING_PREFIX)/lib -luring
+else ifneq ($(wildcard /tmp/liburing-install/include/liburing.h),)
+IOURING_CFLAGS  = -DCOLI_IOURING -I/tmp/liburing-install/include
+IOURING_LDFLAGS = -L/tmp/liburing-install/lib -Wl,-rpath,/tmp/liburing-install/lib -luring
+else ifneq ($(wildcard /usr/include/liburing.h),)
+IOURING_CFLAGS  = -DCOLI_IOURING
+IOURING_LDFLAGS = -luring
+endif
+endif
+
 # --- install ---
 PREFIX     ?= /usr/local
 BINDIR     ?= $(PREFIX)/bin
@@ -182,6 +201,10 @@ endif
 # like compat.h does for Windows). HIP_ARCH=native targets the GPU in this
 # machine; set an explicit arch (e.g. HIP_ARCH=gfx1201) when distributing.
 HIP       ?= 0
+ROCM      ?= 0
+ifeq ($(ROCM),1)
+HIP := 1
+endif
 ROCM_HOME ?= /opt/rocm
 HIPCC     ?= $(ROCM_HOME)/bin/hipcc
 HIP_ARCH  ?= native
@@ -332,11 +355,16 @@ VK_OBJ    = backend_vulkan.o
 VK_SPV    = shaders/qmatmul.spv shaders/qmatmul_gate_up.spv shaders/attention_absorb.spv shaders/rmsnorm.spv
 endif
 
+HY3_INPUT = hy3.c
+
 all: colibri$(EXE)
 
 # phony targets — 'glm' kept for backward compatibility
 colibri: colibri$(EXE)
 glm: colibri$(EXE)
+ifneq ($(EXE),)
+hy3: hy3$(EXE)
+endif
 
 # Config stamp: make only tracks file timestamps, not flag changes. Without this,
 # `make colibri.exe CUDA_DLL=1` after a prior CPU-only build reports "up to date"
@@ -345,7 +373,7 @@ glm: colibri$(EXE)
 # rewrite it ONLY when they change (evaluated here at parse time, so the file's
 # timestamp moves exactly when the config moves). The binary and CUDA/loader objects depend
 # on it, so they relink on a config change and stay put otherwise. (#306)
-BUILD_CONFIG     := $(CC)|$(CFLAGS)|$(LDFLAGS)|CUDA=$(CUDA)|CUDA_DLL=$(CUDA_DLL)|ARCH=$(ARCH)|CUDA_ARCH=$(CUDA_ARCH)|METAL=$(METAL)|HIP=$(HIP)|HIP_ARCH=$(HIP_ARCH)|VK=$(VK)
+BUILD_CONFIG     := $(CC)|$(CFLAGS)|$(LDFLAGS)|CUDA=$(CUDA)|CUDA_DLL=$(CUDA_DLL)|ARCH=$(ARCH)|CUDA_ARCH=$(CUDA_ARCH)|METAL=$(METAL)|HIP=$(HIP)|HIP_ARCH=$(HIP_ARCH)|VK=$(VK)|IOURING=$(IOURING)
 BUILD_CONFIG_OLD := $(shell cat .build-config 2>/dev/null)
 ifneq "$(BUILD_CONFIG)" "$(BUILD_CONFIG_OLD)"
 # $(file ...) writes via make's own primitive (GNU Make >= 4.0, 2013), NOT by
@@ -368,6 +396,9 @@ shaders/%.spv: shaders/%.comp
 	@command -v $(GLSLC) >/dev/null 2>&1 || { echo "glslc not found: install shaderc (for VK=1)" >&2; exit 1; }
 	$(GLSLC) --target-env=vulkan1.2 $< -o $@
 
+hy3$(EXE): $(HY3_INPUT) st.h json.h compat.h uring.h $(CUDA_OBJ) .build-config
+	$(CC) $(CFLAGS) $(IOURING_CFLAGS) $(HY3_INPUT) $(CUDA_OBJ) -o hy3$(EXE) $(LDFLAGS) $(IOURING_LDFLAGS)
+
 # Windows runtime loader object: resolves coli_cuda_* from coli_cuda.dll.
 backend_loader.o: backend_loader.c backend_cuda.h compat.h .build-config
 	$(CC) $(CFLAGS) -c backend_loader.c -o $@
@@ -489,6 +520,9 @@ portable:
 iobench$(EXE): iobench.c compat.h
 	$(CC) $(CFLAGS) iobench.c -o iobench$(EXE) $(LDFLAGS)
 
+tests/test_serve_sentinel$(EXE): tests/test_serve_sentinel.c compat.h
+	$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+
 tests/test_json$(EXE): tests/test_json.c json.h
 	$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
 
@@ -674,12 +708,13 @@ check:
 	$(MAKE) portable
 	$(MAKE) test
 
-install: colibri$(EXE) olmoe$(EXE)
+install: colibri$(EXE) hy3$(EXE) olmoe$(EXE)
 	$(INSTALL) -d $(DESTDIR)$(BINDIR)
 	$(INSTALL) -d $(DESTDIR)$(LIBEXECDIR)
 	$(INSTALL) -d $(DESTDIR)$(LIBEXECDIR)/tools
 	$(INSTALL) -m 755 coli $(DESTDIR)$(BINDIR)/coli
 	$(INSTALL) -m 755 colibri$(EXE) $(DESTDIR)$(LIBEXECDIR)/colibri$(EXE)
+	$(INSTALL) -m 755 hy3$(EXE) $(DESTDIR)$(LIBEXECDIR)/hy3$(EXE)
 	$(INSTALL) -m 755 olmoe$(EXE) $(DESTDIR)$(LIBEXECDIR)/olmoe$(EXE)
 	$(INSTALL) -m 644 resource_plan.py doctor.py autotune.py openai_server.py version.py $(DESTDIR)$(LIBEXECDIR)/
 	$(INSTALL) -m 644 tools/*.py $(DESTDIR)$(LIBEXECDIR)/tools/
@@ -699,7 +734,8 @@ uninstall:
 
 clean:
 	$(PYTHON) tools/clean.py
+	rm -f hy3$(EXE)
 
 bench: iobench$(EXE)
 	@if [ -n "$(ARGS)" ]; then ./iobench$(EXE) $(ARGS); else echo "built iobench$(EXE) — run: ./iobench$(EXE) <file> <MB> <iters> <threads> <direct 0|1>"; fi
-.PHONY: all colibri glm iq3 rans fuzz-rans cuda-test hip-test gpu-compile cuda-bench cuda-dll portable test-c test-python test check clean install uninstall bench
+.PHONY: all colibri glm hy3 iq3 rans fuzz-rans cuda-test hip-test gpu-compile cuda-bench cuda-dll portable test-c test-python test check clean install uninstall bench

c/backend_cuda.cu

@@ -62,6 +62,7 @@ typedef struct {
     float *host_x,*host_y,*host_kv; size_t host_x_cap,host_y_cap,host_kv_cap;
     float *aq,*al,*ar,*ac; size_t aq_cap,al_cap,ar_cap,ac_cap;
     float *pipe_buf[27]; size_t pipe_cap[27];   /* scratch persistenti del resident pipeline */
+    float *gq,*gk,*gv,*gctx,*gsc; size_t gq_cap,gk_cap,gv_cap,gctx_cap,gsc_cap;
     cudaStream_t stream;
     cudaEvent_t ev_done; int ev_done_ok;        /* resident-group issue completion (#431 PR-C0) */
     void *group_desc; size_t group_desc_cap;
@@ -105,7 +106,7 @@ static double ans_now_s(){
 
 static int cuda_ok(cudaError_t err, const char *what) {
     if (err == cudaSuccess) return 1;
-    std::fprintf(stderr, "[CUDA] %s: %s\n", what, cudaGetErrorString(err));
+    std::fprintf(stderr, COLI_ACCEL_TAG " %s: %s\n", what, cudaGetErrorString(err));
     (void)cudaGetLastError();   /* consume the sticky error: a failed call must
                                    not poison the next launch's error check */
     return 0;
@@ -597,6 +598,40 @@ __global__ static void attention_absorb_kernel(float *ctx,const float *q,const f
         for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k)*absorb_scale(wscale,fmt,gs,ng,row,k);ctx[(size_t)h*V+v]=a;}
 }
 
+#define GQA_MAX_NT 8192
+
+__global__ static void gqa_attn_kernel(float *ctx,const float *q,const float *k_cache,const float *v_cache,
+                                       float *scores,int S,int H,int Hkv,int hd,int st0,int pos_base,
+                                       int max_t,float scale,int nrep){
+    int h=blockIdx.x,s=blockIdx.y;
+    if(s>=S)return;
+    int kvh=h/nrep,pos=pos_base+s,nt=pos+1-st0;
+    if(nt<1||nt>GQA_MAX_NT)return;
+    float *sc=scores+((size_t)s*H+h)*GQA_MAX_NT;
+    const float *qv=q+((size_t)s*H+h)*hd;
+    for(int jj=threadIdx.x;jj<nt;jj+=blockDim.x){
+        int t=st0+jj; const float *kv=k_cache+((size_t)kvh*max_t+t)*hd;
+        float dot=0; for(int d=0;d<hd;d++) dot+=qv[d]*kv[d];
+        sc[jj]=dot*scale;
+    }
+    __syncthreads();
+    if(!threadIdx.x){
+        float mx=sc[0]; for(int i=1;i<nt;i++) mx=fmaxf(mx,sc[i]);
+        float sum=0; for(int i=0;i<nt;i++){ sc[i]=expf(sc[i]-mx); sum+=sc[i]; }
+        float inv=sum>0?1.f/sum:0.f; for(int i=0;i<nt;i++) sc[i]*=inv;
+    }
+    __syncthreads();
+    float *cx=ctx+((size_t)s*H+h)*hd;
+    for(int d=threadIdx.x;d<hd;d+=blockDim.x){
+        float acc=0;
+        for(int jj=0;jj<nt;jj++){
+            int t=st0+jj; const float *vv=v_cache+((size_t)kvh*max_t+t)*hd;
+            acc+=sc[jj]*vv[d];
+        }
+        cx[d]=acc;
+    }
+}
+
 __global__ static void attention_absorb_batch_kernel(float *ctx,const float *q,
         const float *latent,const float *rope,const void *weights,const float *wscale,
         int fmt,int S,int H,int Q,int R,int V,int K,int T,float scale,
@@ -791,12 +826,12 @@ extern "C" int coli_cuda_init(const int *devices, int count) {
     for (int i = 0; i < count; i++) {
         int device = devices[i];
         if (device < 0 || device >= available) {
-            std::fprintf(stderr, "[CUDA] invalid device %d (available: 0..%d)\n", device, available - 1);
+            std::fprintf(stderr, COLI_ACCEL_TAG " invalid device %d (available: 0..%d)\n", device, available - 1);
             g_nctx = 0;
             return 0;
         }
         if (find_ctx(device)) {
-            std::fprintf(stderr, "[CUDA] duplicate device %d\n", device);
+            std::fprintf(stderr, COLI_ACCEL_TAG " duplicate device %d\n", device);
             g_nctx = 0;
             return 0;
         }
@@ -817,7 +852,7 @@ extern "C" int coli_cuda_init(const int *devices, int count) {
         }
 #endif
         g_nctx++;
-        std::fprintf(stderr, "[CUDA] device %d: %s, %.1f GB VRAM, sm_%d%d\n",
+        std::fprintf(stderr, COLI_ACCEL_TAG " device %d: %s, %.1f GB VRAM, sm_%d%d\n",
                      device, prop.name, prop.totalGlobalMem / 1e9, prop.major, prop.minor);
     }
     return 1;
@@ -1681,6 +1716,29 @@ extern "C" int coli_cuda_attention_project_ragged(ColiCudaTensor *w,ColiCudaTens
            cuda_ok(cudaStreamSynchronize(dc->stream),"ragged attention synchronize");
 }
 
+extern "C" int coli_cuda_gqa_attention(float *ctx,const float *q,const float *k_cache,const float *v_cache,
+                                       int S,int H,int Hkv,int hd,int st0,int pos_base,int max_t,int device){
+    if(!ctx||!q||!k_cache||!v_cache||S<1||H<1||Hkv<1||hd<1||max_t<1||H%Hkv)return 0;
+    int nrep=H/Hkv,max_pos=pos_base+S-1,nt=max_pos+1-st0;
+    if(nt<1||nt>GQA_MAX_NT)return 0;
+    DeviceContext *dc=find_ctx(device); if(!select_ctx(dc)) return 0;
+    float scale=1.f/sqrtf((float)hd);
+    size_t qb=(size_t)S*H*hd*sizeof(float),kb=(size_t)Hkv*max_t*hd*sizeof(float);
+    size_t cb=qb,sb=(size_t)S*H*GQA_MAX_NT*sizeof(float);
+    if(!reserve(&dc->gq,&dc->gq_cap,qb)||!reserve(&dc->gk,&dc->gk_cap,kb)||
+       !reserve(&dc->gv,&dc->gv_cap,kb)||!reserve(&dc->gctx,&dc->gctx_cap,cb)||
+       !reserve(&dc->gsc,&dc->gsc_cap,sb)) return 0;
+    if(!cuda_ok(cudaMemcpyAsync(dc->gq,q,qb,cudaMemcpyHostToDevice,dc->stream),"gqa q upload")||
+       !cuda_ok(cudaMemcpyAsync(dc->gk,k_cache,kb,cudaMemcpyHostToDevice,dc->stream),"gqa k upload")||
+       !cuda_ok(cudaMemcpyAsync(dc->gv,v_cache,kb,cudaMemcpyHostToDevice,dc->stream),"gqa v upload")) return 0;
+    dim3 grid(H,S);
+    gqa_attn_kernel<<<grid,256,0,dc->stream>>>(dc->gctx,dc->gq,dc->gk,dc->gv,dc->gsc,S,H,Hkv,hd,st0,pos_base,max_t,scale,nrep);
+    if(!cuda_ok(cudaGetLastError(),"gqa attention launch")||
+       !cuda_ok(cudaMemcpyAsync(ctx,dc->gctx,cb,cudaMemcpyDeviceToHost,dc->stream),"gqa ctx download")||
+       !cuda_ok(cudaStreamSynchronize(dc->stream),"gqa attention synchronize")) return 0;
+    return 1;
+}
+
 extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) {
     if (!tensor) return;
     DeviceContext *ctx = find_ctx(tensor->device);

c/backend_cuda.h

@@ -21,6 +21,12 @@ extern "C" {
 
 #define COLI_CUDA_MAX_DEVICES 16
 
+#if defined(COLI_ROCM)
+#define COLI_ACCEL_TAG "[ROCm]"
+#else
+#define COLI_ACCEL_TAG "[CUDA]"
+#endif
+
 /* Opaque, persistent device copy of one resident quantized tensor. */
 typedef struct ColiCudaTensor ColiCudaTensor;
 
@@ -106,6 +112,12 @@ COLI_CUDA_DLLEXPORT int coli_cuda_attention_absorb(ColiCudaTensor *kv_b,float *c
                                const float *latent,const float *rope,int H,int Q,
                                int R,int V,int K,int T,float attention_scale);
 
+/* GQA decode/prefill attention: ctx[S,H,hd] from q[S,H,hd] and float K/V caches. */
+COLI_CUDA_DLLEXPORT int coli_cuda_gqa_attention(float *ctx, const float *q,
+                            const float *k_cache, const float *v_cache,
+                            int S, int H, int Hkv, int hd, int st0, int pos_base,
+                            int max_t, int device);
+
 /* Causal MLA absorption for S contiguous rows from one sequence.  The KV
  * arrays contain T rows ending at the final query; query s attends T-S+s+1
  * rows.  One transfer and one launch replace S host round-trips. */