url/colibri
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #338 from noobdev-ph/feat/gpu-backend-hardening

Browse Source 
GPU backend: failure-path hardening + tests (sticky-error fix, cached-tensor contract, fault-injection hook)
This commit is contained in:
Vincenzo Fornaro
2026-07-20 18:12:00 +02:00
committed by GitHub
parent 93cde63cd3 26d6b2d662
commit e48513c1c6
2 changed files with 87 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
c/backend_cuda.cu
+41 -4
View File
@@ -58,6 +58,8 @@ static std::mutex g_group_stats_mu;
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));
(void)cudaGetLastError(); /* consume the sticky error: a failed call must
not poison the next launch's error check */
return 0;
}
@@ -554,14 +556,20 @@ extern "C" int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor,
extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
const void *weights, const float *scales,
int fmt, int I, int O, int device) {
DeviceContext *ctx = find_ctx(device);
if (!tensor || !weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0;
size_t rb = row_bytes(fmt, I);
if (!rb || (fmt && !scales)) return 0;
if (!tensor) return 0;
if (*tensor) {
/* Cached device copy: usable even when the caller's host pointers are
* gone. CUDA_RELEASE_HOST slots null their host pointers after upload,
* and with the old order (!weights checked first) every later matmul
* on such a slot failed here — the GPU tier silently never computed
* for host-released slab experts. */
ColiCudaTensor *t = *tensor;
return t->fmt == fmt && t->I == I && t->O == O && t->device == device;
}
DeviceContext *ctx = find_ctx(device);
if (!weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0;
size_t rb = row_bytes(fmt, I);
if (!rb || (fmt && !scales)) return 0;
ColiCudaTensor *t = static_cast<ColiCudaTensor *>(std::calloc(1, sizeof(*t)));
if (!t) return 0;
t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O;
@@ -615,10 +623,21 @@ extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor,
cudaMemcpyHostToDevice),"scale refresh");
}
/* Test hook: COLI_GPU_FAIL_AFTER=N makes every GPU COMPUTE entry point report
* failure after N successful calls (N=0: every call fails), exercising the
* engine's CPU fallbacks and host-rematerialization end-to-end without real
* hardware faults. Uploads/queries are not gated. Unset: no effect. */
static long g_gpu_calls;
static int fault_injected(void) {
const char *fa = std::getenv("COLI_GPU_FAIL_AFTER");
return fa && g_gpu_calls++ >= std::atol(fa);
}
extern "C" 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) {
if (fault_injected()) return 0;
if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0;
ColiCudaTensor *t = *tensor;
DeviceContext *ctx = find_ctx(t->device);
@@ -637,6 +656,7 @@ extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor,
extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
ColiCudaTensor *down, float *y,
const float *x, int S) {
if (fault_injected()) return 0;
if (!gate || !up || !down || !x || !y || S < 1 ||
gate->device != up->device || gate->device != down->device ||
gate->I != up->I || gate->O != up->O ||
@@ -665,6 +685,7 @@ extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
extern "C" int coli_cuda_shared_mlp_w4a16(ColiCudaTensor *gate,ColiCudaTensor *up,
ColiCudaTensor *down,float *y,const float *x,int S){
if (fault_injected()) return 0;
if(!gate||!up||!down||!x||!y||S<1||gate->fmt!=2||up->fmt!=2||down->fmt!=2||
gate->device!=up->device||gate->device!=down->device||gate->I!=up->I||
gate->O!=up->O||down->I!=gate->O||down->O!=gate->I)return 0;
@@ -696,6 +717,7 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
ColiCudaTensor *const *downs,
const int *rows, int count,
float *y, const float *x) {
if (fault_injected()) return 0;
if (!gates || !ups || !downs || !rows || !x || !y || count < 1) return 0;
ColiCudaTensor *first=gates[0];
if (!first) return 0;
@@ -912,6 +934,7 @@ extern "C" const float *coli_cuda_expert_group_take(int device) {
extern "C" int coli_cuda_attention_absorb(ColiCudaTensor *w,float *ctx,const float *q,
const float *latent,const float *rope,int H,int Q,
int R,int V,int K,int T,float scale){
if (fault_injected()) return 0;
if(!w||!ctx||!q||!latent||!rope||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>4096||
w->I!=K||w->O!=H*(Q+V))return 0;
DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0;
@@ -965,12 +988,14 @@ static int attention_absorb_batch_run(ColiCudaTensor *w,ColiCudaTensor *proj,flo
extern "C" int coli_cuda_attention_absorb_batch(ColiCudaTensor *w,float *ctx,const float *q,
const float *latent,const float *rope,int S,int H,int Q,int R,int V,int K,int T,
float scale){
if (fault_injected()) return 0;
return attention_absorb_batch_run(w,nullptr,ctx,q,latent,rope,S,H,Q,R,V,K,T,scale);
}
extern "C" int coli_cuda_attention_project_batch(ColiCudaTensor *w,ColiCudaTensor *proj,
float *out,const float *q,const float *latent,const float *rope,int S,int H,int Q,
int R,int V,int K,int T,float scale){
if (fault_injected()) return 0;
return attention_absorb_batch_run(w,proj,out,q,latent,rope,S,H,Q,R,V,K,T,scale);
}
@@ -1161,6 +1186,7 @@ extern "C" int coli_cuda_pipe_download(int device,const void *src,void *dst,size
}
extern "C" int coli_cuda_pipe_rmsnorm(int device,float *y_dev,const float *x_dev,
const float *w_dev,int S,int D,float eps){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device);
if(S<1||D<1||!select_ctx(ctx)) return 0;
pipe_rmsnorm_rows<<<S,256>>>(y_dev,x_dev,w_dev,D,eps,D,D);
@@ -1169,6 +1195,7 @@ extern "C" int coli_cuda_pipe_rmsnorm(int device,float *y_dev,const float *x_dev
extern "C" int coli_cuda_pipe_rmsnorm_s(int device,float *y_dev,const float *x_dev,
const float *w_dev,int S,int D,float eps,
int xstride,int ystride){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device);
if(S<1||D<1||xstride<D||ystride<D||!select_ctx(ctx)) return 0;
pipe_rmsnorm_rows<<<S,256>>>(y_dev,x_dev,w_dev,D,eps,xstride,ystride);
@@ -1177,6 +1204,7 @@ extern "C" int coli_cuda_pipe_rmsnorm_s(int device,float *y_dev,const float *x_d
extern "C" int coli_cuda_pipe_rope(int device,float *v_dev,const int *pos_dev,
int rows,int stride,int offset,int R,int heads,
float theta){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device);
if(rows<1||R<2||R>256||heads<1||!select_ctx(ctx)) return 0;
pipe_rope_rows<<<rows,128>>>(v_dev,pos_dev,0,stride,offset,R,heads,theta);
@@ -1184,6 +1212,7 @@ extern "C" int coli_cuda_pipe_rope(int device,float *v_dev,const int *pos_dev,
}
extern "C" int coli_cuda_pipe_rope_base(int device,float *v_dev,int pos_base,int rows,
int stride,int offset,int R,int heads,float theta){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device);
if(rows<1||R<2||R>256||heads<1||!select_ctx(ctx)) return 0;
pipe_rope_rows<<<rows,128>>>(v_dev,NULL,pos_base,stride,offset,R,heads,theta);
@@ -1277,6 +1306,7 @@ extern "C" int coli_cuda_pipe_copy2d(int device,float *dst,int dpitch,const floa
extern "C" int coli_cuda_attention_project_batch_dev(ColiCudaTensor *w,ColiCudaTensor *proj,
float *out,const float *q_dev,const float *latent_dev,const float *rope_dev,
int S,int H,int Q,int R,int V,int K,int T,float scale){
if (fault_injected()) return 0;
if(!w||!proj||!out||!q_dev||!latent_dev||!rope_dev||S<1||H<1||Q<1||R<1||V<1||
K<1||K>512||T<S||T>8192||w->I!=K||w->O!=H*(Q+V)||
proj->device!=w->device||proj->I!=H*V)return 0;
@@ -1298,17 +1328,20 @@ extern "C" int coli_cuda_attention_project_batch_dev(ColiCudaTensor *w,ColiCudaT
}
extern "C" int coli_cuda_pipe_silu_mul(int device,float *gate_dev,const float *up_dev,
size_t n){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device); if(!n||!select_ctx(ctx)) return 0;
silu_mul<<<(unsigned)((n+255)/256),256>>>(gate_dev,up_dev,n);
return cuda_ok(cudaGetLastError(),"pipe silu mul");
}
extern "C" int coli_cuda_pipe_add(int device,float *x_dev,const float *t_dev,size_t n){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device); if(!n||!select_ctx(ctx)) return 0;
pipe_add_n<<<(unsigned)((n+255)/256),256>>>(x_dev,t_dev,n);
return cuda_ok(cudaGetLastError(),"pipe add");
}
extern "C" int coli_cuda_pipe_rows_add(int device,float *x_dev,const float *partial_dev,
const int *rows_dev,int nrows,int D){
if (fault_injected()) return 0;
DeviceContext *ctx=find_ctx(device); if(nrows<1||D<1||!select_ctx(ctx)) return 0;
pipe_rows_add<<<nrows,256>>>(x_dev,partial_dev,rows_dev,D);
return cuda_ok(cudaGetLastError(),"pipe rows add");
@@ -1317,6 +1350,7 @@ extern "C" int coli_cuda_pipe_rows_add(int device,float *x_dev,const float *part
* coli_cuda_matmul, zero host transfers. */
extern "C" int coli_cuda_pipe_gemm(ColiCudaTensor *t,float *y_dev,const float *x_dev,
int S){
if (fault_injected()) return 0;
if(!t||S<1) return 0;
DeviceContext *ctx=find_ctx(t->device); if(!select_ctx(ctx)) return 0;
dim3 grid((unsigned)t->O,(unsigned)S);
@@ -1336,6 +1370,7 @@ extern "C" int coli_cuda_pipe_peer_copy(int dst_dev,float *dst,int src_dev,
extern "C" int coli_cuda_attention_project_batch_dev_out(ColiCudaTensor *w,ColiCudaTensor *proj,
float *out_dev,const float *q_dev,const float *latent_dev,const float *rope_dev,
int S,int H,int Q,int R,int V,int K,int T,float scale){
if (fault_injected()) return 0;
if(!w||!proj||!out_dev||!q_dev||!latent_dev||!rope_dev||S<1||H<1||Q<1||R<1||V<1||
K<1||K>512||T<S||T>8192||w->I!=K||w->O!=H*(Q+V)||
proj->device!=w->device||proj->I!=H*V)return 0;
@@ -1357,6 +1392,7 @@ extern "C" int coli_cuda_attention_project_batch_dev_out(ColiCudaTensor *w,ColiC
extern "C" int coli_cuda_attention_absorb_batch_dev(ColiCudaTensor *w,float *ctx_dev,
const float *q_dev,const float *latent_dev,const float *rope_dev,
int S,int H,int Q,int R,int V,int K,int T,float scale){
if (fault_injected()) return 0;
if(!w||!ctx_dev||!q_dev||!latent_dev||!rope_dev||S<1||H<1||Q<1||R<1||V<1||
K<1||K>512||T<S||T>8192||w->I!=K||w->O!=H*(Q+V))return 0;
DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0;
@@ -1371,6 +1407,7 @@ extern "C" int coli_cuda_attention_absorb_batch_dev(ColiCudaTensor *w,float *ctx
extern "C" int coli_cuda_attention_absorb_kvdev(ColiCudaTensor *w,float *ctx,const float *q,
const float *latent_dev,const float *rope_dev,int H,int Q,int R,int V,int K,int T,
float scale){
if (fault_injected()) return 0;
if(!w||!ctx||!q||!latent_dev||!rope_dev||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>8192||
w->I!=K||w->O!=H*(Q+V))return 0;
DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0;
c/tests/test_backend_cuda.cu
+46
View File
@@ -50,6 +50,52 @@ int main(int argc, char **argv) {
if (coli_cuda_tensor_upload(&t8, q8, s8, 1, 5, 2, d0)) return 1;
if (ndev > 1 && coli_cuda_tensor_upload(&t8, q8, s8, 1, 4, 2, d1)) return 1;
if (!coli_cuda_matmul(&t8, got, x, q8, s8, 1, 2, 4, 2, d0) || !close_enough(got, want8, 4)) return 1;
/* Cached tensor must stay callable without live host pointers
* (CUDA_RELEASE_HOST slots null theirs after upload) — including
* SUSTAINED reuse, not just the first call. */
for (int rep = 0; rep < 64; rep++)
if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
!close_enough(got, want8, 4)) return 1;
/* A tensor uploaded from a TEMPORARY host buffer must survive the buffer
* being scribbled and freed (the release-host lifecycle). */
{
int8_t *tmpw = static_cast<int8_t *>(std::malloc(8));
float *tmps = static_cast<float *>(std::malloc(2 * sizeof(float)));
if (!tmpw || !tmps) return 2;
for (int i = 0; i < 8; i++) tmpw[i] = q8[i];
tmps[0] = s8[0]; tmps[1] = s8[1];
ColiCudaTensor *tt = nullptr;
if (!coli_cuda_tensor_upload(&tt, tmpw, tmps, 1, 4, 2, d0)) return 1;
for (int i = 0; i < 8; i++) tmpw[i] = 99;
std::free(tmpw); std::free(tmps);
if (!coli_cuda_matmul(&tt, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
!close_enough(got, want8, 4)) return 1;
coli_cuda_tensor_free(tt);
}
/* Upload failures must be graceful and must not corrupt accounting —
* and must not poison LATER healthy launches (sticky-error regression). */
{
size_t c0 = 0, b0 = 0, c1 = 0, b1 = 0;
coli_cuda_stats(-1, &c0, &b0);
ColiCudaTensor *bad = nullptr;
if (coli_cuda_tensor_upload(&bad, q8, s8, 1, 4, 2, 9999)) return 1;
if (coli_cuda_tensor_upload(&bad, q8, s8, 7, 4, 2, d0)) return 1;
if (coli_cuda_tensor_upload(&bad, q8, nullptr, 1, 4, 2, d0)) return 1;
if (coli_cuda_tensor_upload(&bad, nullptr, s8, 1, 4, 2, d0)) return 1;
if (coli_cuda_tensor_upload(&bad, q8, s8, 1, 1 << 20, 1 << 24, d0)) return 1; /* ~16 TB */
if (bad) return 1;
coli_cuda_stats(-1, &c1, &b1);
if (c0 != c1 || b0 != b1) return 1;
/* healthy launch immediately after the failed allocation */
if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
!close_enough(got, want8, 4)) return 1;
}
/* Fault injection hook: on/off, restores cleanly. */
if (setenv("COLI_GPU_FAIL_AFTER", "0", 1)) return 2;
if (coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0)) return 1;
if (unsetenv("COLI_GPU_FAIL_AFTER")) return 2;
if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
!close_enough(got, want8, 4)) return 1;
const int8_t q8b[8]={-1,-2,-3,-4, 1,-2,3,-4};
const float s8b[2]={1.f,.5f},want8b[4]={10.f,15.f,-3.f,-2.5f};
if(!coli_cuda_tensor_update(t8,q8b,s8b)||