cuda: out_prod accepts a quantized src0 — the backward through a resident quantized base (Continuum dd109d3c)

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -5017,7 +5017,12 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                 }
             } break;
         case GGML_OP_OUT_PROD:
-            return op->type == GGML_TYPE_F32 && op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32;
+            // src0 may be F32, or any type with an F32 converter (F16/BF16/quantized) when
+            // contiguous: the backward through a frozen quantized weight (out_prod(W, grad^T))
+            // reads W as stored and dequantizes into a transient pool copy (see out-prod.cu).
+            return op->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 &&
+                (op->src[0]->type == GGML_TYPE_F32 ||
+                 (ggml_is_contiguous(op->src[0]) && ggml_get_to_fp32_cuda(op->src[0]->type) != nullptr));
         case GGML_OP_GET_ROWS:
             {
                 switch (op->src[0]->type) {

ggml/src/ggml-cuda/out-prod.cu

@@ -1,4 +1,5 @@
 #include "out-prod.cuh"
+#include "convert.cuh"
 
 #include <cstdint>
 
@@ -30,10 +31,41 @@ void ggml_cuda_out_prod(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
 
     GGML_TENSOR_BINARY_OP_LOCALS
 
-    GGML_ASSERT(src0->type == GGML_TYPE_F32);
     GGML_ASSERT(src1->type == GGML_TYPE_F32);
     GGML_ASSERT(dst->type  == GGML_TYPE_F32);
 
+    // A quantized (or F16/BF16) src0: the backward of mul_mat(W, x) w.r.t. x is
+    // out_prod(W, grad^T), so training a LoRA on a resident quantized base needs this
+    // op to read W as it is stored. The CPU backend dispatches every quantized type here
+    // (ggml_compute_forward_out_prod_q_f32); this backend only took F32, which made the
+    // backward through a Q4_K_M base impossible on CUDA while it worked on CPU
+    // (Continuum card dd109d3c, 2026-09-26). Dequantize W to a transient F32 copy in the
+    // pool through the same converters mul_mat uses, then run the F32 GEMMs below on it.
+    // The copy is dense (contiguous rows), so the F32 strides are derived from ne0x,
+    // not from the quantized tensor's nb0x.
+    cudaStream_t stream = ctx.stream();
+    ggml_cuda_pool_alloc<float> src0_f32_alloc(ctx.pool());
+    const float * src0_d;
+    int64_t lda, s02, s03;
+    if (src0->type == GGML_TYPE_F32) {
+        src0_d = (const float *) src0->data;
+        lda = nb01 / sizeof(float);
+        s02 = nb02 / sizeof(float);
+        s03 = nb03 / sizeof(float);
+    } else {
+        GGML_ASSERT(ggml_is_contiguous(src0));
+        const to_fp32_cuda_t to_fp32 = ggml_get_to_fp32_cuda(src0->type);
+        GGML_ASSERT(to_fp32 != nullptr);
+        const int64_t n = ggml_nelements(src0);
+        float * tmp = src0_f32_alloc.alloc(n);
+        to_fp32(src0->data, tmp, n, stream);
+        CUDA_CHECK(cudaGetLastError());
+        src0_d = tmp;
+        lda = ne00;
+        s02 = ne00 * ne01;
+        s03 = ne00 * ne01 * ne02;
+    }
+
     GGML_ASSERT(ne01 == ne11);
     GGML_ASSERT(ne0 == ne00);
     GGML_ASSERT(ne1 == ne10);
@@ -44,17 +76,14 @@ void ggml_cuda_out_prod(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
     GGML_ASSERT(ne2 == src1->ne[2]);
     GGML_ASSERT(ne3 == src1->ne[3]);
 
-    const float * src0_d = (const float *) src0->data;
     const float * src1_d = (const float *) src1->data;
     float       *  dst_d = (float       *)  dst->data;
 
-    cudaStream_t   stream = ctx.stream();
     cublasHandle_t handle = ctx.cublas_handle();
 
     const float alpha = 1.0f;
     const float beta = 0.0f;
 
-    const int64_t lda = nb01 / sizeof(float);
     const int64_t ldc = nb1  / sizeof(float);
 
     const bool src1_T = ggml_is_transposed(src1);
@@ -63,8 +92,7 @@ void ggml_cuda_out_prod(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
     GGML_ASSERT(                             (src1_T ?        nb11 :        nb10) == sizeof(float));
 
     // data strides in dimensions 2/3
-    const size_t s02 = nb02 / sizeof(float);
-    const size_t s03 = nb03 / sizeof(float);
+    // s02/s03: src0's strides in dims 2/3, set above per src0's storage (F32 view or the dense copy).
     const size_t s12 = nb12 / sizeof(float);
     const size_t s13 = nb13 / sizeof(float);
     const size_t s2  = nb2  / sizeof(float);