diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f1d04fce9267..4d81bd82b8c2 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/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) { diff --git a/ggml/src/ggml-cuda/out-prod.cu b/ggml/src/ggml-cuda/out-prod.cu index c46e0455de40..02a4a64609dc 100644 --- a/ggml/src/ggml-cuda/out-prod.cu +++ b/ggml/src/ggml-cuda/out-prod.cu @@ -1,4 +1,5 @@ #include "out-prod.cuh" +#include "convert.cuh" #include @@ -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 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);