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 #include #include +#include +#include +#include // 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 #include +#include -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( + 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( + blockDim.x) + * static_cast( + gridDim.x); + + for ( + int64_t linear = + static_cast( + 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( + 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( + std::min( + 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::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::max()); + + const int batch_count = + static_cast(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 +#include + 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 +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(x + i0*snb0 + i1*snb1 + i2*snb2 + i3*snb3); + T * dp = reinterpret_cast(dst + i0*dnb0 + i1*dnb1 + i2*dnb2 + i3*dnb3); + *dp = (T) op((float) *xp); +} + template 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, launch_params, x, dst, k); } +template +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, 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 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((const half *) src0->data, (half *) dst->data, ggml_nelements(src0), stream); + } else { + unary_cuda((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((const half *)src0_d, (half *)dst_d, ggml_nelements(src0), stream); + unary_cuda_strided(src0, dst, stream); } else { - unary_cuda((const float *)src0_d, (float *)dst_d, ggml_nelements(src0), stream); + unary_cuda_strided(src0, dst, stream); } } - void ggml_cuda_op_abs(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_unary(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 #include #include +#include +#include #include #include #include @@ -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 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(seed_before), + static_cast(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 values( + static_cast(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(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 + + 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 + : "", + + node->src[0] + ? node->src[0]->name + : "", + + node->src[1] + ? node->src[1]->name + : "", + + node->src[2] + ? node->src[2]->name + : "", + + (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 #include +#include +#include #include #include @@ -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 #include #include +#include +#include +#include // // 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( + ubatches_total + 1); + + std::vector> + 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 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 direct_host_sum(expected_gradients, 0.0); + std::vector 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 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(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 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(value); + + if (value != 0.0f) { + ++nonzero; + } + + const double value_d = + static_cast(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(n_ctx) / static_cast(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(sentinel_period) / + static_cast(ubatches_total) * + static_cast(n_ctx) / + static_cast(supervised); + + // PRISM_STEP10_V44_GRADIENT_SOURCE_ARBITRATION_BEGIN + if (ok && train) { + std::vector 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 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(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(accumulator_values[index]) + : direct_host_sum[index]; + + const double scaled_value = + raw_value*static_cast(gradient_scale); + + if (!std::isfinite(scaled_value)) { + ok = false; + break; + } + + gradients_out[output_offset + j] = + static_cast(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 #include #include +#include +#include #include #include #include @@ -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(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(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 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 +#include + +// 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> make_test_cases_eval() { std::vector> 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