fix(rebase): adapt to dev's split tensor_upload/_g API + fault-injection hook

Dev split coli_cuda_tensor_upload into two symbols: the plain 7-arg
tensor_upload (no gs) for fmt!=4, and tensor_upload_g (8-arg, with gs)
for fmt=4 grouped. The DLL-side _g sets a g_upload_gs global the plain
upload reads internally. Adapted all call sites:

- colibri.c: qt_cuda_upload already dispatches (_g for fmt=4, plain
  otherwise) — kept dev's version. layer_cuda_shard_kvb was calling the
  plain upload with gs; switched to tensor_upload_g.
- backend_loader.c: the C-API wrapper matches dev's 7-arg signature;
  tensor_upload_g dispatched separately.
- backend_cuda.cu: cache-hit check uses g_upload_gs (not a gs param);
  coli_cuda_matmul dispatches _g when gs>0; the grouped-expert host
  fallback quant_matmul calls pass gs=0,ng=1 (host tier is per-row).
- Kept dev's fault-injection hook, scale_count field, and fmt=4 support
  in the grouped-expert path (all_q4/any_g4 tracking).
- tests: kept dev's comprehensive tensor_upload test (cached reuse,
  temp-buffer survival, upload-failure accounting, fault injection).

Build-verified: colibri.exe + coli_cuda.dll both compile clean.

c/backend_cuda.cu

@@ -592,7 +592,8 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
          * 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 && t->gs == gs;
+        int want_gs = (fmt==4 && g_upload_gs>0) ? g_upload_gs : 0;
+        return t->fmt == fmt && t->I == I && t->O == O && t->device == device && t->gs == want_gs;
     }
     DeviceContext *ctx = find_ctx(device);
     if (!weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0;
@@ -649,7 +650,6 @@ extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor,
     }
     int ng = tensor->ng > 0 ? tensor->ng : 1;
     return !tensor->fmt || cuda_ok(cudaMemcpy(tensor->scales,scales,
-<<<<<<< HEAD
         (tensor->scale_count?tensor->scale_count:(size_t)tensor->O)*sizeof(float),
         cudaMemcpyHostToDevice),"scale refresh");
 }
@@ -662,22 +662,16 @@ 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);
-=======
-        (size_t)tensor->O*ng*sizeof(float),cudaMemcpyHostToDevice),"scale refresh");
->>>>>>> 1a69113 (cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness)
 }
 
 extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor,
                                  float *y, const float *x,
                                  const void *weights, const float *scales,
-<<<<<<< HEAD
-                                 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;
-=======
                                  int fmt, int S, int I, int O, int device, int gs) {
-    if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device, gs)) return 0;
->>>>>>> 1a69113 (cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness)
+    if (fault_injected()) return 0;
+    if (S < 1) return 0;
+    if (gs > 0) { if (!coli_cuda_tensor_upload_g(tensor, weights, scales, fmt, I, O, device, gs)) return 0; }
+    else        { if (!coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0; }
     ColiCudaTensor *t = *tensor;
     DeviceContext *ctx = find_ctx(t->device);
     if (!select_ctx(ctx)) return 0;
@@ -771,16 +765,9 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
                  g->fmt,u->fmt,d->fmt,rows[c],total,
                  g->gs,u->gs,d->gs};
         all_s4&=g->fmt==2&&u->fmt==2&&d->fmt==2;
-<<<<<<< HEAD
         all_q4&=(g->fmt==2||g->fmt==4)&&(u->fmt==2||u->fmt==4)&&(d->fmt==2||d->fmt==4)&&
                 !(g->gs&1)&&!(u->gs&1)&&!(d->gs&1);   /* even gs: a packed byte never straddles groups */
         any_g4|=g->fmt==4||u->fmt==4||d->fmt==4;
-=======
-        /* fmt=4 (grouped int4) experts have per-group scales that the grouped_hidden/
-         * grouped_down kernels don't handle. Fall back to the per-expert path
-         * (coli_cuda_expert_mlp), which uses quant_matmul with correct gs/ng. */
-        if(g->fmt==4||u->fmt==4||d->fmt==4) return 0;
->>>>>>> 1a69113 (cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness)
         total+=rows[c]; if(rows[c]>max_rows) max_rows=rows[c];
     }
     DeviceContext *ctx=find_ctx(device); if(!select_ctx(ctx)) return 0;
@@ -951,12 +938,12 @@ extern "C" int coli_cuda_expert_group_issue(ColiCudaTensor *const *gates,
         float *g16=ctx->gate+(size_t)host[c].offset*I,*u16=ctx->up+(size_t)host[c].offset*I;
         float *x16=ctx->x+(size_t)host[c].offset*D,*y16=ctx->y+(size_t)host[c].offset*D;
         quant_matmul<<<dim3((unsigned)I,(unsigned)r),256,0,ctx->stream>>>(g16,x16,
-            host[c].g,host[c].gs,host[c].gf,r,D,I,row_bytes(host[c].gf,D));
+            host[c].g,host[c].gs,host[c].gf,r,D,I,row_bytes(host[c].gf,D),0,1);
         quant_matmul<<<dim3((unsigned)I,(unsigned)r),256,0,ctx->stream>>>(u16,x16,
-            host[c].u,host[c].us,host[c].uf,r,D,I,row_bytes(host[c].uf,D));
+            host[c].u,host[c].us,host[c].uf,r,D,I,row_bytes(host[c].uf,D),0,1);
         silu_mul<<<(unsigned)(((size_t)r*I+255)/256),256,0,ctx->stream>>>(g16,u16,(size_t)r*I);
         quant_matmul<<<dim3((unsigned)D,(unsigned)r),256,0,ctx->stream>>>(y16,g16,
-            host[c].d,host[c].ds,host[c].df,r,I,D,row_bytes(host[c].df,I));
+            host[c].d,host[c].ds,host[c].df,r,I,D,row_bytes(host[c].df,I),0,1);
     }
     if(!cuda_ok(cudaGetLastError(),"expert group issue launch")||
        !cuda_ok(cudaMemcpyAsync(ctx->host_y,ctx->y,xb,cudaMemcpyDeviceToHost,ctx->stream),

c/backend_cuda.h

@@ -41,7 +41,7 @@ COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor,
         int fmt, int I, int O, int device, int gs);
 COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
                             const void *weights, const float *scales,
-                            int fmt, int I, int O, int device, int gs);
+                            int fmt, int I, int O, int device);
 
 /*
  * y[S,O] = x[S,I] @ W[O,I]^T.

c/backend_loader.c

@@ -325,9 +325,9 @@ int coli_cuda_attention_absorb(ColiCudaTensor *kv_b, float *ctx, const float *q,
 }
 
 int coli_cuda_tensor_upload(ColiCudaTensor **tensor, const void *weights,
-                            const float *scales, int fmt, int I, int O, int device, int gs){
+                            const float *scales, int fmt, int I, int O, int device){
     if(!g_cuda.available) return 0;
-    return g_cuda.tensor_upload(tensor, weights, scales, fmt, I, O, device, gs);
+    return g_cuda.tensor_upload(tensor, weights, scales, fmt, I, O, device);
 }
 
 int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor, const void *weights, const float *scales, int fmt, int I, int O, int device, int gs){

c/colibri.c

@@ -1013,7 +1013,7 @@ static void layer_cuda_shard_kvb(Layer *l,int H,int Q,int V){
         int hn=H/g_cuda_ndev+(d<H%g_cuda_ndev),rows=hn*(Q+V);
         const void *part=weights+(int64_t)h0*(Q+V)*rb;
         const float *scale=l->kv_b.s+(int64_t)h0*(Q+V)*(l->kv_b.gs>0?(l->kv_b.I+l->kv_b.gs-1)/l->kv_b.gs:1);
-        if(!coli_cuda_tensor_upload(&l->kv_b_shard[d],part,scale,l->kv_b.fmt,l->kv_b.I,rows,g_cuda_devices[d],l->kv_b.gs))return;
+        if(!coli_cuda_tensor_upload_g(&l->kv_b_shard[d],part,scale,l->kv_b.fmt,l->kv_b.I,rows,g_cuda_devices[d],l->kv_b.gs))return;
         l->shard_h0[d]=h0;l->shard_hn[d]=hn;l->n_kv_b_shard++;h0+=hn;
     }
     int old=-1;for(int i=0;i<g_cuda_ndev;i++)if(g_cuda_devices[i]==l->kv_b.cuda_device)old=i;