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

fix: GPU backend failure-path hardening + tests

Browse Source 
Three vendor-neutral fixes to backend_cuda.cu, each with test coverage:

1. Upload check-order: a cached device tensor is now usable when the
   caller's host pointers are stale or NULL. CUDA_RELEASE_HOST slots
   null their host pointers after upload; the current engine reaches
   those tensors through direct handles (coli_cuda_expert_mlp etc.), but
   any caller going through coli_cuda_matmul/tensor_upload with a cached
   tensor — as matmul_qt does for QT tensors — hits the !weights check
   before the cached-tensor branch and fails spuriously. This hardens
   the API contract rather than fixing a measured regression; the
   contract is pinned by a 64x sustained-reuse test.

2. Sticky runtime error (real bug, test-caught): a failed allocation
   left the last-error state set, so the NEXT healthy launch's
   cudaGetLastError() check reported 'out of memory' and disabled a
   perfectly good tensor. cuda_ok() now consumes the error on the
   failure path; regression-covered.

3. COLI_GPU_FAIL_AFTER=N test hook: every GPU compute entry point (19
   total: matmul, expert mlp/group, shared mlp, attention ops, pipe ops)
   reports failure after N successful calls, so the engine's CPU
   fallbacks and expert_host_ensure rematerialization can be exercised
   end-to-end without real hardware faults. Unset = zero effect;
   uploads/queries are never gated. Validated on GLM-5.2: total failure
   (N=0) completes coherently with every released expert rematerialized.

Tests (run via make cuda-test on any CUDA GPU; vendor-neutral source):
64x sustained matmul reuse after host pointers are freed; upload from a
scribbled-and-freed temporary; five graceful upload-failure cases with
stats-integrity assertions; healthy-launch-after-failed-alloc (the
sticky-error regression); fault-hook on/off restore.

Verified on AMD RX 9070 XT via the companion HIP PR's compat header
(same test source); a make cuda-test run on NVIDIA hardware would
complete the matrix.
This commit is contained in:
noobdev-ph
2026-07-17 12:39:59 +08:00
parent e7188df16a
commit 288edd7190
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
@@ -45,6 +45,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;
}
@@ -453,14 +455,20 @@ extern "C" void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64
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;
@@ -502,10 +510,21 @@ extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor,
(size_t)tensor->O*sizeof(float),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);
@@ -524,6 +543,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 ||
@@ -552,6 +572,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;
@@ -583,6 +604,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;
@@ -708,6 +730,7 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
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;
@@ -761,12 +784,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);
}
@@ -870,6 +895,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);
@@ -878,6 +904,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);
@@ -886,6 +913,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);
@@ -893,6 +921,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);
@@ -910,6 +939,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;
@@ -931,17 +961,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");
@@ -950,6 +983,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);
@@ -969,6 +1003,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;
@@ -990,6 +1025,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;
@@ -1004,6 +1040,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)||