source/prism_native_q1_lora.patch

diff --git a/ggml/include/ggml-opt.h b/ggml/include/ggml-opt.h
index 1c2ed79..21eb602 100644
--- a/ggml/include/ggml-opt.h
+++ b/ggml/include/ggml-opt.h
@@ -159,6 +159,24 @@ extern "C" {
 
     GGML_API const char * ggml_opt_optimizer_name(enum ggml_opt_optimizer_type);
 
+
+// PRISM_STEP10_GGML_OPT_API_BEGIN
+// Step 10 uses GGML only for deterministic forward/backward. AdamW state and
+// updates are managed by the native packed-Q1 LoRA trainer so they can be
+// checkpointed exactly, clipped, accumulated, and resumed mid-window.
+GGML_API void ggml_opt_configure_gradient_only(
+        ggml_opt_context_t opt_ctx,
+        int32_t            opt_period);
+GGML_API void ggml_opt_reset_gradient_cycle(ggml_opt_context_t opt_ctx);
+GGML_API void ggml_opt_zero_grad_accumulators(ggml_opt_context_t opt_ctx);
+GGML_API struct ggml_tensor * ggml_opt_grad_from_active_graph(
+        ggml_opt_context_t   opt_ctx,
+        struct ggml_tensor * node);
+GGML_API struct ggml_tensor * ggml_opt_grad_acc_from_active_graph(
+        ggml_opt_context_t   opt_ctx,
+        struct ggml_tensor * node);
+// PRISM_STEP10_GGML_OPT_API_END
+
     // ====== Optimization Result ======
 
     GGML_API ggml_opt_result_t ggml_opt_result_init(void);
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 9b60d23..6f930b8 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -11,6 +11,9 @@
 #include <algorithm>
 #include <cfloat>
 #include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
 
 // ggml_compute_forward_dup
 
@@ -1156,45 +1159,38 @@ void ggml_compute_forward_add1(
 static void ggml_compute_forward_acc_f32(
         const ggml_compute_params * params,
         ggml_tensor * dst) {
-
+    // PRISM_Q1_LORA_STRIDED_CPU_ACC_V1
     const ggml_tensor * src0 = dst->src[0];
     const ggml_tensor * src1 = dst->src[1];
 
+    GGML_ASSERT(src0->type == GGML_TYPE_F32);
+    GGML_ASSERT(src1->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type  == GGML_TYPE_F32);
     GGML_ASSERT(ggml_are_same_shape(src0, dst));
     GGML_ASSERT(ggml_is_contiguous(dst) && ggml_is_contiguous(src0));
 
-    // view src0 and dst with these strides and data offset inbytes during acc
-    // nb0 is implicitly element_size because src0 and dst are contiguous
-    size_t nb1     = ((int32_t *) dst->op_params)[0];
-    size_t nb2     = ((int32_t *) dst->op_params)[1];
-    size_t nb3     = ((int32_t *) dst->op_params)[2];
-    size_t offset  = ((int32_t *) dst->op_params)[3];
-    bool   inplace = (bool) ((int32_t *) dst->op_params)[4];
+    size_t nb1 = ((int32_t *) dst->op_params)[0];
+    size_t nb2 = ((int32_t *) dst->op_params)[1];
+    size_t nb3 = ((int32_t *) dst->op_params)[2];
+    size_t offset = ((int32_t *) dst->op_params)[3];
+    bool inplace = (bool) ((int32_t *) dst->op_params)[4];
 
     if (!inplace) {
         if (params->ith == 0) {
-            // memcpy needs to be synchronized across threads to avoid race conditions.
-            // => do it in INIT phase
-            memcpy(
-                ((char *)  dst->data),
-                ((char *) src0->data),
-                ggml_nbytes(dst));
+            memcpy((char *) dst->data, (char *) src0->data, ggml_nbytes(dst));
         }
         ggml_barrier(params->threadpool);
     }
 
     const int ith = params->ith;
     const int nth = params->nth;
-
     const int nr = ggml_nrows(src1);
     const int nc = src1->ne[0];
 
     GGML_TENSOR_LOCALS(int64_t, ne1, src1, ne)
     GGML_TENSOR_LOCALS(size_t,  nb1, src1, nb)
 
-    // src0 and dst as viewed during acc
     const size_t nb0 = ggml_element_size(src0);
-
     const size_t nb00 = nb0;
     const size_t nb01 = nb1;
     const size_t nb02 = nb2;
@@ -1202,34 +1198,54 @@ static void ggml_compute_forward_acc_f32(
 
     GGML_ASSERT(offset + (ne10 == 0 ? 0 : ne10-1)*nb0  + (ne11 == 0 ? 0 : ne11-1)*nb1  + (ne12 == 0 ? 0 : ne12-1)*nb2  + (ne13 == 0 ? 0 : ne13-1)*nb3  < ggml_nbytes(dst));
     GGML_ASSERT(offset + (ne10 == 0 ? 0 : ne10-1)*nb00 + (ne11 == 0 ? 0 : ne11-1)*nb01 + (ne12 == 0 ? 0 : ne12-1)*nb02 + (ne13 == 0 ? 0 : ne13-1)*nb03 < ggml_nbytes(src0));
+    GGML_ASSERT((ne10 == 0 ? 0 : ne10-1)*nb10 + (ne11 == 0 ? 0 : ne11-1)*nb11 + (ne12 == 0 ? 0 : ne12-1)*nb12 + (ne13 == 0 ? 0 : ne13-1)*nb13 < ggml_nbytes(src1));
 
-    GGML_ASSERT(nb10 == sizeof(float));
+    const bool src1_dim0_contiguous = nb10 == sizeof(float);
+    const bool prism_training = std::getenv("PRISM_Q1_LORA_TRAINING") != nullptr;
+    if (!src1_dim0_contiguous) {
+        GGML_ASSERT(prism_training);
+        if (ith == 0) {
+            static int prism_strided_acc_count = 0;
+            if (prism_strided_acc_count < 128) {
+                std::fprintf(stderr,
+                    "PRISM_Q1_LORA_STRIDED_CPU_ACC count=%d dst_name='%s' src1_name='%s' "
+                    "src1_op=%s src1_shape=[%lld,%lld,%lld,%lld] "
+                    "src1_strides=[%zu,%zu,%zu,%zu]\n",
+                    prism_strided_acc_count + 1, dst->name, src1->name, ggml_op_name(src1->op),
+                    (long long) src1->ne[0], (long long) src1->ne[1],
+                    (long long) src1->ne[2], (long long) src1->ne[3],
+                    src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3]);
+            }
+            prism_strided_acc_count++;
+        }
+    }
 
-    // rows per thread
     const int dr = (nr + nth - 1)/nth;
-
-    // row range for this thread
     const int ir0 = dr*ith;
     const int ir1 = MIN(ir0 + dr, nr);
 
     for (int ir = ir0; ir < ir1; ++ir) {
-        // src0 and dst are viewed with shape of src1 and offset
-        // => same indices
         const int i3 = ir/(ne12*ne11);
         const int i2 = (ir - i3*ne12*ne11)/ne11;
-        const int i1 = (ir - i3*ne12*ne11 - i2*ne11);
+        const int i1 = ir - i3*ne12*ne11 - i2*ne11;
+
+        float * dst_row = (float *) ((char *) dst->data + i3*nb3 + i2*nb2 + i1*nb1 + offset);
+        const float * src0_row = (const float *) ((const char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + offset);
+        const char * src1_row = (const char *) src1->data + i3*nb13 + i2*nb12 + i1*nb11;
 
+        if (src1_dim0_contiguous) {
 #ifdef GGML_USE_ACCELERATE
-        vDSP_vadd(
-                (float *) ((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + offset), 1,
-                (float *) ((char *) src1->data + i3*nb13 + i2*nb12 + i1*nb11), 1,
-                (float *) ((char *) dst->data  + i3*nb3  + i2*nb2  + i1*nb1  + offset), 1, nc);
+            vDSP_vadd(src0_row, 1, (const float *) src1_row, 1, dst_row, 1, nc);
 #else
-        ggml_vec_add_f32(nc,
-                (float *) ((char *)  dst->data + i3*nb3  + i2*nb2  + i1*nb1  + offset),
-                (float *) ((char *) src0->data + i3*nb03 + i2*nb02 + i1*nb01 + offset),
-                (float *) ((char *) src1->data + i3*nb13 + i2*nb12 + i1*nb11));
+            ggml_vec_add_f32(nc, dst_row, src0_row, (const float *) src1_row);
 #endif
+        } else {
+            for (int i0 = 0; i0 < nc; ++i0) {
+                float value;
+                memcpy(&value, src1_row + (size_t) i0*nb10, sizeof(value));
+                dst_row[i0] = src0_row[i0] + value;
+            }
+        }
     }
 }
 
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 13e1b8a..0499fb4 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -29,6 +29,7 @@
 #include "ggml-cuda/im2col.cuh"
 #include "ggml-cuda/mmf.cuh"
 #include "ggml-cuda/mmq.cuh"
+#include "ggml-cuda/mulmat-q1-f32.cuh"
 #include "ggml-cuda/mmvf.cuh"
 #include "ggml-cuda/mmvq.cuh"
 #include "ggml-cuda/norm.cuh"
@@ -2542,6 +2543,32 @@ bool ggml_cuda_mul_mat_q1_hopper(ggml_backend_cuda_context & ctx, const ggml_ten
 static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     const bool split = ggml_backend_buft_is_cuda_split(src0->buffer->buft);
 
+    // PRISM_Q1_EXACT_TRAIN_FORWARD_V1
+    // GGML_PREC_F32 is the explicit selector for the
+    // mathematically exact packed-Q1 training path.
+    // Default precision retains the existing inference
+    // MMQ/MMVQ implementation.
+    const int32_t precision =
+        ggml_get_op_params_i32(dst, 0);
+
+    if (
+        !split
+        && precision == GGML_PREC_F32
+        && src0->type == GGML_TYPE_Q1_0
+        && src1->type == GGML_TYPE_F32
+        && dst->type == GGML_TYPE_F32
+        && !ggml_is_transposed(src0)
+        && !ggml_is_transposed(src1)
+    ) {
+        ggml_cuda_mul_mat_q1_f32_exact(
+            ctx,
+            src0,
+            src1,
+            dst);
+
+        return;
+    }
+
     // If src0 is a temporary compute buffer it may have some padding that needs to be cleared for mul_mat_vec_q or mul_mat_q.
     // But if src0 is also a view of another tensor then this cannot be done safely because it may overwrite valid tensor data.
     // Therefore, in such cases use cuBLAS.
@@ -5188,7 +5215,16 @@ 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;
+            // PRISM_Q1_EXACT_TRAIN_BACKWARD_V1
+            return
+                op->type == GGML_TYPE_F32
+                && op->src[1]->type == GGML_TYPE_F32
+                && (
+                    op->src[0]->type == GGML_TYPE_F32
+                    || (
+                        op->src[0]->type == GGML_TYPE_Q1_0
+                        && !ggml_is_transposed(op->src[0])
+                        && op->src[0]->ne[0] % QK1_0 == 0));
         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 499903d..843ce8b 100644
--- a/ggml/src/ggml-cuda/out-prod.cu
+++ b/ggml/src/ggml-cuda/out-prod.cu
@@ -1,83 +1,565 @@
 #include "out-prod.cuh"
 
+#include <algorithm>
 #include <cstdint>
+#include <limits>
 
-void ggml_cuda_out_prod(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
-    const ggml_tensor * src0 = dst->src[0];
-    const ggml_tensor * src1 = dst->src[1];
+// PRISM_Q1_EXACT_TRAIN_BACKWARD_V1
+//
+// Frozen packed-Q1 MUL_MAT backward requires:
+//
+//     dX = OUT_PROD(W_q1, transpose(dY))
+//
+// This implementation computes only dX. It never creates a Q1
+// weight gradient and never expands the full packed matrix.
+
+static __device__ __forceinline__ float q1_out_prod_value(
+        const char * row,
+        const size_t nb00,
+        const int64_t element_index) {
+    const int64_t block_index =
+        element_index / QK1_0;
+
+    const int quant_index =
+        static_cast<int>(
+            element_index % QK1_0);
+
+    const block_q1_0 * block =
+        reinterpret_cast<
+            const block_q1_0 *>(
+                row
+                + block_index
+                * nb00);
+
+    const uint8_t packed =
+        block->qs[
+            quant_index >> 3
+        ];
+
+    const int bit =
+        (
+            packed
+            >> (
+                quant_index
+                & 7)
+        )
+        & 1;
+
+    const float scale =
+        __half2float(
+            block->d);
+
+    return bit
+        ? scale
+        : -scale;
+}
+
+
+static __global__ void out_prod_q1_f32_kernel(
+        const char * __restrict__ src0,
+        const char * __restrict__ src1,
+        char * __restrict__ dst,
+
+        const int64_t ne0,
+        const int64_t ne1,
+        const int64_t ne01,
+        const int64_t ne2,
+        const int64_t ne3,
+
+        const size_t nb00,
+        const size_t nb01,
+        const size_t nb02,
+        const size_t nb03,
+
+        const size_t nb10,
+        const size_t nb11,
+        const size_t nb12,
+        const size_t nb13,
+
+        const size_t nb0,
+        const size_t nb1,
+        const size_t nb2,
+        const size_t nb3,
+
+        const int64_t dps2,
+        const int64_t dps3) {
+    const int64_t total =
+        ne0
+        * ne1
+        * ne2
+        * ne3;
+
+    const int64_t grid_stride =
+        static_cast<int64_t>(
+            blockDim.x)
+        * static_cast<int64_t>(
+            gridDim.x);
+
+    for (
+        int64_t linear =
+            static_cast<int64_t>(
+                blockIdx.x)
+            * blockDim.x
+            + threadIdx.x;
+
+        linear < total;
+
+        linear += grid_stride
+    ) {
+        int64_t remaining = linear;
+
+        const int64_t i0 =
+            remaining % ne0;
+        remaining /= ne0;
+
+        const int64_t i1 =
+            remaining % ne1;
+        remaining /= ne1;
+
+        const int64_t i2 =
+            remaining % ne2;
+        remaining /= ne2;
+
+        const int64_t i3 =
+            remaining;
+
+        const int64_t src0_i2 =
+            i2 / dps2;
+
+        const int64_t src0_i3 =
+            i3 / dps3;
+
+        float sum = 0.0f;
+
+        for (
+            int64_t k = 0;
+            k < ne01;
+            ++k
+        ) {
+            const char * weight_row =
+                src0
+                + src0_i3 * nb03
+                + src0_i2 * nb02
+                + k * nb01;
+
+            const float weight =
+                q1_out_prod_value(
+                    weight_row,
+                    nb00,
+                    i0);
+
+            const float gradient =
+                *reinterpret_cast<
+                    const float *>(
+                        src1
+                        + i3 * nb13
+                        + i2 * nb12
+                        + k * nb11
+                        + i1 * nb10);
+
+            sum += weight * gradient;
+        }
+
+        *reinterpret_cast<float *>(
+            dst
+            + i3 * nb3
+            + i2 * nb2
+            + i1 * nb1
+            + i0 * nb0) = sum;
+    }
+}
+
+
+static void ggml_cuda_out_prod_q1_f32(
+        ggml_backend_cuda_context & ctx,
+        ggml_tensor * dst) {
+    const ggml_tensor * src0 =
+        dst->src[0];
+
+    const ggml_tensor * src1 =
+        dst->src[1];
 
     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);
+    GGML_ASSERT(
+        src0->type
+        == GGML_TYPE_Q1_0);
+
+    GGML_ASSERT(
+        src1->type
+        == GGML_TYPE_F32);
+
+    GGML_ASSERT(
+        dst->type
+        == GGML_TYPE_F32);
+
+    GGML_ASSERT(
+        !ggml_is_transposed(src0));
+
+    GGML_ASSERT(
+        ne00
+        % QK1_0
+        == 0);
+
+    GGML_ASSERT(
+        nb00
+        == sizeof(block_q1_0));
+
+    GGML_ASSERT(
+        ne01
+        == ne11);
+
+    GGML_ASSERT(
+        ne0
+        == ne00);
+
+    GGML_ASSERT(
+        ne1
+        == ne10);
+
+    GGML_ASSERT(
+        ne2
+        % ne02
+        == 0);
 
-    GGML_ASSERT(ne01 == ne11);
-    GGML_ASSERT(ne0 == ne00);
-    GGML_ASSERT(ne1 == ne10);
+    GGML_ASSERT(
+        ne3
+        % ne03
+        == 0);
 
-    GGML_ASSERT(ne2 % src0->ne[2] == 0);
-    GGML_ASSERT(ne3 % src0->ne[3] == 0);
+    GGML_ASSERT(
+        ne2
+        == ne12);
 
-    GGML_ASSERT(ne2 == src1->ne[2]);
-    GGML_ASSERT(ne3 == src1->ne[3]);
+    GGML_ASSERT(
+        ne3
+        == ne13);
 
-    const float * src0_d = (const float *) src0->data;
-    const float * src1_d = (const float *) src1->data;
-    float       *  dst_d = (float       *)  dst->data;
+    GGML_ASSERT(
+        nb0
+        == sizeof(float));
 
-    cudaStream_t   stream = ctx.stream();
-    cublasHandle_t handle = ctx.cublas_handle();
+    const int64_t dps2 =
+        ne2 / ne02;
+
+    const int64_t dps3 =
+        ne3 / ne03;
+
+    const int64_t total =
+        ne0
+        * ne1
+        * ne2
+        * ne3;
+
+    constexpr int threads = 256;
+
+    const int64_t required_blocks =
+        (
+            total
+            + threads
+            - 1
+        )
+        / threads;
+
+    const int blocks =
+        static_cast<int>(
+            std::min<int64_t>(
+                required_blocks,
+                65535));
+
+    cudaStream_t stream =
+        ctx.stream();
+
+    out_prod_q1_f32_kernel<<<
+        blocks,
+        threads,
+        0,
+        stream
+    >>>(
+        reinterpret_cast<
+            const char *>(src0->data),
+
+        reinterpret_cast<
+            const char *>(src1->data),
+
+        reinterpret_cast<
+            char *>(dst->data),
+
+        ne0,
+        ne1,
+        ne01,
+        ne2,
+        ne3,
+
+        nb00,
+        nb01,
+        nb02,
+        nb03,
+
+        nb10,
+        nb11,
+        nb12,
+        nb13,
+
+        nb0,
+        nb1,
+        nb2,
+        nb3,
+
+        dps2,
+        dps3);
+
+    CUDA_CHECK(
+        cudaGetLastError());
+}
+
+
+void ggml_cuda_out_prod(
+        ggml_backend_cuda_context & ctx,
+        ggml_tensor * dst) {
+    const ggml_tensor * src0 =
+        dst->src[0];
+
+    const ggml_tensor * src1 =
+        dst->src[1];
+
+    if (
+        src0->type
+            == GGML_TYPE_Q1_0
+        && src1->type
+            == GGML_TYPE_F32
+        && dst->type
+            == GGML_TYPE_F32
+    ) {
+        ggml_cuda_out_prod_q1_f32(
+            ctx,
+            dst);
+
+        return;
+    }
+
+    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);
+
+    GGML_ASSERT(
+        ne01
+        == ne11);
+
+    GGML_ASSERT(
+        ne0
+        == ne00);
+
+    GGML_ASSERT(
+        ne1
+        == ne10);
+
+    GGML_ASSERT(
+        ne2
+        % src0->ne[2]
+        == 0);
+
+    GGML_ASSERT(
+        ne3
+        % src0->ne[3]
+        == 0);
+
+    GGML_ASSERT(
+        ne2
+        == src1->ne[2]);
+
+    GGML_ASSERT(
+        ne3
+        == src1->ne[3]);
+
+    const float * src0_d =
+        reinterpret_cast<
+            const float *>(src0->data);
+
+    const float * src1_d =
+        reinterpret_cast<
+            const float *>(src1->data);
+
+    float * dst_d =
+        reinterpret_cast<
+            float *>(dst->data);
+
+    cudaStream_t stream =
+        ctx.stream();
+
+    cublasHandle_t handle =
+        ctx.cublas_handle();
 
     const float alpha = 1.0f;
     const float beta = 0.0f;
 
-    CUBLAS_CHECK(cublasSetStream(handle, stream));
-
-    const int64_t lda = nb01 / sizeof(float);
-    const int64_t ldc = nb1  / sizeof(float);
-
-    const bool src1_T = ggml_is_transposed(src1);
-    const cublasOperation_t src1_cublas_op =  src1_T ? CUBLAS_OP_N : CUBLAS_OP_T;
-    const int64_t           ldb            = (src1_T ?        nb10 :        nb11) /  sizeof(float);
-    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);
-    const size_t s12 = nb12 / sizeof(float);
-    const size_t s13 = nb13 / sizeof(float);
-    const size_t s2  = nb2  / sizeof(float);
-    const size_t s3  = nb3  / sizeof(float);
-
-    // dps == dst per src0, used for group query attention
-    const int64_t dps2 = ne2 / ne02;
-    const int64_t dps3 = ne3 / ne03;
-
-    if (dps2 == 1 && ne2 > 1) {
-        // src0 has uniform stride s02 along dim 2; batch the inner loop with a strided GEMM
-        GGML_ASSERT(ne2 <= std::numeric_limits<int>::max());
-        const int batch_count = (int) ne2;
-        for (int64_t i3 = 0; i3 < ne3; ++i3) {
+    CUBLAS_CHECK(
+        cublasSetStream(
+            handle,
+            stream));
+
+    const int64_t lda =
+        nb01 / sizeof(float);
+
+    const int64_t ldc =
+        nb1 / sizeof(float);
+
+    const bool src1_T =
+        ggml_is_transposed(src1);
+
+    const cublasOperation_t src1_cublas_op =
+        src1_T
+        ? CUBLAS_OP_N
+        : CUBLAS_OP_T;
+
+    const int64_t ldb =
+        (
+            src1_T
+            ? nb10
+            : nb11
+        )
+        / sizeof(float);
+
+    GGML_ASSERT(
+        (
+            src1_T
+            ? nb11
+            : nb10
+        )
+        == sizeof(float));
+
+    const size_t s02 =
+        nb02 / sizeof(float);
+
+    const size_t s03 =
+        nb03 / sizeof(float);
+
+    const size_t s12 =
+        nb12 / sizeof(float);
+
+    const size_t s13 =
+        nb13 / sizeof(float);
+
+    const size_t s2 =
+        nb2 / sizeof(float);
+
+    const size_t s3 =
+        nb3 / sizeof(float);
+
+    const int64_t dps2 =
+        ne2 / ne02;
+
+    const int64_t dps3 =
+        ne3 / ne03;
+
+    if (
+        dps2 == 1
+        && ne2 > 1
+    ) {
+        GGML_ASSERT(
+            ne2
+            <= std::numeric_limits<int>::max());
+
+        const int batch_count =
+            static_cast<int>(ne2);
+
+        for (
+            int64_t i3 = 0;
+            i3 < ne3;
+            ++i3
+        ) {
             CUBLAS_CHECK(
-                cublasSgemmStridedBatched(handle, CUBLAS_OP_N, src1_cublas_op,
-                        ne0, ne1, ne01,
-                        &alpha, src0_d + (i3/dps3)*s03, lda, s02,
-                                src1_d +  i3     *s13, ldb, s12,
-                        &beta,  dst_d  +  i3     *s3,  ldc, s2,
-                        batch_count));
+                cublasSgemmStridedBatched(
+                    handle,
+                    CUBLAS_OP_N,
+                    src1_cublas_op,
+
+                    ne0,
+                    ne1,
+                    ne01,
+
+                    &alpha,
+
+                    src0_d
+                        + (
+                            i3 / dps3)
+                        * s03,
+
+                    lda,
+                    s02,
+
+                    src1_d
+                        + i3 * s13,
+
+                    ldb,
+                    s12,
+
+                    &beta,
+
+                    dst_d
+                        + i3 * s3,
+
+                    ldc,
+                    s2,
+
+                    batch_count));
         }
     } else {
-        // Fallback: ne2 == 1 (no batching benefit) or dps2 > 1 (src0 broadcast along dim 2
-        // with non-uniform stride; would need cublasSgemmBatched with pointer arrays).
-        for (int64_t i3 = 0; i3 < ne3; ++i3) {
-            for (int64_t i2 = 0; i2 < ne2; ++i2) {
+        for (
+            int64_t i3 = 0;
+            i3 < ne3;
+            ++i3
+        ) {
+            for (
+                int64_t i2 = 0;
+                i2 < ne2;
+                ++i2
+            ) {
                 CUBLAS_CHECK(
-                    cublasSgemm(handle, CUBLAS_OP_N, src1_cublas_op,
-                            ne0, ne1, ne01,
-                            &alpha, src0_d + (i3/dps3)*s03 + (i2/dps2)*s02, lda,
-                                    src1_d +  i3      *s13 +  i2      *s12, ldb,
-                            &beta,  dst_d  +  i3      *s3  +  i2      *s2,  ldc));
+                    cublasSgemm(
+                        handle,
+                        CUBLAS_OP_N,
+                        src1_cublas_op,
+
+                        ne0,
+                        ne1,
+                        ne01,
+
+                        &alpha,
+
+                        src0_d
+                            + (
+                                i3 / dps3)
+                            * s03
+                            + (
+                                i2 / dps2)
+                            * s02,
+
+                        lda,
+
+                        src1_d
+                            + i3 * s13
+                            + i2 * s12,
+
+                        ldb,
+
+                        &beta,
+
+                        dst_d
+                            + i3 * s3
+                            + i2 * s2,
+
+                        ldc));
             }
         }
     }
diff --git a/ggml/src/ggml-cuda/unary.cu b/ggml/src/ggml-cuda/unary.cu
index 4cb805f..f555039 100644
--- a/ggml/src/ggml-cuda/unary.cu
+++ b/ggml/src/ggml-cuda/unary.cu
@@ -1,6 +1,9 @@
 #include "unary.cuh"
 #include "convert.cuh"
 
+#include <cstdio>
+#include <cstdlib>
+
 static __device__ __forceinline__ float op_abs(float x) {
     return fabsf(x);
 }
@@ -127,6 +130,34 @@ static __global__ void unary_op_kernel(const T * x, T * dst, const int k) {
     dst[i] = (T)op((float)x[i]);
 }
 
+// PRISM_Q1_LORA_STRIDED_UNARY_KERNEL_V2
+template <float (*op)(float), typename T>
+static __global__ void unary_op_strided_kernel(
+        const char * x,
+        char * dst,
+        int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3,
+        size_t snb0, size_t snb1, size_t snb2, size_t snb3,
+        size_t dnb0, size_t dnb1, size_t dnb2, size_t dnb3,
+        int64_t n) {
+    ggml_cuda_pdl_lc();
+    const int64_t linear = (int64_t) blockDim.x*blockIdx.x + threadIdx.x;
+    if (linear >= n) {
+        return;
+    }
+
+    int64_t value = linear;
+    const int64_t i0 = value % ne0; value /= ne0;
+    const int64_t i1 = value % ne1; value /= ne1;
+    const int64_t i2 = value % ne2; value /= ne2;
+    const int64_t i3 = value % ne3;
+
+    ggml_cuda_pdl_sync();
+
+    const T * xp = reinterpret_cast<const T *>(x + i0*snb0 + i1*snb1 + i2*snb2 + i3*snb3);
+    T * dp = reinterpret_cast<T *>(dst + i0*dnb0 + i1*dnb1 + i2*dnb2 + i3*dnb3);
+    *dp = (T) op((float) *xp);
+}
+
 template <float (*op)(float), typename T>
 static void unary_cuda(const T * x, T * dst, const int k, cudaStream_t stream) {
     const int num_blocks = (k + CUDA_NEG_BLOCK_SIZE - 1) / CUDA_NEG_BLOCK_SIZE;
@@ -134,26 +165,62 @@ static void unary_cuda(const T * x, T * dst, const int k, cudaStream_t stream) {
     ggml_cuda_kernel_launch(unary_op_kernel<op, T>, launch_params, x, dst, k);
 }
 
+template <float (*op)(float), typename T>
+static void unary_cuda_strided(const ggml_tensor * src0, ggml_tensor * dst, cudaStream_t stream) {
+    const int64_t n = ggml_nelements(src0);
+    const int64_t num_blocks = (n + CUDA_NEG_BLOCK_SIZE - 1) / CUDA_NEG_BLOCK_SIZE;
+    const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params((dim3) num_blocks, CUDA_NEG_BLOCK_SIZE, 0, stream);
+    ggml_cuda_kernel_launch(unary_op_strided_kernel<op, T>, launch_params,
+        (const char *) src0->data, (char *) dst->data,
+        src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
+        src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+        dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], n);
+}
+
 template <float (*op)(float)>
 void ggml_cuda_op_unary(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+    // PRISM_Q1_LORA_STRIDED_UNARY_DISPATCH_V2
     const ggml_tensor * src0 = dst->src[0];
-    const void * src0_d = src0->data;
-    void * dst_d = dst->data;
     cudaStream_t stream = ctx.stream();
 
-    GGML_ASSERT(ggml_is_contiguous(src0));
-
     GGML_ASSERT(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16);
     GGML_ASSERT( dst->type == GGML_TYPE_F32 ||  dst->type == GGML_TYPE_F16);
     GGML_ASSERT(src0->type == dst->type);
+    GGML_ASSERT(ggml_are_same_shape(src0, dst));
+
+    const bool contiguous = ggml_is_contiguous(src0) && ggml_is_contiguous(dst);
+    if (contiguous) {
+        if (src0->type == GGML_TYPE_F16) {
+            unary_cuda<op>((const half *) src0->data, (half *) dst->data, ggml_nelements(src0), stream);
+        } else {
+            unary_cuda<op>((const float *) src0->data, (float *) dst->data, ggml_nelements(src0), stream);
+        }
+        return;
+    }
+
+    const bool prism_training = std::getenv("PRISM_Q1_LORA_TRAINING") != nullptr;
+    GGML_ASSERT(prism_training);
+
+    static int prism_strided_unary_count = 0;
+    if (prism_strided_unary_count < 128) {
+        std::fprintf(stderr,
+            "PRISM_Q1_LORA_STRIDED_UNARY count=%d src_name='%s' src_op=%s "
+            "shape=[%lld,%lld,%lld,%lld] src_strides=[%zu,%zu,%zu,%zu] "
+            "dst_strides=[%zu,%zu,%zu,%zu]\n",
+            prism_strided_unary_count + 1, src0->name, ggml_op_name(src0->op),
+            (long long) src0->ne[0], (long long) src0->ne[1],
+            (long long) src0->ne[2], (long long) src0->ne[3],
+            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3]);
+    }
+    prism_strided_unary_count++;
 
     if (src0->type == GGML_TYPE_F16) {
-        unary_cuda<op>((const half *)src0_d, (half *)dst_d, ggml_nelements(src0), stream);
+        unary_cuda_strided<op, half>(src0, dst, stream);
     } else {
-        unary_cuda<op>((const float *)src0_d, (float *)dst_d, ggml_nelements(src0), stream);
+        unary_cuda_strided<op, float>(src0, dst, stream);
     }
 }
-
 void ggml_cuda_op_abs(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
     ggml_cuda_op_unary<op_abs>(ctx, dst);
 }
diff --git a/ggml/src/ggml-opt.cpp b/ggml/src/ggml-opt.cpp
index 53903de..f7a6e90 100644
--- a/ggml/src/ggml-opt.cpp
+++ b/ggml/src/ggml-opt.cpp
@@ -8,6 +8,8 @@
 #include <algorithm>
 #include <cmath>
 #include <cstdint>
+#include <cstdlib>
+#include <cstring>
 #include <cinttypes>
 #include <map>
 #include <random>
@@ -325,8 +327,37 @@ static void ggml_opt_build(ggml_opt_context_t opt_ctx) {
 
     const enum ggml_opt_optimizer_type optimizer = opt_ctx->optimizer;
 
-    const bool accumulate = opt_ctx->build_type_alloc >= GGML_OPT_BUILD_TYPE_GRAD &&
-        !(opt_ctx->static_graphs && opt_ctx->build_type_alloc == GGML_OPT_BUILD_TYPE_OPT && opt_ctx->opt_period == 1);
+    // PRISM_STEP10_V46_EXTERNAL_ACCUM_POLICY_BEGIN
+    const char * prism_external_accum_env =
+        std::getenv("PRISM_STEP10_EXTERNAL_GRAD_ACCUM");
+
+    const bool prism_external_host_accum =
+        !opt_ctx->static_graphs
+        && prism_external_accum_env
+        && prism_external_accum_env[0] == '1'
+        && prism_external_accum_env[1] == '\0';
+
+    // Dynamic Step 10 graphs expose direct parameter gradients.
+    // Accumulation, clipping, AdamW, checkpointing, and resume remain
+    // exclusively owned by the deterministic host trainer.
+    const bool accumulate =
+        !prism_external_host_accum
+        && opt_ctx->build_type_alloc >= GGML_OPT_BUILD_TYPE_GRAD
+        && !(opt_ctx->static_graphs
+             && opt_ctx->build_type_alloc == GGML_OPT_BUILD_TYPE_OPT
+             && opt_ctx->opt_period == 1);
+
+    if (prism_external_host_accum) {
+        static bool prism_external_accum_logged = false;
+        if (!prism_external_accum_logged) {
+            fprintf(
+                stderr,
+                "PRISM_STEP10_EXTERNAL_HOST_ACCUM=1 "
+                "ggml_parameter_accumulate=0\n");
+            prism_external_accum_logged = true;
+        }
+    }
+    // PRISM_STEP10_V46_EXTERNAL_ACCUM_POLICY_END
 
     const bool need_momenta = opt_ctx->build_type_alloc == GGML_OPT_BUILD_TYPE_OPT &&
         opt_ctx->optimizer == GGML_OPT_OPTIMIZER_TYPE_ADAMW;
@@ -631,8 +662,193 @@ struct ggml_tensor * ggml_opt_ncorrect(ggml_opt_context_t opt_ctx) {
 }
 
 struct ggml_tensor * ggml_opt_grad_acc(ggml_opt_context_t opt_ctx, struct ggml_tensor * node) {
-    return ggml_graph_get_grad_acc(opt_ctx->gb_opt, node);
+    if (!opt_ctx || !node) {
+        return nullptr;
+    }
+    struct ggml_cgraph * graph = opt_ctx->allocated_graph;
+    if (!graph) {
+        switch (opt_ctx->build_type) {
+            case GGML_OPT_BUILD_TYPE_FORWARD: graph = opt_ctx->gf;      break;
+            case GGML_OPT_BUILD_TYPE_GRAD:    graph = opt_ctx->gb_grad; break;
+            case GGML_OPT_BUILD_TYPE_OPT:     graph = opt_ctx->gb_opt;  break;
+        }
+    }
+    return graph ? ggml_graph_get_grad_acc(graph, node) : nullptr;
+}
+
+
+// PRISM_STEP10_GGML_OPT_IMPL_BEGIN
+// PRISM_STEP10_V451_FORCE_OPT_CONFIG_BEGIN
+void ggml_opt_configure_gradient_only(
+        ggml_opt_context_t opt_ctx,
+        int32_t            opt_period) {
+    GGML_ASSERT(opt_ctx != nullptr);
+    GGML_ASSERT(opt_period >= 1);
+    GGML_ASSERT(
+        opt_ctx->build_type_alloc
+        == GGML_OPT_BUILD_TYPE_OPT);
+
+    const char * force_opt_env =
+        std::getenv("PRISM_STEP10_FORCE_OPT_BACKWARD");
+
+    const bool force_opt =
+        force_opt_env
+        && force_opt_env[0] == '1'
+        && force_opt_env[1] == '\0';
+
+    opt_ctx->build_type =
+        force_opt
+            ? GGML_OPT_BUILD_TYPE_OPT
+            : GGML_OPT_BUILD_TYPE_GRAD;
+
+    opt_ctx->opt_period =
+        force_opt
+            ? 1
+            : opt_period;
+
+    opt_ctx->opt_i = 0;
+
+    fprintf(
+        stderr,
+        "PRISM_STEP10_FORCE_OPT_BACKWARD=%d "
+        "build_type_alloc=OPT opt_period=%d\n",
+        force_opt ? 1 : 0,
+        opt_ctx->opt_period);
+}
+// PRISM_STEP10_V451_FORCE_OPT_CONFIG_END
+
+void ggml_opt_reset_gradient_cycle(ggml_opt_context_t opt_ctx) {
+    GGML_ASSERT(opt_ctx != nullptr);
+    GGML_ASSERT(opt_ctx->build_type_alloc == GGML_OPT_BUILD_TYPE_OPT);
+    opt_ctx->opt_i      = 0;
+    opt_ctx->build_type = GGML_OPT_BUILD_TYPE_GRAD;
+}
+
+// PRISM_STEP10_V47_LOSS_SEED_PRESERVE_BEGIN
+void ggml_opt_zero_grad_accumulators(
+        ggml_opt_context_t opt_ctx) {
+    GGML_ASSERT(opt_ctx != nullptr);
+    GGML_ASSERT(opt_ctx->allocated_graph != nullptr);
+    GGML_ASSERT(opt_ctx->loss != nullptr);
+
+    ggml_tensor * loss_grad =
+        ggml_graph_get_grad_acc(
+            opt_ctx->allocated_graph,
+            opt_ctx->loss);
+
+    GGML_ASSERT(
+        loss_grad != nullptr
+        && "active backward graph has no loss gradient accumulator");
+
+    GGML_ASSERT(
+        loss_grad->type == GGML_TYPE_F32
+        && ggml_is_scalar(loss_grad));
+
+    float seed_before = NAN;
+
+    if (loss_grad->buffer) {
+        ggml_backend_tensor_get(
+            loss_grad,
+            &seed_before,
+            0,
+            sizeof(seed_before));
+    }
+
+    size_t zeroed_accumulators = 0;
+    size_t preserved_loss_entries = 0;
+
+    for (ggml_tensor * tensor : opt_ctx->grad_accs) {
+        if (!tensor || !tensor->buffer) {
+            continue;
+        }
+
+        if (tensor == loss_grad) {
+            ++preserved_loss_entries;
+            continue;
+        }
+
+        std::vector<uint8_t> zeros(
+            ggml_nbytes(tensor),
+            0);
+
+        ggml_backend_tensor_set(
+            tensor,
+            zeros.data(),
+            0,
+            zeros.size());
+
+        ++zeroed_accumulators;
+    }
+
+    const float loss_seed = 1.0f;
+
+    ggml_backend_tensor_set(
+        loss_grad,
+        &loss_seed,
+        0,
+        sizeof(loss_seed));
+
+    float seed_after = NAN;
+
+    ggml_backend_tensor_get(
+        loss_grad,
+        &seed_after,
+        0,
+        sizeof(seed_after));
+
+    fprintf(
+        stderr,
+        "PRISM_STEP10_LOSS_SEED_BEFORE=%.9g\n"
+        "PRISM_STEP10_LOSS_SEED_AFTER=%.9g\n"
+        "PRISM_STEP10_NONLOSS_ACCUMULATORS_ZEROED=%zu\n"
+        "PRISM_STEP10_LOSS_ACCUMULATOR_PRESERVED=%zu\n",
+        static_cast<double>(seed_before),
+        static_cast<double>(seed_after),
+        zeroed_accumulators,
+        preserved_loss_entries);
+
+    GGML_ASSERT(
+        seed_after == 1.0f
+        && "failed to restore loss backward seed");
+}
+// PRISM_STEP10_V47_LOSS_SEED_PRESERVE_END
+
+static ggml_cgraph * prism_step10_active_graph(
+        ggml_opt_context_t opt_ctx) {
+    if (!opt_ctx) {
+        return nullptr;
+    }
+
+    if (opt_ctx->allocated_graph) {
+        return opt_ctx->allocated_graph;
+    }
+
+    switch (opt_ctx->build_type) {
+        case GGML_OPT_BUILD_TYPE_FORWARD:
+            return opt_ctx->gf;
+        case GGML_OPT_BUILD_TYPE_GRAD:
+            return opt_ctx->gb_grad;
+        case GGML_OPT_BUILD_TYPE_OPT:
+            return opt_ctx->gb_opt;
+    }
+
+    return nullptr;
+}
+
+struct ggml_tensor * ggml_opt_grad_from_active_graph(
+        ggml_opt_context_t   opt_ctx,
+        struct ggml_tensor * node) {
+    ggml_cgraph * graph = prism_step10_active_graph(opt_ctx);
+    return graph && node ? ggml_graph_get_grad(graph, node) : nullptr;
+}
+
+struct ggml_tensor * ggml_opt_grad_acc_from_active_graph(
+        ggml_opt_context_t   opt_ctx,
+        struct ggml_tensor * node) {
+    ggml_cgraph * graph = prism_step10_active_graph(opt_ctx);
+    return graph && node ? ggml_graph_get_grad_acc(graph, node) : nullptr;
 }
+// PRISM_STEP10_GGML_OPT_IMPL_END
 
 // ====== Optimization Result ======
 
@@ -658,11 +874,15 @@ void ggml_opt_result_ndata(ggml_opt_result_t result, int64_t * ndata) {
 void ggml_opt_result_loss(ggml_opt_result_t result, double * loss, double * unc) {
     const int64_t nbatches = result->loss.size(); // Number of physical batches.
 
+    // PRISM_STEP10_V43_NULL_SAFE_RESULT_LOSS_BEGIN
     if (nbatches == 0) {
         *loss = 0.0;
-        *unc  = NAN;
+        if (unc) {
+            *unc = NAN;
+        }
         return;
     }
+    // PRISM_STEP10_V43_NULL_SAFE_RESULT_LOSS_END
 
     double sum         = 0.0;
     double sum_squared = 0.0;
@@ -822,6 +1042,108 @@ void ggml_opt_eval(ggml_opt_context_t opt_ctx, ggml_opt_result_t result) {
     }
 
     ggml_backend_sched_graph_compute(opt_ctx->backend_sched, opt_ctx->allocated_graph_copy);
+
+    // PRISM_Q1_LORA_CHAIN_AUDIT_HELPER_V1
+    // PRISM_Q1_LORA_ATTENTION_CHAIN_AUDIT_V1
+    // Retired diagnostic identities: the safe post-compute parameter-gradient
+    // audit below supersedes the earlier view-chain debug reader. No callback
+    // reads from dynamic graphs are reintroduced.
+    // PRISM_Q1_LORA_SAFE_GRADIENT_AUDIT_V1
+    //
+    // Dynamic llama optimizer graphs are released later in this
+    // function. Inspect gradients here, after synchronous compute
+    // and while gb_opt still contains its valid mapped parameter
+    // nodes. Do not call ggml_opt_grad_acc() from the later epoch
+    // callback for dynamic graphs.
+    if (
+        getenv("PRISM_Q1_LORA_GRAD_AUDIT") != nullptr
+        && opt_ctx->allocated_graph == opt_ctx->gb_opt
+        && opt_ctx->gb_opt != nullptr
+    ) {
+        for (
+            int i = 0;
+            i < opt_ctx->gb_opt->n_nodes;
+            ++i
+        ) {
+            struct ggml_tensor * node =
+                opt_ctx->gb_opt->nodes[i];
+
+            if (
+                node == nullptr
+                || !(node->flags & GGML_TENSOR_FLAG_PARAM)
+            ) {
+                continue;
+            }
+
+            const bool is_lora_a =
+                strstr(node->name, ".lora_a") != nullptr;
+
+            const bool is_lora_b =
+                strstr(node->name, ".lora_b") != nullptr;
+
+            if (!is_lora_a && !is_lora_b) {
+                continue;
+            }
+
+            struct ggml_tensor * grad =
+                ggml_graph_get_grad(
+                    opt_ctx->gb_opt,
+                    node);
+
+            if (grad == nullptr) {
+                fprintf(
+                    stderr,
+                    "PRISM_Q1_LORA_OPT_GRAD "
+                    "name=%s present=0 max_abs=0\n",
+                    node->name);
+
+                continue;
+            }
+
+            if (grad->type != GGML_TYPE_F32) {
+                fprintf(
+                    stderr,
+                    "PRISM_Q1_LORA_OPT_GRAD "
+                    "name=%s present=1 "
+                    "type=%s max_abs=nan\n",
+                    node->name,
+                    ggml_type_name(grad->type));
+
+                continue;
+            }
+
+            const int64_t count =
+                ggml_nelements(grad);
+
+            std::vector<float> values(
+                static_cast<size_t>(count));
+
+            ggml_backend_tensor_get(
+                grad,
+                values.data(),
+                0,
+                ggml_nbytes(grad));
+
+            double maximum = 0.0;
+
+            for (const float value : values) {
+                maximum = std::max(
+                    maximum,
+                    std::abs(
+                        static_cast<double>(value)));
+            }
+
+            fprintf(
+                stderr,
+                "PRISM_Q1_LORA_OPT_GRAD "
+                "name=%s present=1 "
+                "max_abs=%.12g nelements=%" PRId64 "\n",
+                node->name,
+                maximum,
+                count);
+        }
+    }
+
     opt_ctx->iter += opt_ctx->allocated_graph == opt_ctx->gb_opt;
     opt_ctx->opt_i = (opt_ctx->opt_i + 1) % opt_ctx->opt_period;
 
diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
index de0615f..b1b93f8 100644
--- a/ggml/src/ggml.c
+++ b/ggml/src/ggml.c
@@ -143,6 +143,8 @@ static void ggml_print_backtrace_symbols(void) {
 }
 #elif defined(__APPLE__)
 #include <execinfo.h>
+
+
 static void ggml_print_backtrace_symbols(void) {
     void * trace[100];
     int nptrs = backtrace(trace, sizeof(trace)/sizeof(trace[0]));
@@ -1999,10 +2001,30 @@ struct ggml_tensor * ggml_get_tensor(struct ggml_context * ctx, const char * nam
 
 // ggml_dup
 
+
+// PRISM_Q1_LORA_TRAINING_NO_INPLACE_V2
+//
+// GGML backward graph construction rejects arithmetic operations
+// whose output aliases a source tensor.
+//
+// During explicit Prism native Q1-LoRA training, standard GGML
+// *_inplace constructors therefore return out-of-place tensors.
+//
+// Normal inference behavior remains unchanged.
+static bool ggml_prism_q1_lora_training_no_inplace(void) {
+    const char * value =
+        getenv("PRISM_Q1_LORA_TRAINING");
+
+    return value != NULL
+        && strcmp(value, "0") != 0;
+}
+
 static struct ggml_tensor * ggml_dup_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * a,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_DUP;
@@ -2032,6 +2054,8 @@ static struct ggml_tensor * ggml_add_impl(
         bool                  inplace) {
     GGML_ASSERT(ggml_can_repeat(b, a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_ADD;
@@ -2119,6 +2143,8 @@ static struct ggml_tensor * ggml_add1_impl(
     GGML_ASSERT(ggml_is_scalar(b));
     GGML_ASSERT(ggml_is_padded_1d(a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_ADD1;
@@ -2158,6 +2184,8 @@ static struct ggml_tensor * ggml_acc_impl(
     GGML_ASSERT(a->type == GGML_TYPE_F32);
     GGML_ASSERT(b->type == GGML_TYPE_F32);
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     int32_t params[] = { nb1, nb2, nb3, offset, inplace ? 1 : 0 };
@@ -2201,6 +2229,8 @@ static struct ggml_tensor * ggml_sub_impl(
         bool                  inplace) {
     GGML_ASSERT(ggml_can_repeat(b, a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_SUB;
@@ -2233,6 +2263,8 @@ static struct ggml_tensor * ggml_mul_impl(
         bool                  inplace) {
     GGML_ASSERT(ggml_can_repeat(b, a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_MUL;
@@ -2265,6 +2297,8 @@ static struct ggml_tensor * ggml_div_impl(
         bool                  inplace) {
     GGML_ASSERT(ggml_can_repeat(b, a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_DIV;
@@ -2294,6 +2328,8 @@ static struct ggml_tensor * ggml_sqr_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * a,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_SQR;
@@ -2320,6 +2356,8 @@ static struct ggml_tensor * ggml_sqrt_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * a,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_SQRT;
@@ -2346,6 +2384,8 @@ static struct ggml_tensor * ggml_log_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * a,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_LOG;
@@ -2396,6 +2436,8 @@ static struct ggml_tensor * ggml_sin_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * a,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_SIN;
@@ -2422,6 +2464,8 @@ static struct ggml_tensor * ggml_cos_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * a,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_COS;
@@ -2723,6 +2767,8 @@ struct ggml_tensor * ggml_leaky_relu(
         struct ggml_tensor  * a,
         float                 negative_slope,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params(result, &negative_slope, sizeof(negative_slope));
@@ -3097,6 +3143,8 @@ static struct ggml_tensor * ggml_norm_impl(
         struct ggml_tensor  * a,
         float                 eps,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params(result, &eps, sizeof(eps));
@@ -3128,6 +3176,8 @@ static struct ggml_tensor * ggml_rms_norm_impl(
         struct ggml_tensor  * a,
         float                 eps,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params(result, &eps, sizeof(eps));
@@ -3178,6 +3228,8 @@ static struct ggml_tensor * ggml_group_norm_impl(
         int                   n_groups,
         float                 eps,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params_i32(result, 0, n_groups);
@@ -3212,6 +3264,8 @@ static struct ggml_tensor * ggml_l2_norm_impl(
         struct ggml_tensor  * a,
         float                 eps,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params_f32(result, 0, eps);
@@ -3359,7 +3413,30 @@ static struct ggml_tensor * ggml_scale_impl(
         float                 s,
         float                 b,
         bool                  inplace) {
-    GGML_ASSERT(ggml_is_padded_1d(a));
+    // PRISM_Q1_LORA_TRAINING_CONTIGUOUS_SCALE_V1
+    // Preserve the stock padded-1D contract outside explicit training.
+    if (!ggml_is_padded_1d(a)) {
+        const bool prism_training = getenv("PRISM_Q1_LORA_TRAINING") != NULL;
+        if (prism_training && !inplace) {
+            static int prism_scale_materialize_count = 0;
+            if (prism_scale_materialize_count < 128) {
+                fprintf(stderr,
+                    "PRISM_Q1_LORA_SCALE_MATERIALIZE count=%d name='%s' op=%s "
+                    "shape=[%" PRId64 ",%" PRId64 ",%" PRId64 ",%" PRId64 "] "
+                    "strides=[%zu,%zu,%zu,%zu]\n",
+                    prism_scale_materialize_count + 1,
+                    a->name, ggml_op_name(a->op),
+                    a->ne[0], a->ne[1], a->ne[2], a->ne[3],
+                    a->nb[0], a->nb[1], a->nb[2], a->nb[3]);
+            }
+            prism_scale_materialize_count++;
+            a = ggml_cont(ctx, a);
+        } else {
+            GGML_ASSERT(ggml_is_padded_1d(a));
+        }
+    }
+
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
 
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
@@ -3416,6 +3493,8 @@ static struct ggml_tensor * ggml_set_impl(
     GGML_ASSERT(ggml_nelements(a) >= ggml_nelements(b));
 
     // make a view of the destination
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     GGML_ASSERT(offset < (size_t)(1 << 30));
@@ -3956,6 +4035,8 @@ static struct ggml_tensor * ggml_diag_mask_inf_impl(
         struct ggml_tensor  * a,
         int                   n_past,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     int32_t params[] = { n_past };
@@ -3988,6 +4069,8 @@ static struct ggml_tensor * ggml_diag_mask_zero_impl(
         struct ggml_tensor  * a,
         int                   n_past,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     int32_t params[] = { n_past };
@@ -4037,6 +4120,8 @@ static struct ggml_tensor * ggml_soft_max_impl(
         GGML_ASSERT(mask);
     }
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     float params[] = { scale, max_bias };
@@ -4104,6 +4189,8 @@ static struct ggml_tensor * ggml_soft_max_ext_back_impl(
         float                 scale,
         float                 max_bias,
         bool                  inplace) {
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     result->op     = GGML_OP_SOFT_MAX_BACK;
@@ -4169,6 +4256,8 @@ static struct ggml_tensor * ggml_rope_impl(
         GGML_ASSERT(c->ne[0] >= n_dims / 2);
     }
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     int32_t params[15] = { /*n_past*/ 0, n_dims, mode, /*n_ctx*/ 0, n_ctx_orig };
@@ -5235,6 +5324,8 @@ static struct ggml_tensor * ggml_fill_impl(
     GGML_ASSERT(a->type == GGML_TYPE_F32 || a->type == GGML_TYPE_F16);
     GGML_ASSERT(ggml_is_contiguous(a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params_f32(result, 0, c);
@@ -5666,6 +5757,8 @@ static struct ggml_tensor * ggml_add_rel_pos_impl(
     GGML_ASSERT(pw->ne[0]*pw->ne[0] == a->ne[0]);
     GGML_ASSERT(pw->ne[1]*pw->ne[2] == a->ne[1]);
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
     ggml_set_op_params_i32(result, 0, inplace ? 1 : 0);
 
@@ -5836,6 +5929,8 @@ static struct ggml_tensor * ggml_unary_impl(
         bool                  inplace) {
     GGML_ASSERT(ggml_is_contiguous_rows(a));
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     ggml_set_op_params_i32(result, 0, (int32_t) op);
@@ -5871,6 +5966,8 @@ static struct ggml_tensor * ggml_map_custom1_impl(
         bool                       inplace) {
     GGML_ASSERT(n_tasks == GGML_N_TASKS_MAX || n_tasks > 0);
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     struct ggml_map_custom1_op_params params = {
@@ -5916,6 +6013,8 @@ static struct ggml_tensor * ggml_map_custom2_impl(
         bool                       inplace) {
     GGML_ASSERT(n_tasks == GGML_N_TASKS_MAX || n_tasks > 0);
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     struct ggml_map_custom2_op_params params = {
@@ -5965,6 +6064,8 @@ static struct ggml_tensor * ggml_map_custom3_impl(
         bool                       inplace) {
     GGML_ASSERT(n_tasks == GGML_N_TASKS_MAX || n_tasks > 0);
 
+    inplace = inplace && !ggml_prism_q1_lora_training_no_inplace();
+
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
     struct ggml_map_custom3_op_params params = {
@@ -6510,7 +6611,12 @@ static void ggml_compute_backward(
                 ggml_add_or_set(ctx, cgraph, isrc0, grad);
             }
             if (src1_needs_grads) {
-                ggml_sub_or_set(ctx, cgraph, isrc1, grad);
+                // PRISM_Q1_LORA_SUB_BACKWARD_BROADCAST_V1
+                struct ggml_tensor * tmp = grad;
+                if (!ggml_are_same_shape(grad, src1)) {
+                    tmp = ggml_repeat_back(ctx, tmp, src1);
+                }
+                ggml_sub_or_set(ctx, cgraph, isrc1, tmp);
             }
         } break;
         case GGML_OP_MUL: {
@@ -6530,7 +6636,12 @@ static void ggml_compute_backward(
                 ggml_add_or_set(ctx, cgraph, isrc0, ggml_div(ctx, grad, src1));
             }
             if (src1_needs_grads) {
-                ggml_sub_or_set(ctx, cgraph, isrc1, ggml_mul(ctx, grad, ggml_div(ctx, tensor, src1)));
+                // PRISM_Q1_LORA_DIV_BACKWARD_BROADCAST_V1
+                struct ggml_tensor * tmp = ggml_mul(ctx, grad, ggml_div(ctx, tensor, src1));
+                if (!ggml_are_same_shape(tmp, src1)) {
+                    tmp = ggml_repeat_back(ctx, tmp, src1);
+                }
+                ggml_sub_or_set(ctx, cgraph, isrc1, tmp);
             }
         } break;
         case GGML_OP_SQR: {
@@ -6583,6 +6694,139 @@ static void ggml_compute_backward(
                 ggml_add_or_set(ctx, cgraph, isrc0, ggml_repeat(ctx, grad, src0));
             }
         } break;
+        case GGML_OP_L2_NORM: {
+            if (src0_needs_grads) {
+                // PRISM_Q1_LORA_L2_NORM_BACKWARD_V3
+                float eps;
+                memcpy(&eps, tensor->op_params, sizeof(float));
+
+                // y = x / ||x|| ; dx = (g - y * sum(g*y)) / ||x||
+                struct ggml_tensor * gy = ggml_mul(ctx, grad, tensor);
+                struct ggml_tensor * projection = ggml_sum_rows(ctx, gy);
+                projection = ggml_mul(ctx, tensor, projection);
+                struct ggml_tensor * numerator = ggml_sub(ctx, grad, projection);
+
+                struct ggml_tensor * norm_squared = ggml_sum_rows(ctx, ggml_sqr(ctx, src0));
+                struct ggml_tensor * eps_tensor = ggml_fill(ctx, norm_squared, eps);
+                struct ggml_tensor * norm = ggml_sqrt(ctx, ggml_add(ctx, norm_squared, eps_tensor));
+                ggml_add_or_set(ctx, cgraph, isrc0, ggml_div(ctx, numerator, norm));
+            }
+        } break;
+        case GGML_OP_TRI: {
+            if (src0_needs_grads) {
+                // PRISM_Q1_LORA_TRI_BACKWARD_V1
+                const enum ggml_tri_type type = (enum ggml_tri_type) ggml_get_op_params_i32(tensor, 0);
+                ggml_add_or_set(ctx, cgraph, isrc0, ggml_tri(ctx, ggml_cont(ctx, grad), type));
+            }
+        } break;
+        case GGML_OP_PAD: {
+            if (src0_needs_grads) {
+                // PRISM_Q1_LORA_PAD_BACKWARD_V2
+                GGML_ASSERT(ggml_get_op_params_i32(tensor, 8) == 0 && "circular PAD backward is not implemented");
+                const int64_t lp0 = ggml_get_op_params_i32(tensor, 0);
+                const int64_t lp1 = ggml_get_op_params_i32(tensor, 2);
+                const int64_t lp2 = ggml_get_op_params_i32(tensor, 4);
+                const int64_t lp3 = ggml_get_op_params_i32(tensor, 6);
+                struct ggml_tensor * grad_cont = ggml_cont(ctx, grad);
+                const size_t offset =
+                    (size_t) lp0 * grad_cont->nb[0] +
+                    (size_t) lp1 * grad_cont->nb[1] +
+                    (size_t) lp2 * grad_cont->nb[2] +
+                    (size_t) lp3 * grad_cont->nb[3];
+                struct ggml_tensor * crop = ggml_view_4d(ctx, grad_cont,
+                    src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
+                    grad_cont->nb[1], grad_cont->nb[2], grad_cont->nb[3], offset);
+                ggml_add_or_set(ctx, cgraph, isrc0, ggml_cont(ctx, crop));
+            }
+        } break;
+        case GGML_OP_CUMSUM: {
+            if (src0_needs_grads) {
+                // PRISM_Q1_LORA_CUMSUM_BACKWARD_V1
+                const int64_t n = grad->ne[0];
+                struct ggml_tensor * order = ggml_argsort(ctx, ggml_arange(ctx, 0.0f, (float) n, 1.0f), GGML_SORT_ORDER_DESC);
+                order = ggml_repeat_4d(ctx, order, n, grad->ne[2], grad->ne[3], 1);
+                struct ggml_tensor * transposed = ggml_cont(ctx, ggml_transpose(ctx, grad));
+                struct ggml_tensor * reversed = ggml_get_rows(ctx, transposed, order);
+                reversed = ggml_cont(ctx, ggml_transpose(ctx, reversed));
+                struct ggml_tensor * accumulated = ggml_cumsum(ctx, reversed);
+                transposed = ggml_cont(ctx, ggml_transpose(ctx, accumulated));
+                reversed = ggml_get_rows(ctx, transposed, order);
+                reversed = ggml_cont(ctx, ggml_transpose(ctx, reversed));
+                ggml_add_or_set(ctx, cgraph, isrc0, reversed);
+            }
+        } break;
+        case GGML_OP_SOLVE_TRI: {
+            // PRISM_Q1_LORA_SOLVE_TRI_BACKWARD_V4
+            // Forward contract is A X = B with lower-triangular A.
+            // For dB solve A^T dB = grad. Convert the upper-triangular
+            // system into a lower-triangular one with J A^T J.
+            const int64_t n = src0->ne[0];
+            struct ggml_tensor * order = ggml_argsort(ctx, ggml_arange(ctx, 0.0f, (float) n, 1.0f), GGML_SORT_ORDER_DESC);
+            order = ggml_repeat_4d(ctx, order, n, src0->ne[2], src0->ne[3], 1);
+
+            struct ggml_tensor * a_t = ggml_cont(ctx, ggml_transpose(ctx, src0));
+            struct ggml_tensor * a_t_reverse_rows = ggml_get_rows(ctx, a_t, order);
+            struct ggml_tensor * a_t_reverse_rows_t = ggml_cont(ctx, ggml_transpose(ctx, a_t_reverse_rows));
+            struct ggml_tensor * a_star_rows = ggml_get_rows(ctx, a_t_reverse_rows_t, order);
+            struct ggml_tensor * a_star = ggml_cont(ctx, ggml_transpose(ctx, a_star_rows));
+
+            struct ggml_tensor * grad_reversed = ggml_get_rows(ctx, ggml_cont(ctx, grad), order);
+            struct ggml_tensor * db_reversed = ggml_solve_tri(ctx, a_star, ggml_cont(ctx, grad_reversed), true, true, false);
+            struct ggml_tensor * db = ggml_get_rows(ctx, db_reversed, order);
+            db = ggml_cont(ctx, db);
+
+            if (src1_needs_grads) {
+                ggml_add_or_set(ctx, cgraph, isrc1, db);
+            }
+            if (src0_needs_grads) {
+                // In GGML storage convention this is dB * X^T.
+                struct ggml_tensor * da = ggml_neg(ctx, ggml_mul_mat(ctx, tensor, db));
+                ggml_add_or_set(ctx, cgraph, isrc0, da);
+            }
+        } break;
+        case GGML_OP_FILL: {
+            // PRISM_Q1_LORA_FILL_BACKWARD_V1
+            // The output is constant with respect to its shape-source tensor.
+        } break;
+        case GGML_OP_DIAG: {
+            if (src0_needs_grads) {
+                // PRISM_Q1_LORA_DIAG_BACKWARD_V1
+                struct ggml_tensor * ones = ggml_fill(ctx, src0, 1.0f);
+                struct ggml_tensor * identity = ggml_diag(ctx, ones);
+                struct ggml_tensor * masked = ggml_mul(ctx, grad, identity);
+                struct ggml_tensor * diagonal = ggml_sum_rows(ctx, masked);
+                diagonal = ggml_cont(ctx, ggml_transpose(ctx, diagonal));
+                ggml_add_or_set(ctx, cgraph, isrc0, ggml_reshape(ctx, diagonal, src0));
+            }
+        } break;
+        case GGML_OP_CONCAT: {
+            // PRISM_Q1_LORA_CONCAT_BACKWARD_V2_MATERIALIZED
+            const int dim = ggml_get_op_params_i32(tensor, 0);
+            GGML_ASSERT(dim >= 0 && dim < GGML_MAX_DIMS);
+            struct ggml_tensor * grad_cont = ggml_cont(ctx, grad);
+            const size_t src1_offset = (size_t) src0->ne[dim] * grad_cont->nb[dim];
+
+            struct ggml_tensor * src0_view = ggml_view_4d(ctx, grad_cont,
+                src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
+                grad_cont->nb[1], grad_cont->nb[2], grad_cont->nb[3], 0);
+            struct ggml_tensor * src1_view = ggml_view_4d(ctx, grad_cont,
+                src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
+                grad_cont->nb[1], grad_cont->nb[2], grad_cont->nb[3], src1_offset);
+
+            struct ggml_tensor * src0_grad = ggml_cont(ctx, src0_view);
+            struct ggml_tensor * src1_grad = ggml_cont(ctx, src1_view);
+            GGML_ASSERT(src0_grad->view_src == NULL);
+            GGML_ASSERT(src1_grad->view_src == NULL);
+            GGML_ASSERT(ggml_are_same_shape(src0_grad, src0));
+            GGML_ASSERT(ggml_are_same_shape(src1_grad, src1));
+
+            if (src0_needs_grads) {
+                ggml_add_or_set(ctx, cgraph, isrc0, src0_grad);
+            }
+            if (src1_needs_grads) {
+                ggml_add_or_set(ctx, cgraph, isrc1, src1_grad);
+            }
+        } break;
         case GGML_OP_RMS_NORM: {
             if (src0_needs_grads) {
                 float eps;
@@ -6760,6 +7004,76 @@ static void ggml_compute_backward(
                 // noop
             }
         } break;
+        case GGML_OP_SET_ROWS: {
+            // PRISM_Q1_LORA_SET_ROWS_BACKWARD_V1
+            //
+            // Forward source ordering:
+            //
+            //   src0 = rows inserted into the destination
+            //   src1 = row indices
+            //   src2 = original destination tensor
+            //
+            // Forward:
+            //
+            //   tensor = src2
+            //   tensor[src1] = src0
+            //
+            // Therefore:
+            //
+            //   dsrc0 = get_rows(grad, src1)
+            //
+            //   dsrc2 = grad with the overwritten rows zeroed
+            //
+            // Row indices are discrete and have no gradient.
+
+            if (src0_needs_grads) {
+                struct ggml_tensor * inserted_rows_grad =
+                    ggml_get_rows(
+                        ctx,
+                        grad,
+                        src1);
+
+                ggml_add_or_set(
+                    ctx,
+                    cgraph,
+                    isrc0,
+                    inserted_rows_grad);
+            }
+
+            GGML_ASSERT(
+                !src1_needs_grads
+                && "SET_ROWS indices are not differentiable");
+
+            if (src2_needs_grads) {
+                // Create zeros with the exact shape of the inserted
+                // rows. No second-order backward graph is requested,
+                // so this dependency is only used to build values.
+                struct ggml_tensor * zero_rows =
+                    ggml_scale(
+                        ctx,
+                        src0,
+                        0.0f);
+
+                // Do not overwrite the incoming gradient tensor.
+                // First duplicate it, then clear rows overwritten by
+                // the forward SET_ROWS operation.
+                struct ggml_tensor * destination_grad =
+                    ggml_set_rows(
+                        ctx,
+                        ggml_dup(
+                            ctx,
+                            grad),
+                        zero_rows,
+                        src1);
+
+                ggml_add_or_set(
+                    ctx,
+                    cgraph,
+                    isrc2,
+                    destination_grad);
+            }
+        } break;
+
         case GGML_OP_DIAG_MASK_INF: {
             if (src0_needs_grads) {
                 /* ggml_diag_mask_inf_impl() shouldn't be here */
@@ -6869,6 +7183,40 @@ static void ggml_compute_backward(
                         ggml_add_or_set(ctx, cgraph, isrc0, ggml_silu_back(ctx, grad, src0));
                     }
                 } break;
+                case GGML_UNARY_OP_SIGMOID: {
+                    // PRISM_Q1_LORA_SIGMOID_BACKWARD_V1
+                    //
+                    // y = sigmoid(x)
+                    //
+                    // dy/dx = y * (1 - y)
+                    //       = sigmoid(x) * sigmoid(-x)
+                    //
+                    // `tensor` is the already-created forward
+                    // sigmoid output.
+                    if (src0_needs_grads) {
+                        struct ggml_tensor * one_minus_y =
+                            ggml_sigmoid(
+                                ctx,
+                                ggml_neg(
+                                    ctx,
+                                    src0));
+
+                        struct ggml_tensor * local_derivative =
+                            ggml_mul(
+                                ctx,
+                                tensor,
+                                one_minus_y);
+
+                        ggml_add_or_set(
+                            ctx,
+                            cgraph,
+                            isrc0,
+                            ggml_mul(
+                                ctx,
+                                grad,
+                                local_derivative));
+                    }
+                } break;
                 case GGML_UNARY_OP_EXP: {
                     if (src0_needs_grads) {
                         ggml_add_or_set(ctx, cgraph, isrc0, ggml_mul(ctx, tensor, grad));
@@ -6922,6 +7270,7 @@ static void ggml_compute_backward(
         } //break;
     }
 
+    // PRISM_Q1_LORA_BACKWARD_SHAPE_AUDIT_V1
     GGML_ASSERT(!src0_needs_grads || ggml_are_same_shape(src0, cgraph->grads[isrc0]));
     GGML_ASSERT(!src1_needs_grads || ggml_are_same_shape(src1, cgraph->grads[isrc1]));
     GGML_ASSERT(!src2_needs_grads || ggml_are_same_shape(src2, cgraph->grads[isrc2]));
@@ -7105,8 +7454,52 @@ void ggml_build_backward_expand(
         }
 
         // inplace operations are currently not supported
-        GGML_ASSERT(!node->view_src || node->op == GGML_OP_CPY || node->op == GGML_OP_VIEW ||
-            node->op == GGML_OP_RESHAPE || node->op == GGML_OP_PERMUTE || node->op == GGML_OP_TRANSPOSE);
+        if (
+            node->view_src
+            && node->op != GGML_OP_CPY
+            && node->op != GGML_OP_VIEW
+            && node->op != GGML_OP_RESHAPE
+            && node->op != GGML_OP_PERMUTE
+            && node->op != GGML_OP_TRANSPOSE
+            && node->op != GGML_OP_SET_ROWS
+        ) {
+            // PRISM_Q1_LORA_INPLACE_DIAGNOSTIC_V2
+            GGML_LOG_ERROR(
+                "PRISM_Q1_LORA_INPLACE_BLOCKER "
+                "node='%s' op=%d "
+                "view_src='%s' "
+                "src0='%s' src1='%s' src2='%s' "
+                "shape=[%lld,%lld,%lld,%lld] "
+                "flags=%d\n",
+                node->name,
+                (int) node->op,
+
+                node->view_src
+                    ? node->view_src->name
+                    : "<null>",
+
+                node->src[0]
+                    ? node->src[0]->name
+                    : "<null>",
+
+                node->src[1]
+                    ? node->src[1]->name
+                    : "<null>",
+
+                node->src[2]
+                    ? node->src[2]->name
+                    : "<null>",
+
+                (long long) node->ne[0],
+                (long long) node->ne[1],
+                (long long) node->ne[2],
+                (long long) node->ne[3],
+                (int) node->flags);
+
+            GGML_ABORT(
+                "native Q1-LoRA backward encountered "
+                "a remaining in-place graph node");
+        }
 
         const size_t ihash = ggml_hash_find(&cgraph->visited_hash_set, node);
         GGML_ASSERT(ihash != GGML_HASHSET_FULL);
diff --git a/include/llama.h b/include/llama.h
index 646ba13..1f4ca1f 100644
--- a/include/llama.h
+++ b/include/llama.h
@@ -1590,6 +1590,33 @@ extern "C" {
             ggml_opt_epoch_callback   callback_train,
             ggml_opt_epoch_callback   callback_eval);
 
+
+// PRISM_STEP10_LLAMA_MASKED_API_BEGIN
+struct llama_opt_masked_stats {
+    double   loss;
+    uint32_t supervised_tokens;
+    uint32_t ubatches;
+    float    gradient_scale;
+};
+
+// Execute one fixed-length sequence through the real model graph. labels are
+// already causally shifted. loss_mask selects the output rows that contribute
+// to cross entropy. When train=true, gradients for params are copied into the
+// caller-provided contiguous F32 buffer; no optimizer update is performed.
+LLAMA_API bool llama_opt_masked_sequence(
+        struct llama_context          * ctx,
+        const llama_token             * tokens,
+        const llama_token             * labels_sparse,
+        const uint8_t                 * loss_mask,
+        uint32_t                        n_tokens,
+        bool                            train,
+        struct ggml_tensor           ** params,
+        size_t                          n_params,
+        float                         * gradients_out,
+        size_t                          gradients_count,
+        struct llama_opt_masked_stats * stats);
+// PRISM_STEP10_LLAMA_MASKED_API_END
+
 #ifdef __cplusplus
 }
 #endif
diff --git a/src/llama-adapter.cpp b/src/llama-adapter.cpp
index 3e0fe66..c3ebc98 100644
--- a/src/llama-adapter.cpp
+++ b/src/llama-adapter.cpp
@@ -6,6 +6,8 @@
 
 #include <map>
 #include <cassert>
+#include <cstdlib>
+#include <cstring>
 #include <sstream>
 #include <stdexcept>
 
@@ -372,6 +374,28 @@ static void llama_adapter_lora_init_impl(llama_model & model, const char * path_
         ggml_tensor * tensor_b = ggml_dup_tensor(dev_ctx, w.b);
         ggml_set_name(tensor_a, w.a->name);
         ggml_set_name(tensor_b, w.b->name);
+
+        // PRISM_Q1_LORA_TRAINABLE_ADAPTER_V2
+        const char * prism_training_env =
+            std::getenv(
+                "PRISM_Q1_LORA_TRAINING");
+
+        const bool prism_q1_lora_training =
+            prism_training_env != nullptr
+            && std::strcmp(
+                prism_training_env,
+                "0") != 0;
+
+        if (prism_q1_lora_training) {
+            ggml_set_param(tensor_a);
+            ggml_set_param(tensor_b);
+
+            LLAMA_LOG_INFO(
+                "PRISM_Q1_LORA_PARAM "
+                "weight=%s a_param=1 b_param=1\n",
+                name.c_str());
+        }
+
         adapter.ab_map[name] = llama_adapter_lora_weight(tensor_a, tensor_b);
     }
 
diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index 5f43818..571ed25 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -22,6 +22,9 @@
 #include <limits>
 #include <numeric>
 #include <stdexcept>
+#include <iomanip>
+#include <iostream>
+#include <cstdlib>
 
 //
 // llama_context
@@ -197,6 +200,36 @@ llama_context::llama_context(
     cparams.fused_gdn_ch = true;
     cparams.auto_fgdn    = true;
 
+    // PRISM_Q1_LORA_UNFUSED_GDN_TRAINING_V1
+    //
+    // The fused GGML_OP_GATED_DELTA_NET is an inference kernel
+    // and currently has no generic ggml_compute_backward case.
+    //
+    // Qwen3.5 already has a graph-composed non-fused path.
+    // Enable it only for explicit native-LoRA SSM validation.
+    if (
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr
+        && getenv("PRISM_Q1_LORA_UNFUSED_GDN") != nullptr
+    ) {
+        cparams.fused_gdn_ar = false;
+        cparams.fused_gdn_ch = false;
+        cparams.auto_fgdn    = false;
+
+        LLAMA_LOG_INFO(
+            "PRISM_Q1_LORA_UNFUSED_GDN enabled=1\n");
+    }
+
+    // PRISM_Q1_LORA_NO_FLASH_TRAINING_V2
+    // Flash-attention backward is not used by the native packed-Q1 trainer.
+    if (
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_NO_KV_CACHE") != nullptr
+    ) {
+        cparams.flash_attn = false;
+        cparams.auto_fa = false;
+        LLAMA_LOG_INFO("PRISM_Q1_LORA_NO_FLASH_TRAINING enabled=1\n");
+    }
+
     // with causal attention, the batch size is limited by the context size
     cparams.n_batch = cparams.causal_attn ? std::min(cparams.n_ctx, params.n_batch) : params.n_batch;
 
@@ -3541,6 +3574,21 @@ static void llama_set_param(struct ggml_tensor * tensor, llama_opt_param_filter
 }
 
 void llama_context::opt_init(struct llama_model * model, struct llama_opt_params lopt_params) {
+
+    // PRISM_STEP10_UNFUSED_GDN_TRAINING_BEGIN
+    // GGML's fused GATED_DELTA_NET node has a forward implementation but the
+    // generic backward builder cannot differentiate it. Keep fused GDN enabled
+    // for normal inference, but rebuild the scheduler with the decomposed graph
+    // before creating the optimizer context.
+    if (cparams.fused_gdn_ar || cparams.fused_gdn_ch || cparams.auto_fgdn) {
+        cparams.fused_gdn_ar = false;
+        cparams.fused_gdn_ch = false;
+        cparams.auto_fgdn    = false;
+        sched_need_reserve   = true;
+        LLAMA_LOG_INFO("%s: PRISM_STEP10_TRAINING_GDN_MODE=UNFUSED\n", __func__);
+        sched_reserve();
+    }
+    // PRISM_STEP10_UNFUSED_GDN_TRAINING_END
     GGML_ASSERT(!opt_ctx);
     model->hparams.n_ctx_train = lopt_params.n_ctx_train > 0 ? lopt_params.n_ctx_train : n_ctx();
     const uint32_t n_batch     = std::min(this->n_batch(),  model->hparams.n_ctx_train);
@@ -3690,6 +3738,546 @@ void llama_context::opt_epoch_iter(
     }
 }
 
+
+// PRISM_STEP10_MASKED_SEQUENCE_IMPL_BEGIN
+bool llama_context::opt_masked_sequence(
+        const llama_token             * tokens,
+        const llama_token             * labels_sparse,
+        const uint8_t                 * loss_mask,
+        uint32_t                        n_tokens,
+        bool                            train,
+        struct ggml_tensor           ** params,
+        size_t                          n_params,
+        float                         * gradients_out,
+        size_t                          gradients_count,
+        struct llama_opt_masked_stats * stats) {
+    if (!opt_ctx || !tokens || !labels_sparse || !loss_mask || !stats) {
+        return false;
+    }
+
+    const uint32_t n_ctx    = llama_model_n_ctx_train(&model);
+    const uint32_t n_batch  = std::min(this->n_batch(),  n_ctx);
+    const uint32_t n_ubatch = std::min(this->n_ubatch(), n_batch);
+    if (n_tokens != n_ctx || n_ctx == 0 || n_ctx % n_batch != 0 || n_batch % n_ubatch != 0) {
+        return false;
+    }
+
+    uint32_t supervised = 0;
+    for (uint32_t i = 0; i < n_ctx; ++i) {
+        supervised += loss_mask[i] != 0;
+    }
+    if (supervised == 0) {
+        return false;
+    }
+
+    size_t expected_gradients = 0;
+    for (size_t i = 0; i < n_params; ++i) {
+        if (!params || !params[i]) {
+            return false;
+        }
+        expected_gradients += ggml_nelements(params[i]);
+    }
+    if (train && (!gradients_out || gradients_count != expected_gradients)) {
+        return false;
+    }
+
+// PRISM_STEP10_V451_OPT_SNAPSHOT_BEGIN
+    const uint32_t ubatches_total =
+        n_ctx / n_ubatch;
+
+    const char * force_opt_env =
+        std::getenv("PRISM_STEP10_FORCE_OPT_BACKWARD");
+
+    const bool force_opt_backward =
+        force_opt_env
+        && force_opt_env[0] == '1'
+        && force_opt_env[1] == '\0';
+
+    const int32_t sentinel_period =
+        force_opt_backward
+            ? 1
+            : static_cast<int32_t>(
+                ubatches_total + 1);
+
+    std::vector<std::vector<float>>
+        opt_probe_snapshots;
+
+    if (train && force_opt_backward) {
+        opt_probe_snapshots.resize(n_params);
+
+        for (size_t i = 0; i < n_params; ++i) {
+            if (!params[i]
+                || params[i]->type != GGML_TYPE_F32) {
+                return false;
+            }
+
+            const size_t count =
+                ggml_nelements(params[i]);
+
+            opt_probe_snapshots[i].resize(
+                count,
+                0.0f);
+
+            ggml_backend_tensor_get(
+                params[i],
+                opt_probe_snapshots[i].data(),
+                0,
+                count*sizeof(float));
+        }
+    }
+
+    ggml_opt_configure_gradient_only(
+        opt_ctx,
+        sentinel_period);
+    // PRISM_STEP10_V451_OPT_SNAPSHOT_END
+
+    ggml_opt_result_t result = ggml_opt_result_init();
+    llama_batch batch = llama_batch_init(n_batch, 0, 1);
+    // PRISM_STEP10_V44_HOST_GRADIENT_CAPTURE_BEGIN
+    std::vector<size_t> param_offsets(n_params + 1, 0);
+    for (size_t i = 0; i < n_params; ++i) {
+        param_offsets[i + 1] =
+            param_offsets[i] + ggml_nelements(params[i]);
+    }
+
+    std::vector<double> direct_host_sum(expected_gradients, 0.0);
+    std::vector<ggml_tensor *> persistent_accumulators(n_params, nullptr);
+    bool accumulators_zeroed = false;
+    uint32_t physical_ubatch_index = 0;
+    // PRISM_STEP10_V44_HOST_GRADIENT_CAPTURE_END
+    bool ok = true;
+
+    memory->clear(true);
+
+    for (uint32_t pos_ctx = 0; ok && pos_ctx < n_ctx; pos_ctx += n_batch) {
+        batch.n_tokens = n_batch;
+        for (uint32_t pos_batch = 0; pos_batch < n_batch; ++pos_batch) {
+            batch.token   [pos_batch]    = tokens[pos_ctx + pos_batch];
+            batch.pos     [pos_batch]    = pos_ctx + pos_batch;
+            batch.n_seq_id[pos_batch]    = 1;
+            batch.seq_id  [pos_batch][0] = 0;
+            batch.logits  [pos_batch]    = true;
+        }
+
+        if (!balloc->init(batch, model.vocab, nullptr, model.hparams.n_embd_inp(),
+                          cparams.kv_unified ? LLAMA_MAX_SEQ : cparams.n_seq_max, true)) {
+            ok = false;
+            break;
+        }
+
+        const uint32_t n_tokens_all = balloc->get_n_tokens();
+        n_queued_tokens += n_tokens_all;
+        embd_seq.clear();
+
+        auto mctx = memory->init_batch(*balloc, cparams.n_ubatch, true);
+        if (!mctx || mctx->get_status() != LLAMA_MEMORY_STATUS_SUCCESS) {
+            ok = false;
+            break;
+        }
+
+        if (output_reserve(n_tokens_all) < n_tokens_all) {
+            ok = false;
+            break;
+        }
+
+        uint32_t pos_batch = 0;
+        do {
+            const auto & ubatch = mctx->get_ubatch();
+            n_outputs = ubatch.n_tokens;
+
+            if (!mctx->apply()) {
+                ok = false;
+                break;
+            }
+
+            auto * res = gf_res_prev.get();
+            const auto gparams = graph_params(
+                res, ubatch, mctx.get(), ctx_type_to_graph_type(cparams.ctx_type));
+            res->reset();
+            auto * gf = model.build_graph(gparams);
+
+            struct ggml_context * ctx_compute_opt = nullptr;
+            {
+                const size_t size_gf = ggml_graph_size(gf);
+                const size_t size_meta =
+                    4*size_gf*ggml_tensor_overhead() +
+                    2*ggml_graph_overhead_custom(size_gf, true);
+                struct ggml_init_params init_params = {
+                    /*.mem_size   =*/ size_meta,
+                    /*.mem_buffer =*/ nullptr,
+                    /*.no_alloc   =*/ true,
+                };
+                ctx_compute_opt = ggml_init(init_params);
+            }
+
+            ggml_opt_prepare_alloc(
+                opt_ctx, ctx_compute_opt, gf, res->get_inp_tokens(), res->get_logits());
+            ggml_opt_alloc(opt_ctx, train);
+
+            // PRISM_STEP10_V44_ACTIVE_GRADIENT_LOOKUP_BEGIN
+            std::vector<ggml_tensor *> direct_tensors(n_params, nullptr);
+
+            if (train) {
+                if (!accumulators_zeroed) {
+                    ggml_opt_zero_grad_accumulators(opt_ctx);
+                    accumulators_zeroed = true;
+                }
+
+                for (size_t i = 0; i < n_params; ++i) {
+                    direct_tensors[i] =
+                        ggml_opt_grad_from_active_graph(opt_ctx, params[i]);
+
+                    ggml_tensor * accumulator =
+                        ggml_opt_grad_acc_from_active_graph(opt_ctx, params[i]);
+
+                    if (accumulator) {
+                        persistent_accumulators[i] = accumulator;
+                    }
+
+                    if (!direct_tensors[i] && !accumulator) {
+                        std::cerr
+                            << "PRISM_STEP10_GRAD_LOOKUP_FAILED "
+                            << "param_index=" << i
+                            << " name='" << params[i]->name << "'\n";
+                        ok = false;
+                        break;
+                    }
+                }
+            }
+            // PRISM_STEP10_V44_ACTIVE_GRADIENT_LOOKUP_END
+            if (!ok) {
+                ggml_free(ctx_compute_opt);
+                break;
+            }
+
+            res->set_inputs(&ubatch);
+            struct ggml_tensor * labels = ggml_opt_labels(opt_ctx);
+            if (!labels || labels->ne[1] != n_ubatch) {
+                ggml_free(ctx_compute_opt);
+                ok = false;
+                break;
+            }
+            ggml_set_zero(labels);
+            const float one = 1.0f;
+            for (uint32_t pos_ubatch = 0; pos_ubatch < n_ubatch; ++pos_ubatch) {
+                const uint32_t index = pos_ctx + pos_batch + pos_ubatch;
+                if (!loss_mask[index]) {
+                    continue;
+                }
+                const llama_token label = labels_sparse[index];
+                if (label < 0 || label >= labels->ne[0]) {
+                    ggml_free(ctx_compute_opt);
+                    ok = false;
+                    break;
+                }
+                const size_t offset =
+                    (static_cast<size_t>(pos_ubatch)*labels->ne[0] + label)*sizeof(float);
+                ggml_backend_tensor_set(labels, &one, offset, sizeof(float));
+            }
+            if (!ok) {
+                break;
+            }
+
+            ggml_opt_eval(opt_ctx, result);
+
+            // PRISM_STEP10_V44_IMMEDIATE_DIRECT_COPY_BEGIN
+            if (train) {
+                for (size_t i = 0; i < n_params; ++i) {
+                    ggml_tensor * direct = direct_tensors[i];
+                    if (!direct || direct->type != GGML_TYPE_F32) {
+                        continue;
+                    }
+
+                    const size_t count = ggml_nelements(params[i]);
+                    std::vector<float> values(count, 0.0f);
+                    ggml_backend_tensor_get(
+                        direct,
+                        values.data(),
+                        0,
+                        count*sizeof(float));
+
+                    size_t nonzero = 0;
+                    double max_abs = 0.0;
+                    double l2_squared = 0.0;
+
+                    for (size_t j = 0; j < count; ++j) {
+                        const float value = values[j];
+                        if (!std::isfinite(value)) {
+                            ok = false;
+                            break;
+                        }
+
+                        direct_host_sum[param_offsets[i] + j] +=
+                            static_cast<double>(value);
+
+                        if (value != 0.0f) {
+                            ++nonzero;
+                        }
+
+                        const double value_d =
+                            static_cast<double>(value);
+
+                        max_abs = std::max(
+                            max_abs,
+                            std::abs(value_d));
+
+                        l2_squared += value_d*value_d;
+                    }
+
+                    std::cerr
+                        << std::setprecision(17)
+                        << "PRISM_STEP10_UBATCH_DIRECT_GRAD "
+                        << "ubatch=" << physical_ubatch_index
+                        << " param_index=" << i
+                        << " name='" << params[i]->name << "'"
+                        << " nonzero=" << nonzero
+                        << " max_abs=" << max_abs
+                        << " l2=" << std::sqrt(l2_squared)
+                        << "\n";
+
+                    if (!ok) {
+                        break;
+                    }
+                }
+            }
+
+            ++physical_ubatch_index;
+            // PRISM_STEP10_V44_IMMEDIATE_DIRECT_COPY_END
+
+// PRISM_STEP10_V451_OPT_RESTORE_BEGIN
+            if (train && force_opt_backward) {
+                for (size_t i = 0; i < n_params; ++i) {
+                    const size_t count =
+                        ggml_nelements(params[i]);
+
+                    ggml_backend_tensor_set(
+                        params[i],
+                        opt_probe_snapshots[i].data(),
+                        0,
+                        count*sizeof(float));
+                }
+
+                std::cerr
+                    << "PRISM_STEP10_OPT_PARAMS_RESTORED "
+                    << "ubatch="
+                    << physical_ubatch_index
+                    << "\n";
+            }
+            // PRISM_STEP10_V451_OPT_RESTORE_END
+
+            ggml_free(ctx_compute_opt);
+            pos_batch += ubatch.n_tokens;
+        } while (mctx->next());
+    }
+
+    // PRISM_STEP10_V43_RESULT_ACCOUNTING_BEGIN
+    int64_t result_ndata = 0;
+    ggml_opt_result_ndata(result, &result_ndata);
+    std::cerr << "PRISM_STEP10_RESULT_NDATA=" << result_ndata << "\n";
+
+    double loss = 0.0;
+    double loss_uncertainty = NAN;
+
+    if (ok && result_ndata <= 0) {
+        std::cerr << "PRISM_STEP10_EMPTY_OPT_RESULT=1\n";
+        ok = false;
+    }
+
+    if (ok) {
+        ggml_opt_result_loss(result, &loss, &loss_uncertainty);
+        std::cerr << std::setprecision(17)
+                  << "PRISM_STEP10_RAW_LOSS=" << loss << "\n"
+                  << "PRISM_STEP10_LOSS_UNCERTAINTY=" << loss_uncertainty << "\n";
+        loss *= static_cast<double>(n_ctx) / static_cast<double>(supervised);
+    }
+    // PRISM_STEP10_V43_RESULT_ACCOUNTING_END
+
+    // The backward graph used 1/sentinel_period for each ubatch. Convert the
+    // accumulated gradient to an average over supervised output rows.
+    const float gradient_scale =
+        static_cast<float>(sentinel_period) /
+        static_cast<float>(ubatches_total) *
+        static_cast<float>(n_ctx) /
+        static_cast<float>(supervised);
+
+    // PRISM_STEP10_V44_GRADIENT_SOURCE_ARBITRATION_BEGIN
+    if (ok && train) {
+        std::vector<float> accumulator_values(expected_gradients, 0.0f);
+
+        size_t accumulator_nonzero = 0;
+        size_t direct_nonzero = 0;
+        double accumulator_l2_squared = 0.0;
+        double direct_l2_squared = 0.0;
+
+        for (size_t i = 0; i < n_params; ++i) {
+            const size_t count = ggml_nelements(params[i]);
+
+            if (persistent_accumulators[i] &&
+                persistent_accumulators[i]->type == GGML_TYPE_F32) {
+                std::vector<float> values(count, 0.0f);
+                ggml_backend_tensor_get(
+                    persistent_accumulators[i],
+                    values.data(),
+                    0,
+                    count*sizeof(float));
+
+                for (size_t j = 0; j < count; ++j) {
+                    const float value = values[j];
+                    if (!std::isfinite(value)) {
+                        ok = false;
+                        break;
+                    }
+
+                    accumulator_values[param_offsets[i] + j] = value;
+
+                    const double value_d =
+                        static_cast<double>(value);
+
+                    if (value != 0.0f) {
+                        ++accumulator_nonzero;
+                    }
+
+                    accumulator_l2_squared += value_d*value_d;
+                }
+            }
+
+            for (size_t j = 0; j < count; ++j) {
+                const double value =
+                    direct_host_sum[param_offsets[i] + j];
+
+                if (value != 0.0) {
+                    ++direct_nonzero;
+                }
+
+                direct_l2_squared += value*value;
+            }
+
+            if (!ok) {
+                break;
+            }
+        }
+
+        const double accumulator_l2 =
+            std::sqrt(accumulator_l2_squared);
+
+        const double direct_l2 =
+            std::sqrt(direct_l2_squared);
+
+        const bool use_accumulator =
+            accumulator_nonzero > 0 &&
+            std::isfinite(accumulator_l2) &&
+            accumulator_l2 > 0.0;
+
+        const bool use_direct =
+            !use_accumulator &&
+            direct_nonzero > 0 &&
+            std::isfinite(direct_l2) &&
+            direct_l2 > 0.0;
+
+        const char * source_name =
+            use_accumulator
+                ? "ACCUMULATOR"
+                : use_direct
+                    ? "DIRECT_HOST_SUM"
+                    : "NONE";
+
+        std::cerr
+            << std::setprecision(17)
+            << "PRISM_STEP10_GRADIENT_SOURCE="
+            << source_name << "\n"
+            << "PRISM_STEP10_ACCUMULATOR_NONZERO="
+            << accumulator_nonzero << "\n"
+            << "PRISM_STEP10_ACCUMULATOR_L2="
+            << accumulator_l2 << "\n"
+            << "PRISM_STEP10_DIRECT_SUM_NONZERO="
+            << direct_nonzero << "\n"
+            << "PRISM_STEP10_DIRECT_SUM_L2="
+            << direct_l2 << "\n";
+
+        if (!use_accumulator && !use_direct) {
+            std::cerr
+                << "PRISM_STEP10_ZERO_USEFUL_GRADIENT=1\n";
+            ok = false;
+        }
+
+        size_t output_offset = 0;
+
+        for (size_t i = 0; ok && i < n_params; ++i) {
+            const size_t count = ggml_nelements(params[i]);
+
+            size_t param_nonzero = 0;
+            double param_max_abs = 0.0;
+            double param_l2_squared = 0.0;
+
+            for (size_t j = 0; j < count; ++j) {
+                const size_t index = param_offsets[i] + j;
+
+                const double raw_value =
+                    use_accumulator
+                        ? static_cast<double>(accumulator_values[index])
+                        : direct_host_sum[index];
+
+                const double scaled_value =
+                    raw_value*static_cast<double>(gradient_scale);
+
+                if (!std::isfinite(scaled_value)) {
+                    ok = false;
+                    break;
+                }
+
+                gradients_out[output_offset + j] =
+                    static_cast<float>(scaled_value);
+
+                if (scaled_value != 0.0) {
+                    ++param_nonzero;
+                }
+
+                param_max_abs = std::max(
+                    param_max_abs,
+                    std::abs(scaled_value));
+
+                param_l2_squared +=
+                    scaled_value*scaled_value;
+            }
+
+            std::cerr
+                << std::setprecision(17)
+                << "PRISM_STEP10_PARAM_GRAD "
+                << "param_index=" << i
+                << " name='" << params[i]->name << "'"
+                << " source=" << source_name
+                << " nonzero=" << param_nonzero
+                << " max_abs=" << param_max_abs
+                << " l2=" << std::sqrt(param_l2_squared)
+                << "\n";
+
+            if (param_nonzero == 0 ||
+                !std::isfinite(param_max_abs) ||
+                param_max_abs <= 0.0) {
+                std::cerr
+                    << "PRISM_STEP10_ZERO_PARAM_GRAD "
+                    << "param_index=" << i
+                    << " name='" << params[i]->name << "'\n";
+                ok = false;
+            }
+
+            output_offset += count;
+        }
+    }
+    // PRISM_STEP10_V44_GRADIENT_SOURCE_ARBITRATION_END
+
+    stats->loss              = loss;
+    stats->supervised_tokens = supervised;
+    stats->ubatches          = ubatches_total;
+    stats->gradient_scale    = gradient_scale;
+
+    ggml_opt_reset_gradient_cycle(opt_ctx);
+    llama_batch_free(batch);
+    ggml_opt_result_free(result);
+    return ok && std::isfinite(loss);
+}
+// PRISM_STEP10_MASKED_SEQUENCE_IMPL_END
+
 void llama_context::opt_epoch(
         ggml_opt_dataset_t        dataset,
         ggml_opt_result_t         result_train,
@@ -4453,6 +5041,34 @@ void llama_opt_epoch(
         callback_eval);
 }
 
+
+// PRISM_STEP10_MASKED_SEQUENCE_C_API_BEGIN
+bool llama_opt_masked_sequence(
+        struct llama_context          * ctx,
+        const llama_token             * tokens,
+        const llama_token             * labels_sparse,
+        const uint8_t                 * loss_mask,
+        uint32_t                        n_tokens,
+        bool                            train,
+        struct ggml_tensor           ** params,
+        size_t                          n_params,
+        float                         * gradients_out,
+        size_t                          gradients_count,
+        struct llama_opt_masked_stats * stats) {
+    return ctx && ctx->opt_masked_sequence(
+        tokens,
+        labels_sparse,
+        loss_mask,
+        n_tokens,
+        train,
+        params,
+        n_params,
+        gradients_out,
+        gradients_count,
+        stats);
+}
+// PRISM_STEP10_MASKED_SEQUENCE_C_API_END
+
 //
 // ext
 //
diff --git a/src/llama-context.h b/src/llama-context.h
index de93d0c..1856e71 100644
--- a/src/llama-context.h
+++ b/src/llama-context.h
@@ -218,6 +218,21 @@ struct llama_context {
             ggml_opt_epoch_callback callback_train,
             ggml_opt_epoch_callback callback_eval);
 
+
+    // PRISM_STEP10_LLAMA_CONTEXT_API_BEGIN
+    bool opt_masked_sequence(
+            const llama_token             * tokens,
+            const llama_token             * labels_sparse,
+            const uint8_t                 * loss_mask,
+            uint32_t                        n_tokens,
+            bool                            train,
+            struct ggml_tensor           ** params,
+            size_t                          n_params,
+            float                         * gradients_out,
+            size_t                          gradients_count,
+            struct llama_opt_masked_stats * stats);
+    // PRISM_STEP10_LLAMA_CONTEXT_API_END
+
     void opt_epoch_iter(
             ggml_opt_dataset_t               dataset,
             ggml_opt_result_t                result,
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index b3d9863..685da71 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -15,6 +15,8 @@
 #include <cassert>
 #include <cmath>
 #include <cstring>
+#include <cstdlib>
+#include <cstdio>
 #include <numeric>
 #include <sstream>
 #include <unordered_set>
@@ -692,31 +694,51 @@ void llm_graph_input_attn_cross::set_input(const llama_ubatch * ubatch) {
 }
 
 void llm_graph_input_mem_hybrid::set_input(const llama_ubatch * ubatch) {
-    mctx->get_attn()->set_input_k_idxs(inp_attn->self_k_idxs, ubatch);
-    mctx->get_attn()->set_input_v_idxs(inp_attn->self_v_idxs, ubatch);
+    const bool prism_no_kv =
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_NO_KV_CACHE") != nullptr;
 
-    mctx->get_attn()->set_input_kq_mask(inp_attn->self_kq_mask, ubatch, cparams.causal_attn);
+    const bool k_allocated = inp_attn->self_k_idxs && inp_attn->self_k_idxs->buffer;
+    const bool v_allocated = inp_attn->self_v_idxs && inp_attn->self_v_idxs->buffer;
+    const bool mask_allocated = inp_attn->self_kq_mask && inp_attn->self_kq_mask->buffer;
+
+    // PRISM_Q1_LORA_HYBRID_NO_KV_INPUT_SKIP_V2
+    if (!prism_no_kv || k_allocated) {
+        mctx->get_attn()->set_input_k_idxs(inp_attn->self_k_idxs, ubatch);
+    }
+    if (!prism_no_kv || v_allocated) {
+        mctx->get_attn()->set_input_v_idxs(inp_attn->self_v_idxs, ubatch);
+    }
+    if (!prism_no_kv || mask_allocated) {
+        mctx->get_attn()->set_input_kq_mask(inp_attn->self_kq_mask, ubatch, cparams.causal_attn);
+    }
+
+    if (prism_no_kv && (!k_allocated || !v_allocated || !mask_allocated)) {
+        static int prism_skip_count = 0;
+        if (prism_skip_count < 32) {
+            fprintf(stderr,
+                "PRISM_Q1_LORA_SKIP_NO_KV_INPUTS class=hybrid count=%d "
+                "k_allocated=%d v_allocated=%d mask_allocated=%d\n",
+                prism_skip_count + 1, (int) k_allocated, (int) v_allocated, (int) mask_allocated);
+        }
+        prism_skip_count++;
+    }
 
     if (inp_attn->self_k_rot) {
         mctx->get_attn()->set_input_k_rot(inp_attn->self_k_rot);
     }
-
     if (inp_attn->self_v_rot) {
         mctx->get_attn()->set_input_v_rot(inp_attn->self_v_rot);
     }
 
     const int64_t n_rs = mctx->get_recr()->get_n_rs();
-
     if (inp_rs->s_copy) {
         GGML_ASSERT(ggml_backend_buffer_is_host(inp_rs->s_copy->buffer));
         int32_t * data = (int32_t *) inp_rs->s_copy->data;
-
-        // assuming copy destinations ALWAYS happen ONLY on the cells between head and head+n
         for (uint32_t i = 0; i < n_rs; ++i) {
             data[i] = mctx->get_recr()->s_copy(i);
         }
     }
-
     if (inp_rs->s_write_rows) {
         mctx->get_recr()->set_input_s_write_rows(inp_rs->s_write_rows, inp_rs->s_write_rows_conv);
     }
@@ -724,29 +746,30 @@ void llm_graph_input_mem_hybrid::set_input(const llama_ubatch * ubatch) {
 
 bool llm_graph_input_mem_hybrid::can_reuse(const llm_graph_params & params) {
     const auto * mctx = static_cast<const llama_memory_hybrid_context *>(params.mctx);
-
     this->mctx = mctx;
-
     bool res = true;
 
-    res &= inp_attn->self_k_idxs->ne[0] == params.ubatch.n_tokens;
-  //res &= inp_attn->self_v_idxs->ne[0] == params.ubatch.n_tokens; // TODO: need to move this to the unified cache and check there
+    const bool prism_no_kv =
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_NO_KV_CACHE") != nullptr;
 
-    res &= can_reuse_kq_mask(inp_attn->self_kq_mask, mctx->get_attn(), params.ubatch, params.cparams);
+    // PRISM_Q1_LORA_HYBRID_NO_KV_REUSE_V2
+    if (!prism_no_kv || (inp_attn->self_k_idxs && inp_attn->self_k_idxs->buffer)) {
+        res &= inp_attn->self_k_idxs->ne[0] == params.ubatch.n_tokens;
+    }
+    if (!prism_no_kv || (inp_attn->self_kq_mask && inp_attn->self_kq_mask->buffer)) {
+        res &= can_reuse_kq_mask(inp_attn->self_kq_mask, mctx->get_attn(), params.ubatch, params.cparams);
+    }
 
     res &= inp_rs->s_copy->ne[0] == mctx->get_recr()->get_n_rs();
-
     res &= inp_rs->s_copy_main->ne[0]  == params.ubatch.n_seqs;
     res &= inp_rs->s_copy_extra->ne[0] == mctx->get_recr()->get_n_rs() - params.ubatch.n_seqs;
-
     if (inp_rs->s_write_rows) {
         res &= inp_rs->s_write_rows->ne[0] == rs_n_write_rows(mctx->get_recr(), params.ubatch);
         res &= !inp_rs->s_write_rows_conv || inp_rs->s_write_rows_conv->ne[0] == rs_n_write_rows(mctx->get_recr(), params.ubatch);
     }
-
     res &= inp_rs->head == mctx->get_recr()->get_head();
     res &= inp_rs->rs_z == mctx->get_recr()->get_rs_z();
-
     return res;
 }
 
@@ -754,22 +777,28 @@ bool llm_graph_input_mem_hybrid::can_reuse(const llm_graph_params & params) {
 // Instead of creating a hybrid input, the graph can simply create 2 separate inputs.
 // Refactoring is required in the future.
 void llm_graph_input_mem_hybrid_k::set_input(const llama_ubatch * ubatch) {
-    mctx->get_attn()->set_input_k_idxs(inp_attn->self_k_idxs, ubatch);
+    const bool prism_no_kv =
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_NO_KV_CACHE") != nullptr;
+    const bool k_allocated = inp_attn->self_k_idxs && inp_attn->self_k_idxs->buffer;
+    const bool mask_allocated = inp_attn->self_kq_mask && inp_attn->self_kq_mask->buffer;
 
-    mctx->get_attn()->set_input_kq_mask(inp_attn->self_kq_mask, ubatch, cparams.causal_attn);
+    // PRISM_Q1_LORA_HYBRID_K_NO_KV_INPUT_SKIP_V2
+    if (!prism_no_kv || k_allocated) {
+        mctx->get_attn()->set_input_k_idxs(inp_attn->self_k_idxs, ubatch);
+    }
+    if (!prism_no_kv || mask_allocated) {
+        mctx->get_attn()->set_input_kq_mask(inp_attn->self_kq_mask, ubatch, cparams.causal_attn);
+    }
 
     const int64_t n_rs = mctx->get_recr()->get_n_rs();
-
     if (inp_rs->s_copy) {
         GGML_ASSERT(ggml_backend_buffer_is_host(inp_rs->s_copy->buffer));
         int32_t * data = (int32_t *) inp_rs->s_copy->data;
-
-        // assuming copy destinations ALWAYS happen ONLY on the cells between head and head+n
         for (uint32_t i = 0; i < n_rs; ++i) {
             data[i] = mctx->get_recr()->s_copy(i);
         }
     }
-
     if (inp_rs->s_write_rows) {
         mctx->get_recr()->set_input_s_write_rows(inp_rs->s_write_rows, inp_rs->s_write_rows_conv);
     }
@@ -777,28 +806,30 @@ void llm_graph_input_mem_hybrid_k::set_input(const llama_ubatch * ubatch) {
 
 bool llm_graph_input_mem_hybrid_k::can_reuse(const llm_graph_params & params) {
     const auto * mctx = static_cast<const llama_memory_hybrid_context *>(params.mctx);
-
     this->mctx = mctx;
-
     bool res = true;
 
-    res &= inp_attn->self_k_idxs->ne[0] == params.ubatch.n_tokens;
+    const bool prism_no_kv =
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_NO_KV_CACHE") != nullptr;
 
-    res &= can_reuse_kq_mask(inp_attn->self_kq_mask, mctx->get_attn(), params.ubatch, params.cparams);
+    // PRISM_Q1_LORA_HYBRID_K_NO_KV_REUSE_V1
+    if (!prism_no_kv || (inp_attn->self_k_idxs && inp_attn->self_k_idxs->buffer)) {
+        res &= inp_attn->self_k_idxs->ne[0] == params.ubatch.n_tokens;
+    }
+    if (!prism_no_kv || (inp_attn->self_kq_mask && inp_attn->self_kq_mask->buffer)) {
+        res &= can_reuse_kq_mask(inp_attn->self_kq_mask, mctx->get_attn(), params.ubatch, params.cparams);
+    }
 
     res &= inp_rs->s_copy->ne[0] == mctx->get_recr()->get_n_rs();
-
     res &= inp_rs->s_copy_main->ne[0]  == params.ubatch.n_seqs;
     res &= inp_rs->s_copy_extra->ne[0] == mctx->get_recr()->get_n_rs() - params.ubatch.n_seqs;
-
     if (inp_rs->s_write_rows) {
         res &= inp_rs->s_write_rows->ne[0] == rs_n_write_rows(mctx->get_recr(), params.ubatch);
         res &= !inp_rs->s_write_rows_conv || inp_rs->s_write_rows_conv->ne[0] == rs_n_write_rows(mctx->get_recr(), params.ubatch);
     }
-
     res &= inp_rs->head == mctx->get_recr()->get_head();
     res &= inp_rs->rs_z == mctx->get_recr()->get_rs_z();
-
     return res;
 }
 
@@ -1132,28 +1163,99 @@ ggml_tensor * llm_graph_context::build_lora_mm(
           ggml_tensor * w,
           ggml_tensor * cur,
           ggml_tensor * w_s) const {
-    ggml_tensor * res = ggml_mul_mat(ctx0, w, cur);
+    // PRISM_Q1_LORA_REAL_LOADER_V2
+    //
+    // Normal inference remains unchanged unless the explicit
+    // training environment flag is enabled.
+    const char * prism_training_env =
+        std::getenv(
+            "PRISM_Q1_LORA_TRAINING");
+
+    const bool prism_q1_lora_training =
+        prism_training_env != nullptr
+        && std::strcmp(
+            prism_training_env,
+            "0") != 0;
+
+    ggml_tensor * res =
+        ggml_mul_mat(
+            ctx0,
+            w,
+            cur);
+
+    // Q1 inference normally uses the fast MMQ route.
+    // Native LoRA training requires the exact packed-Q1 × F32
+    // activation route so backward-X matches the linear forward.
+    if (
+        prism_q1_lora_training
+        && w->type == GGML_TYPE_Q1_0
+    ) {
+        ggml_mul_mat_set_prec(
+            res,
+            GGML_PREC_F32);
+    }
 
     for (const auto & lora : *loras) {
-        llama_adapter_lora_weight * lw = lora.first->get_weight(w);
+        llama_adapter_lora_weight * lw =
+            lora.first->get_weight(w);
+
         if (lw == nullptr) {
             continue;
         }
 
-        const float adapter_scale = lora.second;
-        const float scale = lw->get_scale(lora.first->alpha, adapter_scale);
-
-        ggml_tensor * ab_cur = ggml_mul_mat(
-                ctx0, lw->b,
-                ggml_mul_mat(ctx0, lw->a, cur)
-                );
-
-        ab_cur = ggml_scale(ctx0, ab_cur, scale);
-        res = ggml_add(ctx0, res, ab_cur);
+        const float adapter_scale =
+            lora.second;
+
+        const float scale =
+            lw->get_scale(
+                lora.first->alpha,
+                adapter_scale);
+
+        if (prism_q1_lora_training) {
+            LLAMA_LOG_INFO(
+                "PRISM_Q1_LORA_GRAPH "
+                "weight=%s exact_q1=%d "
+                "a_param=%d b_param=%d\n",
+                w->name,
+                w->type == GGML_TYPE_Q1_0 ? 1 : 0,
+                (
+                    lw->a->flags
+                    & GGML_TENSOR_FLAG_PARAM
+                ) ? 1 : 0,
+                (
+                    lw->b->flags
+                    & GGML_TENSOR_FLAG_PARAM
+                ) ? 1 : 0);
+        }
+
+        ggml_tensor * ab_cur =
+            ggml_mul_mat(
+                ctx0,
+                lw->b,
+                ggml_mul_mat(
+                    ctx0,
+                    lw->a,
+                    cur));
+
+        ab_cur =
+            ggml_scale(
+                ctx0,
+                ab_cur,
+                scale);
+
+        res =
+            ggml_add(
+                ctx0,
+                res,
+                ab_cur);
     }
 
     if (w_s) {
-        res = ggml_mul(ctx0, res, w_s);
+        res =
+            ggml_mul(
+                ctx0,
+                res,
+                w_s);
     }
 
     return res;
diff --git a/src/models/delta-net-base.cpp b/src/models/delta-net-base.cpp
index e65f55b..7e23d7a 100644
--- a/src/models/delta-net-base.cpp
+++ b/src/models/delta-net-base.cpp
@@ -152,10 +152,10 @@ std::pair<ggml_tensor *, ggml_tensor *> llm_build_delta_net_base::build_delta_ne
     attn = ggml_tri(ctx0, kb, GGML_TRI_TYPE_LOWER);
     cb(attn, "attn", il);
 
-    ggml_tensor * identity;
-    identity = ggml_view_1d(ctx0, attn, CS, 0);
-    identity = ggml_fill   (ctx0, identity, 1.0f);
-    identity = ggml_diag   (ctx0, identity);
+    // PRISM_Q1_LORA_INDEPENDENT_DELTA_IDENTITY_V1
+    ggml_tensor * identity = ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, CS);
+    identity = ggml_fill(ctx0, identity, 1.0f);
+    identity = ggml_diag(ctx0, identity);
 
     ggml_tensor * lhs = ggml_add(ctx0, attn, identity);
     cb(lhs, "dnet_add_ch_lhs", il);
diff --git a/src/models/qwen35.cpp b/src/models/qwen35.cpp
index 1d16143..f227f82 100644
--- a/src/models/qwen35.cpp
+++ b/src/models/qwen35.cpp
@@ -1,6 +1,42 @@
 #include "models.h"
 #include "llama-memory-recurrent.h"
 
+#include <cstdio>
+#include <cstdlib>
+
+// PRISM_Q1_LORA_GENERIC_SSM_CONV_V2
+static ggml_tensor * prism_q1_lora_generic_ssm_conv(
+        ggml_context * ctx,
+        ggml_tensor  * input,
+        ggml_tensor  * kernel) {
+    GGML_ASSERT(input->type == GGML_TYPE_F32);
+    GGML_ASSERT(kernel->type == GGML_TYPE_F32);
+    GGML_ASSERT(input->ne[1] == kernel->ne[1]);
+    GGML_ASSERT(input->ne[0] >= kernel->ne[0]);
+
+    const int64_t kernel_size = kernel->ne[0];
+    const int64_t channels = input->ne[1];
+    const int64_t tokens = input->ne[0] - kernel_size + 1;
+    const int64_t sequences = input->ne[2];
+
+    ggml_tensor * result = nullptr;
+    for (int64_t tap = 0; tap < kernel_size; ++tap) {
+        ggml_tensor * window = ggml_view_3d(ctx, input,
+            tokens, channels, sequences,
+            input->nb[1], input->nb[2], (size_t) tap*input->nb[0]);
+        window = ggml_cont(ctx, ggml_transpose(ctx, window));
+
+        ggml_tensor * tap_weight = ggml_view_2d(ctx, kernel,
+            1, channels, kernel->nb[1], (size_t) tap*kernel->nb[0]);
+        tap_weight = ggml_cont(ctx, ggml_transpose(ctx, tap_weight));
+        tap_weight = ggml_repeat(ctx, tap_weight, window);
+
+        ggml_tensor * term = ggml_mul(ctx, window, tap_weight);
+        result = result ? ggml_add(ctx, result, term) : term;
+    }
+    return result;
+}
+
 void llama_model_qwen35::load_arch_hparams(llama_model_loader & ml) {
     ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS,       hparams.f_norm_rms_eps);
     ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS,    hparams.rope_sections, 4, true);
@@ -393,9 +429,25 @@ ggml_tensor * llama_model_qwen35::graph::build_layer_attn(
     // Attention computation
     const float kq_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale;
 
-    cur = build_attn(inp,
-                nullptr, nullptr, nullptr,
-                Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
+    if (
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_NO_KV_CACHE") != nullptr
+    ) {
+        // PRISM_Q1_LORA_NO_KV_ATTENTION_V2
+        static int prism_no_kv_attention_count = 0;
+        if (prism_no_kv_attention_count < 128) {
+            fprintf(stderr, "PRISM_Q1_LORA_NO_KV_ATTENTION enabled=1 layer=%d\n", il);
+        }
+        prism_no_kv_attention_count++;
+        llm_graph_input_attn_no_cache * inp_no_cache = build_attn_inp_no_cache();
+        cur = build_attn(inp_no_cache,
+                    nullptr, nullptr, nullptr,
+                    Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
+    } else {
+        cur = build_attn(inp,
+                    nullptr, nullptr, nullptr,
+                    Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
+    }
     cb(cur, "attn_pregate", il);
 
     ggml_tensor * gate_sigmoid = ggml_sigmoid(ctx0, gate);
@@ -498,7 +550,20 @@ ggml_tensor * llama_model_qwen35::graph::build_layer_attn_linear(
         cb(state, "state_predelta", il);
     }
 
-    ggml_tensor * conv_output_proper = ggml_ssm_conv(ctx0, conv_input, conv_kernel);
+    ggml_tensor * conv_output_proper;
+    if (
+        getenv("PRISM_Q1_LORA_TRAINING") != nullptr &&
+        getenv("PRISM_Q1_LORA_TRAINING_GENERIC_SSM_CONV") != nullptr
+    ) {
+        static int prism_generic_conv_count = 0;
+        if (prism_generic_conv_count < 128) {
+            fprintf(stderr, "PRISM_Q1_LORA_GENERIC_SSM_CONV enabled=1 layer=%d\n", il);
+        }
+        prism_generic_conv_count++;
+        conv_output_proper = prism_q1_lora_generic_ssm_conv(ctx0, conv_input, conv_kernel);
+    } else {
+        conv_output_proper = ggml_ssm_conv(ctx0, conv_input, conv_kernel);
+    }
     cb(conv_output_proper, "conv_output_raw", il);
 
     ggml_tensor * conv_output_silu = ggml_silu(ctx0, conv_output_proper);
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 4194388..45fd390 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -367,3 +367,42 @@ llama_build(test-dspark-loop.cpp)
 # file's header comment). Not wired into `ctest`: needs real multi-GB GGUFs
 # passed on the command line; run manually on a GPU host.
 llama_build(test-dspark-real-eval.cpp)
+
+# PRISM_Q1_LORA_OPTIMIZER_TEST_V1
+llama_build(test-q1-lora-opt.cpp)
+# PRISM_Q1_LORA_INTERNAL_INCLUDE_V1
+target_include_directories(test-q1-lora-opt PRIVATE ${PROJECT_SOURCE_DIR}/ggml/src)
+
+# PRISM_REAL_BONSAI_Q1_LORA_TARGET_V1
+llama_build(test-bonsai-q1-lora-layers.cpp)
+target_include_directories(test-bonsai-q1-lora-layers PRIVATE ${PROJECT_SOURCE_DIR}/ggml/src)
+
+# PRISM_Q1_LORA_LOADER_BLOCK_TARGET_V2
+llama_build(test-q1-lora-loader-blocks.cpp)
+target_include_directories(
+    test-q1-lora-loader-blocks
+    PRIVATE
+        ${PROJECT_SOURCE_DIR}/src
+        ${PROJECT_SOURCE_DIR}/ggml/src
+)
+
+# PRISM_Q1_LORA_COMPLETE_BACKWARD_TARGET_V1
+llama_build(test-q1-lora-full-backward.cpp)
+target_include_directories(
+    test-q1-lora-full-backward
+    PRIVATE
+        ${PROJECT_SOURCE_DIR}/src
+        ${PROJECT_SOURCE_DIR}/ggml/src
+)
+llama_build(test-q1-lora-multitarget.cpp)
+llama_build(test-q1-lora-dataset.cpp)
+llama_build(test-q1-lora-step10.cpp)
+llama_build(test-q1-lora-step11.cpp)
+target_include_directories(test-q1-lora-multitarget PRIVATE ${PROJECT_SOURCE_DIR}/src)
+target_include_directories(test-q1-lora-step10 PRIVATE ${PROJECT_SOURCE_DIR}/src)
+
+# PRISM_STEP10_V481_ROUNDTRIP_TARGET
+add_executable(test-q1-lora-roundtrip test-q1-lora-roundtrip.cpp)
+target_link_libraries(test-q1-lora-roundtrip PRIVATE llama)
+target_include_directories(test-q1-lora-roundtrip PRIVATE ${CMAKE_SOURCE_DIR}/src)
+target_include_directories(test-q1-lora-step11 PRIVATE ${PROJECT_SOURCE_DIR}/src)
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 7512e34..b94b91c 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -4348,6 +4348,126 @@ struct test_mul_mat_id_fusion : public test_case {
     }
 };
 
+
+// PRISM_Q1_EXACT_TRAIN_GRAPH_TEST_V1
+//
+// This operation proves that setting GGML_PREC_F32 on a packed-Q1
+// MUL_MAT selects the exact training-forward CUDA path.
+//
+// Only X is a gradient probe. The packed Q1 base is frozen.
+
+struct test_q1_exact_train_x : public test_case {
+    const int64_t k;
+    const int64_t m;
+    const int64_t batch;
+
+    test_q1_exact_train_x(
+            int64_t k = 256,
+            int64_t m = 16,
+            int64_t batch = 2)
+        : k(k),
+          m(m),
+          batch(batch) {
+    }
+
+    std::string op_desc(
+            ggml_tensor * tensor) override {
+        GGML_UNUSED(tensor);
+        return "Q1_EXACT_TRAIN_X";
+    }
+
+    std::string vars() override {
+        return
+            VAR_TO_STR(k)
+            + ","
+            + VAR_TO_STR(m)
+            + ","
+            + VAR_TO_STR(batch);
+    }
+
+    // PRISM_Q1_EXACT_CROSS_BACKEND_TOLERANCE_V1
+    // This threshold applies only to test mode's
+    // generic cross-backend Q1 comparison. The exact
+    // dequantized oracle is validated independently in
+    // Step 2, and gradient mode retains max_maa_err().
+    double max_nmse_err() override {
+        return 1.0000000000e-04;
+    }
+
+    double max_maa_err() override {
+        return 2e-3;
+    }
+
+    float grad_eps() override {
+        return 3e-2f;
+    }
+
+    bool grad_precise() override {
+        return true;
+    }
+
+    int64_t grad_nmax() override {
+        return 1024;
+    }
+
+    bool run_whole_graph() override {
+        return true;
+    }
+
+    ggml_tensor * build_graph(
+            ggml_context * ctx) override {
+        GGML_ASSERT(
+            k
+            % ggml_blck_size(
+                GGML_TYPE_Q1_0)
+            == 0);
+
+        ggml_tensor * x =
+            ggml_new_tensor_2d(
+                ctx,
+                GGML_TYPE_F32,
+                k,
+                batch);
+
+        ggml_set_name(
+            x,
+            "q1_exact_train.x");
+
+        ggml_set_param(x);
+
+        ggml_tensor * base_weight =
+            ggml_new_tensor_2d(
+                ctx,
+                GGML_TYPE_Q1_0,
+                k,
+                m);
+
+        ggml_set_name(
+            base_weight,
+            "q1_exact_train.base_q1");
+
+        // Deliberately frozen:
+        // no ggml_set_param(base_weight).
+
+        ggml_tensor * output =
+            ggml_mul_mat(
+                ctx,
+                base_weight,
+                x);
+
+        ggml_set_name(
+            output,
+            "q1_exact_train.output");
+
+        // This is the opt-in training-forward selector.
+        ggml_mul_mat_set_prec(
+            output,
+            GGML_PREC_F32);
+
+        return output;
+    }
+};
+
 // GGML_OP_OUT_PROD
 struct test_out_prod : public test_case {
     const ggml_type type_a;
@@ -7635,6 +7755,10 @@ static const ggml_type other_types[] = {
 // Test cases for evaluation: should try to cover edge cases while using small input sizes to keep the runtime low
 static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     std::vector<std::unique_ptr<test_case>> test_cases;
+
+    // PRISM_Q1_EXACT_TRAIN_GRAPH_TEST_V1
+    test_cases.emplace_back(
+        new test_q1_exact_train_x());
     std::default_random_engine rng(0);
 
     // unary ops