danielkrause/colibri-strix
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Windows native port: serve-mode pipe fix + RAM detection + POSIX guards, AVX-VNNI kernel, gated CUDA DLL (#131, fixes #123)

Browse Source 
Rebased onto current dev, split into 3 logical parts (all validated):
1. CPU portability (serve-mode _O_BINARY pipe fix — stock main hangs on MinGW without it; RAM detection cap 0->9/layer; POSIX guards for select/mmap/madvise; warmup script).
2. AVX-VNNI 128-bit int8/int4 dot kernel (Alder Lake+/Meteor Lake+), bit-identical to AVX2 (author-verified on Meteor Lake; compiles out to AVX2 elsewhere) + _mm256_extracti128_si256 typo fix that blocked -march=native.
3. CUDA DLL via LoadLibrary, gated behind CUDA_DLL=1 (host never links cudart; silent CPU fallback if absent; author-verified on RTX 5070 Ti).

Validated here: make check 59/59, oracle 32/32 TF, Windows cross-compile clean + glm.exe loads+runs via WSL interop. Fixes the #123 Windows build failure.
This commit is contained in:
woolcoxm
2026-07-13 14:54:30 -04:00
committed by GitHub
parent afc259c599
commit 2319b942d2
10 changed files with 644 additions and 30 deletions
Download Patch File Download Diff File Expand all files Collapse all files
c/backend_cuda.h
+25 -15
View File
@@ -4,6 +4,16 @@
#include <stddef.h>
#include <stdint.h>
/* COLI_CUDA_DLLEXPORT marks functions exported from coli_cuda.dll on Windows.
* Define COLI_CUDA_BUILDING_DLL when compiling the .cu into the DLL (so the
* functions are __declspec(dllexport)); the host loader does NOT include this
* header's declarations — it resolves symbols at runtime via GetProcAddress. */
#if defined(_WIN32) && defined(COLI_CUDA_BUILDING_DLL)
#define COLI_CUDA_DLLEXPORT __declspec(dllexport)
#else
#define COLI_CUDA_DLLEXPORT
#endif
#ifdef __cplusplus
extern "C" {
#endif
@@ -14,18 +24,18 @@ extern "C" {
typedef struct ColiCudaTensor ColiCudaTensor;
/* Devices are CUDA ordinals, not positions in the input list. */
int coli_cuda_init(const int *devices, int count);
void coli_cuda_shutdown(void);
int coli_cuda_device_count(void);
int coli_cuda_device_at(int index);
int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes);
COLI_CUDA_DLLEXPORT int coli_cuda_init(const int *devices, int count);
COLI_CUDA_DLLEXPORT void coli_cuda_shutdown(void);
COLI_CUDA_DLLEXPORT int coli_cuda_device_count(void);
COLI_CUDA_DLLEXPORT int coli_cuda_device_at(int index);
COLI_CUDA_DLLEXPORT int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes);
/* device < 0 returns aggregate statistics for all configured devices. */
void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes);
void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows,
COLI_CUDA_DLLEXPORT void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes);
COLI_CUDA_DLLEXPORT void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows,
double *h2d_ms, double *kernel_ms, double *d2h_ms);
/* Upload without executing, so capacity failures happen during model startup. */
int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
const void *weights, const float *scales,
int fmt, int I, int O, int device);
@@ -35,7 +45,7 @@ int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
* The first successful call uploads W and its row scales; later calls reuse it.
* Returns 1 on success and 0 when CUDA is not initialized or the format is invalid.
*/
int coli_cuda_matmul(ColiCudaTensor **tensor,
COLI_CUDA_DLLEXPORT int coli_cuda_matmul(ColiCudaTensor **tensor,
float *y, const float *x,
const void *weights, const float *scales,
int fmt, int S, int I, int O, int device);
@@ -43,25 +53,25 @@ int coli_cuda_matmul(ColiCudaTensor **tensor,
/* Fused expert pipeline: y = down(silu(gate(x)) * up(x)). All three tensors
* must already be resident on one device. Activations cross PCIe once in
* each direction instead of once per matrix. */
int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
COLI_CUDA_DLLEXPORT int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
ColiCudaTensor *down, float *y, const float *x, int S);
/* Packed group of same-shaped experts. Inputs and outputs contain sum(rows)
* consecutive [D] rows in call order. */
int coli_cuda_expert_group(ColiCudaTensor *const *gates,
COLI_CUDA_DLLEXPORT int coli_cuda_expert_group(ColiCudaTensor *const *gates,
ColiCudaTensor *const *ups,
ColiCudaTensor *const *downs,
const int *rows, int count,
float *y, const float *x);
/* Decode-only MLA weight-absorption core for one token. kv_b is [H*(Q+V),K]. */
int coli_cuda_attention_absorb(ColiCudaTensor *kv_b,float *ctx,const float *q,
COLI_CUDA_DLLEXPORT int coli_cuda_attention_absorb(ColiCudaTensor *kv_b,float *ctx,const float *q,
const float *latent,const float *rope,int H,int Q,
int R,int V,int K,int T,float attention_scale);
void coli_cuda_tensor_free(ColiCudaTensor *tensor);
size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor);
int coli_cuda_tensor_device(const ColiCudaTensor *tensor);
COLI_CUDA_DLLEXPORT void coli_cuda_tensor_free(ColiCudaTensor *tensor);
COLI_CUDA_DLLEXPORT size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor);
COLI_CUDA_DLLEXPORT int coli_cuda_tensor_device(const ColiCudaTensor *tensor);
#ifdef __cplusplus
}