[Feature]: Provide the necessary Windows coli_cuda.dll files or update instructions to build them

[SoftologyPro]

Problem

Building the CUDA dll on Windows.

Following #837 (comment)

1. Install https://github.com/msys2/msys2-installer/releases/download/2026-06-11/msys2-x86_64-20260611.exe (also the link to msys has a trailing \ here https://github.com/JustVugg/colibri/blob/main/docs/windows.md so it 404s)
2. When install finishes and the CLI for MSYS64 appears, run this inside it pacman -S make mingw-w64-ucrt-x86_64-gcc
3. Windows key, "developer", start the "Developer Command Prompt for VS 2022"
4. From the dev prompt run
   set MSYS2_PATH_TYPE=inherit
   C:\msys64\msys2_shell.cmd -ucrt64
5. From the second MSYS CLI that appears
   cd /d/Tests/Colibri/Colibri
   make cuda-dll CUDA_ARCH=native

It gets to the last make command and reports
make: *** No rule to make target 'cuda-dll'. Stop.

Proposed solution

Devs build them as required once for all Windows users and include them with the release zip. Please do this and save all of the Windows users having to go through the same mess I am :)

Or at least give full working steps to build the DLLs for the various architectures (at a minimum 30xx 40xx and 50xx GPUs)

Please help all of us poor Windows users.

Alternatives considered

No response

Scope and compatibility

No response

[JustVugg]

Thank you, and the "no rule to make target" has a plain cause: the release zip ships the engines and the Python only, no Makefile and no backend_cuda.cu, so make cuda-dll in the unzipped folder has nothing to build from. The steps from #837 are right, with one line they did not say: they need a git clone, and every make runs inside its c/ directory:

git clone https://github.com/JustVugg/colibri
cd colibri/c                (in the MSYS2 shell started from the VS developer prompt: cd /d/Tests/colibri/c)
make cuda-dll CUDA_ARCH=native
make colibri.exe CUDA_DLL=1 ARCH=native

docs/windows.md now says so, and the two things you tripped on there (the msys2 link with the stray backslash, CUDA-ARCH for CUDA_ARCH) are fixed in #1406.

On the request itself: agreed, and it was decided in August and not done. The pieces exist: CUDA_ARCH=portable builds one DLL with SASS for sm_80/86/89/90/120 plus PTX (30xx, 40xx, 50xx, and JIT for anything newer; it needs CUDA 12.8 or later), and engines built with CUDA_DLL=1 load coli_cuda.dll at runtime and stay CPU-only when it is absent. What has never existed is a card to prove the portable DLL runs, since CI has no GPU and builds it for sm_80 only as a compile check.

So, in two steps:

1. The CI job "CUDA build (Windows, MSVC host)" moves to CUDA 12.8 and CUDA_ARCH=portable, and publishes coli_cuda.dll as a workflow artifact. You mentioned 30xx, 40xx and 50xx: if you can run that artifact with a CUDA_DLL=1 colibri.exe on each and paste the [CUDA] banner and one decode, that is the validation nobody else in this tracker can give.
2. Once it holds, release.yml puts the DLL, the loader-enabled engines and the CUDA runtime DLL in the Windows zip, so the next release needs no toolchain at all.

I will link step 1 here when the artifact is up.

[SoftologyPro]

Error when running make cuda-dll CUDA_ARCH=native

$ make cuda-dll CUDA_ARCH=native
"nvcc" -O3 -std=c++17 -ftz=false -arch=native -Xcompiler=-W3 -shared -DCOLI_CUDA_BUILDING_DLL \
        -L"C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v12.8/lib/x64" -lcudart \
        backend_cuda.cu -o coli_cuda.dll
backend_cuda.cu
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1893): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.nc.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1899): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.nc.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1905): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cg.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1911): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cg.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1917): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.ca.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1923): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.ca.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1929): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cs.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1935): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cs.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1941): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.lu.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1947): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.lu.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1953): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cv.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1959): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cv.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1964): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wb.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1968): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wb.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1972): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cg.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1976): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cg.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1980): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cs.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1984): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cs.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1988): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wt.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(1992): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wt.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_fp16.hpp(3451): error: asm operand type size(8) does not match type/size implied by constraint 'r'
                    : "r"(address), "h"(*(reinterpret_cast<const unsigned short *>(&(val))))
                      ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1822): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.nc.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1828): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.nc.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1834): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cg.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1840): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cg.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1846): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.ca.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1852): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.ca.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1858): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cs.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr));
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1864): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cs.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr));
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1870): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.lu.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1876): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.lu.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1882): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cv.b32 %0, [%1];"  : "=r"(*(reinterpret_cast<unsigned int *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                               ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1888): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("ld.global.cv.b16 %0, [%1];"  : "=h"(*(reinterpret_cast<unsigned short *>(&(ret)))) : "r"(ptr) : "memory");
                                                                                                 ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1894): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wb.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1898): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wb.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1902): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cg.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1906): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cg.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1910): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cs.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1914): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.cs.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1918): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wt.b32 [%0], %1;"  :: "r"(ptr), "r"(*(reinterpret_cast<const unsigned int *>(&(value)))) : "memory");
                                            ^

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\include\cuda_bf16.hpp(1922): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("st.global.wt.b16 [%0], %1;"  :: "r"(ptr),  "h"(*(reinterpret_cast<const unsigned short *>(&(value)))) : "memory");
                                            ^

backend_cuda.cu(438): warning #767-D: conversion from pointer to smaller integer
      int vec=!(I&3)&&!((size_t)wrow&3)&&!((size_t)xs&15);
                        ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

backend_cuda.cu(438): warning #767-D: conversion from pointer to smaller integer
      int vec=!(I&3)&&!((size_t)wrow&3)&&!((size_t)xs&15);
                                           ^

backend_cuda.cu(946): warning #767-D: conversion from pointer to smaller integer
      int vec=!(D&3)&&!(((size_t)gr|(size_t)ur)&3)&&!((size_t)x&15);
                         ^

backend_cuda.cu(946): warning #767-D: conversion from pointer to smaller integer
      int vec=!(D&3)&&!(((size_t)gr|(size_t)ur)&3)&&!((size_t)x&15);
                                    ^

backend_cuda.cu(946): warning #767-D: conversion from pointer to smaller integer
      int vec=!(D&3)&&!(((size_t)gr|(size_t)ur)&3)&&!((size_t)x&15);
                                                      ^

backend_cuda.cu(978): warning #767-D: conversion from pointer to smaller integer
      int vec=!(I&3)&&!((size_t)row&3)&&!((size_t)x&15);
                        ^

backend_cuda.cu(978): warning #767-D: conversion from pointer to smaller integer
      int vec=!(I&3)&&!((size_t)row&3)&&!((size_t)x&15);
                                          ^

41 errors detected in the compilation of "backend_cuda.cu".
make: *** [Makefile:748: cuda-dll] Error 2

make colibri.exe CUDA_DLL=1 ARCH=native
seems to run file and compile the colibri.exe OK

[JustVugg]

That error is a known one, and it is my instructions that caused it: asm operand type size(8) does not match type/size implied by constraint 'r' in cuda_fp16.hpp is what nvcc prints when its host compiler is the 32-bit cl.exe. "Developer Command Prompt for VS 2022", the one I named in #837, opens the x86 environment. The DLL needs the 64-bit one:

• Start menu: "x64 Native Tools Command Prompt for VS 2022" (or run "C:\Program Files\Microsoft Visual Studio\2022\Community\VC\Auxiliary\Build\vcvars64.bat" in a cmd window; adjust Community to your edition),
• then the same two lines as before from that window: set MSYS2_PATH_TYPE=inherit and C:\msys64\msys2_shell.cmd -ucrt64,
• then in the MSYS2 shell: cd /d/Tests/colibri/c and make cuda-dll CUDA_ARCH=native.

You can tell which one you have with cl alone: the banner says for x64 or for x86. I am adding that check to the Makefile so the next person gets one line instead of forty errors, and the note to docs/windows.md.

Separately, #1407 (in CI now) makes the CI build the portable DLL on CUDA 12.8.1 for sm_80/86/89/90/120 and publishes it as an artifact with the loader-built colibri.exe, qwen36.exe, kimi_k3.exe and cudart64_12.dll; I will post the link here as soon as it is up, which would spare you the toolchain entirely for the test.

[JustVugg]

The portable DLL exists now: built by CI on CUDA 12.8.1 with CUDA_ARCH=portable (SASS for sm_80/86/89/90/120 plus PTX), from #1407.

Artifact (GitHub login needed to download): https://github.com/JustVugg/colibri/actions/runs/34422449152 -> coli_cuda-windows-portable. It holds coli_cuda.dll (exports the 55 coli_cuda_* symbols, depends only on cudart64_12.dll and KERNEL32), cudart64_12.dll, and colibri.exe, qwen36.exe, kimi_k3.exe built with the runtime loader. It expires in 30 days; the release will carry the same files once this is confirmed on a card.

To test without any toolchain: unzip the v1.10.2 Windows release, copy the five files from the artifact over it (the three exe replace the CPU-only ones; the two DLLs sit next to them), then in PowerShell from that folder:

$env:COLI_CUDA="1"; $env:COLI_GPU="0"; $env:CUDA_DENSE="1"; $env:CUDA_EXPERT_GB="8"
python coli doctor --model <your GLM-5.2 container>
python coli chat --model <your GLM-5.2 container>

coli doctor should stop saying "engine is CPU-only" and name the DLL it will load; the engine prints a [CUDA] banner at start (device, dense placement, expert tier). What settles it is that banner plus one decode with tok/s, on each of the 30xx / 40xx / 50xx boxes you have. A driver that is too old for the 12.8 runtime shows up as an init failure naming cudart, so please include the driver version too.

[SoftologyPro]

OK, that worked. I now have a coli_cuda.dll.
I copy that and the colibri.exe that were built into the unzipped release download.
When starting it shows GPU

The UI shows the GPU now detected

I enter a simple test prompt. Task manager shows full disk activity and physical RAM fills up.
It then starts answering the promptr, but no GPU activity and it is at around 5 seconds per word.

I will try your built dll next.

[JustVugg]

That worked better than the words after it suggest, and your banner says exactly where the time went. Two lines out of it:

[CUDA] device 0: NVIDIA GeForce RTX 5090, 34.2 GB VRAM, sm_120
[CUDA] hot expert tier: 53/53 experts, VRAM 1.13 GB (budget 30.7 GB, reserve 2.0 GB)

The first is the news for this issue: your DLL initialised on a 50-series card, so the x64 build recipe on CUDA 12.8 is correct end to end. The second is why the GPU then sat idle: the tier was given a 30.7 GB budget and filled 1.13 GB of it, 53 experts out of 19,456. The reason is the line above it, [PIN] auto: ... max_pin 1.1 GB -> pinning 1.1 GB [CAPPED by the LRU reserve]: in this mode the VRAM tier is fed from the RAM pin plan, and the planner held 101 GB of your RAM back as an LRU reserve and let the pin have 1.1 GB, so the card got the same 1.1 GB. With 53 experts on the GPU and an empty RAM cache, every token streams ~12 GB of experts from disk, which is the full-disk, 5 s/word you saw, and resident dense on CPU means the 12.5 GB dense trunk stayed on the CPU too.

That is a planner defect, not your build, and I am opening it as its own issue. The recipe that bypasses it today is the one the big-VRAM reports use:

$env:COLI_CUDA="1"; $env:CUDA_DENSE="1"; $env:CUDA_EXPERT_GB="28"; $env:PIN_GB="all"; $env:PIN="auto"

PIN_GB=all makes the VRAM prefix come from the VRAM budget instead of the RAM pin cap; CUDA_DENSE=1 moves the dense trunk to the card; 28 leaves the 2 GB reserve and the dense weights room on a 34 GB card. Expect the tier line to read on the order of 1,300 experts, and the first request to still be slow (cold page cache) with the second and later ones warming as the RAM cache fills; 120 GB of RAM holds most of the model.

If you can paste the [CUDA] lines and a decode tok/s from that run, it is the datapoint this issue needed, and the artifact DLL should behave the same as yours.

[SoftologyPro]

The portable DLL exists now: built by CI on CUDA 12.8.1 with CUDA_ARCH=portable (SASS for sm_80/86/89/90/120 plus PTX), from #1407.

Artifact (GitHub login needed to download): https://github.com/JustVugg/colibri/actions/runs/34422449152 -> coli_cuda-windows-portable. It holds coli_cuda.dll (exports the 55 coli_cuda_* symbols, depends only on cudart64_12.dll and KERNEL32), cudart64_12.dll, and colibri.exe, qwen36.exe, kimi_k3.exe built with the runtime loader. It expires in 30 days; the release will carry the same files once this is confirmed on a card.

To test without any toolchain: unzip the v1.10.2 Windows release, copy the five files from the artifact over it (the three exe replace the CPU-only ones; the two DLLs sit next to them), then in PowerShell from that folder:

I cannot see the exe or DLL in those links? Can you give direct links or tell me what link and path?

[JustVugg]

Direct link to the artifact zip (you need to be signed in to GitHub, artifacts are not public downloads):

https://github.com/JustVugg/colibri/actions/runs/34422449152/artifacts/10131473232

On the run page it is at the very bottom of the Summary, in the "Artifacts" box, named coli_cuda-windows-portable (2.5 MB). The zip is flat, five files:

coli_cuda.dll
cudart64_12.dll
colibri.exe
qwen36.exe
kimi_k3.exe

Copy all five into the folder of the unzipped v1.10.2 Windows release, next to the existing colibri.exe (overwrite the three exe), and run coli from there. If you have the GitHub CLI: gh run download 34422449152 -n coli_cuda-windows-portable -R JustVugg/colibri fetches the same files.

[SoftologyPro]

Direct link to the artifact zip (you need to be signed in to GitHub, artifacts are not public downloads):

https://github.com/JustVugg/colibri/actions/runs/34422449152/artifacts/10131473232

On the run page it is at the very bottom of the Summary, in the "Artifacts" box, named coli_cuda-windows-portable (2.5 MB). The zip is flat, five files:

coli_cuda.dll
cudart64_12.dll
colibri.exe
qwen36.exe
kimi_k3.exe

Copy all five into the folder of the unzipped v1.10.2 Windows release, next to the existing colibri.exe (overwrite the three exe), and run coli from there. If you have the GitHub CLI: gh run download 34422449152 -n coli_cuda-windows-portable -R JustVugg/colibri fetches the same files.

OK, Downloaded those and overwrite the existing files in the windows release.
Same problem. Starts showing GPU, GPU shows in UI, when the prompt runs it goes full disk activity as the 128 GB physical RAM is filed up.
The prompt response then starts at around 5 seconds a word. No GPU activity is shown in Task Manager.

coli doctor

[JustVugg]

Same result is expected: the DLL was never the problem, and your own coli doctor shows it. It plans VRAM 30.6 GB hot tier · ~1442 experts, while the engine banner from your earlier run placed 53. That gap is the bug I filed as #1409 (the tier is fed from the RAM pin plan, which the LRU reserve capped at 1.1 GB), and it is the same with your DLL or the artifact one. The 128 GB filling up is the RAM expert cache warming on the first request, which is normal; what is not normal is the card holding 53 experts.

The workaround has to be set in the same window before you start coli web (or coli chat). Your prompt is cmd, so:

set COLI_CUDA=1
set CUDA_DENSE=1
set CUDA_EXPERT_GB=28
set PIN=auto
set PIN_GB=all
coli chat --model D:\Tests\Colibri\Colibri\glm52_i4

(PowerShell form: $env:COLI_CUDA="1"; $env:CUDA_DENSE="1"; $env:CUDA_EXPERT_GB="28"; $env:PIN="auto"; $env:PIN_GB="all".)

What to look for at start: [CUDA] mode: should no longer say resident dense on CPU, and [CUDA] hot expert tier: should read on the order of 1,400 experts and ~28 GB, not 53 / 1.13 GB. The first prompt is still slow (cold disk), the second one is the number. If the tier line still says 53 with those variables set, paste the [PIN] and [CUDA] lines and I will take it from there; if it says ~1,400 and the GPU still shows no load, that is a different bug and also worth the paste.