diff --git a/common/common.cpp b/common/common.cpp index 84a28d0fe..a0edd82b2 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1482,6 +1482,20 @@ std::string common_get_model_endpoint() { return model_endpoint; } +char * common_get_model_or_exit(int argc, char * argv[]) { + if (argc > 1) { + return argv[1]; + } + + char * path = getenv("LLAMACPP_TEST_MODELFILE"); + if (!path || strlen(path) == 0) { + fprintf(stderr, "\033[33mWARNING: No model file provided. Skipping this test. Set LLAMACPP_TEST_MODELFILE= to silence this warning and run this test.\n\033[0m"); + exit(EXIT_SUCCESS); + } + + return path; +} + common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) { auto * mem = llama_get_memory(ctx); if (mem == nullptr) { diff --git a/common/common.h b/common/common.h index 5ea07d19d..e6d5d892e 100644 --- a/common/common.h +++ b/common/common.h @@ -295,10 +295,6 @@ struct common_params_sampling { bool backend_sampling = false; - bool has_logit_bias() const { - return !logit_bias.empty(); - } - // print the parameters into a string std::string print() const; }; @@ -935,6 +931,9 @@ void common_set_adapter_lora(struct llama_context * ctx, std::vector merged = params.logit_bias; + + int32_t n_suppress = 0; + const llama_token * suppress = llama_vocab_get_suppress_tokens(vocab, &n_suppress); + for (int32_t i = 0; i < n_suppress; ++i) { + merged.push_back({ suppress[i], -INFINITY }); + } + + if (!merged.empty()) { + samplers.push_back(llama_sampler_init_logit_bias(llama_vocab_n_tokens(vocab), merged.size(), merged.data())); + } } if (params.mirostat == 0) { diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 81c234586..9d26a264a 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -469,6 +469,8 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st return (src0->type == GGML_TYPE_F32 || ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) && src1->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; + case GGML_OP_CONV_2D: + return ggml_is_contiguous(op->src[0]); default: return true; } diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index e64b8ad44..14fbc4a88 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -984,6 +984,13 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI1_0; }; +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK2_0; + static constexpr int qr = QR2_0; + static constexpr int qi = QI2_0; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK4_0; diff --git a/ggml/src/ggml-cuda/conv2d.cu b/ggml/src/ggml-cuda/conv2d.cu index 142dd6690..14774d4a5 100644 --- a/ggml/src/ggml-cuda/conv2d.cu +++ b/ggml/src/ggml-cuda/conv2d.cu @@ -126,6 +126,7 @@ void ggml_cuda_op_conv2d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { const float * X_D = (const float *) input->data; float * Y_D = (float *) dst->data; + GGML_ASSERT(ggml_is_contiguous(input)); GGML_ASSERT(ggml_is_contiguous(kernel)); GGML_ASSERT(kernel->type == GGML_TYPE_F16 || kernel->type == GGML_TYPE_F32); diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 946e02af5..360c614a4 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -459,6 +459,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: return dequantize_block_cont_cuda; + case GGML_TYPE_Q2_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -514,6 +516,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: return dequantize_block_cont_cuda; + case GGML_TYPE_Q2_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -572,6 +576,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: return dequantize_block_cont_cuda; + case GGML_TYPE_Q2_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -629,6 +635,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { return convert_unary_cuda; case GGML_TYPE_Q1_0: return dequantize_block_cuda; + case GGML_TYPE_Q2_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -652,6 +660,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) { return convert_unary_cuda; case GGML_TYPE_Q1_0: return dequantize_block_cuda; + case GGML_TYPE_Q2_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -675,6 +685,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) { return convert_unary_cuda; case GGML_TYPE_Q1_0: return dequantize_block_cuda; + case GGML_TYPE_Q2_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index 8ab5ad8e7..651524be7 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -23,6 +23,26 @@ static __device__ __forceinline__ void dequantize_q1_0(const void * vx, const in v.y = (2*bit_1 - 1) * d; } +static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_q2_0 * x = (const block_q2_0 *) vx; + + const float d = x[ib].d; + + // Q2_0: 2 bits per element, 4 elements per byte. + // Stored code c in {0,1,2,3} maps to symbol s = c - 1 in {-1, 0, +1, +2}. + const int byte_index_0 = iqs / 4; + const int bit_offset_0 = (iqs % 4) * 2; + + const int byte_index_1 = (iqs + 1) / 4; + const int bit_offset_1 = ((iqs + 1) % 4) * 2; + + const int c0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 0x3; + const int c1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 0x3; + + v.x = (c0 - 1) * d; + v.y = (c1 - 1) * d; +} + static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_q4_0 * x = (const block_q4_0 *) vx; diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index a9ec4f697..6b36b8fbc 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -320,6 +320,10 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); break; + case GGML_TYPE_Q2_0: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; case GGML_TYPE_Q4_0: get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index ea488c78a..e2459196c 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1836,6 +1836,20 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor ggml_cuda_mul_mat_vec_f(ctx, src0, src1, nullptr, dst); return; } + // A transposed vector can still use MMVQ (i.e. ne01 == 1) + if (ne01 == 1 && ne11 > MMVF_MAX_BATCH_SIZE && ne2 == 1 && ne3 == 1 + && src0->type == GGML_TYPE_F32 + && ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst) + && ggml_cuda_should_use_mmvf(src1->type, cc, src1->ne, src1->nb, /*ne11 =*/ 1)) { + ggml_tensor dst_vec = *dst; + dst_vec.ne[0] = ne11; + dst_vec.ne[1] = 1; + dst_vec.nb[1] = dst_vec.nb[0]*ne11; + dst_vec.nb[2] = dst_vec.nb[1]; + dst_vec.nb[3] = dst_vec.nb[1]; + ggml_cuda_mul_mat_vec_f(ctx, src1, src0, nullptr, &dst_vec); + return; + } if (ggml_cuda_should_use_mmf(src0->type, cc, warp_size, src0->ne, src0->nb, ne11, /*mul_mat_id =*/ false)) { ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst); return; @@ -4815,6 +4829,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_F32: case GGML_TYPE_F16: case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4853,6 +4868,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_BF16: case GGML_TYPE_I32: case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5102,7 +5118,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_IM2COL: case GGML_OP_IM2COL_3D: case GGML_OP_CONV_2D: - return true; + return (ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1])); case GGML_OP_CONV_2D_DW: return op->src[0]->type == GGML_TYPE_F32; case GGML_OP_CONV_TRANSPOSE_2D: diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 0037bac3d..9f9fd1973 100644 --- a/ggml/src/ggml-cuda/mmq-config-ampere.cuh +++ b/ggml/src/ggml-cuda/mmq-config-ampere.cuh @@ -16,6 +16,23 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); CASE(GGML_TYPE_Q1_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q4_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q4_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); diff --git a/ggml/src/ggml-cuda/mmq-config-cdna.cuh b/ggml/src/ggml-cuda/mmq-config-cdna.cuh index 46ec6aa9d..4a8d89f72 100644 --- a/ggml/src/ggml-cuda/mmq-config-cdna.cuh +++ b/ggml/src/ggml-cuda/mmq-config-cdna.cuh @@ -7,6 +7,14 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); CASE(GGML_TYPE_Q1_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q2_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q4_0, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q4_0, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); diff --git a/ggml/src/ggml-cuda/mmq-config-pascal.cuh b/ggml/src/ggml-cuda/mmq-config-pascal.cuh index 8f0faac88..e7d4a9a3f 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal.cuh @@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh index de4db0a3d..8324d9e1a 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh @@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh index e420a32f0..180b2d937 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh @@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q4_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh index 122623060..676f27fea 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh @@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q4_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh index a224ecafc..9293d9d55 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh @@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q1_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false); CASE(GGML_TYPE_Q4_0, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); CASE(GGML_TYPE_Q4_0, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index 7fb242096..8ed704c28 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -95,6 +95,87 @@ template static __device__ __forceinline_ } } +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_0( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + + constexpr int blocks_per_iter = MMQ_ITER_K / QK2_0; + constexpr int threads_per_row = blocks_per_iter * QI2_0; + constexpr int nrows = warp_size / threads_per_row; + constexpr int scale_entries_per_block = QK2_0 / QK8_1; + constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; + + const int txi = threadIdx.x % threads_per_row; + const int kbx = txi / QI2_0; + const int kqsx = txi % QI2_0; + +#pragma unroll + for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { + int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; + + if (fallback) { + i = min(i, i_max); + } + + const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + kbx; + const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 4; + + const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; + +#pragma unroll + for (int j = 0; j < 4; ++j) { + const int q = qxi[j]; + + // unpack even and odd crumbs into byte values + const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0); + const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2); + // unshuffle values + const int qx = __byte_perm(qe, qo, 0x5140); + const int qy = __byte_perm(qe, qo, 0x7362); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_qs[i*sram_stride + dst_offset + j*2+0] = qx; + x_qs[i*sram_stride + dst_offset + j*2+1] = qy; +#else + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+0] = qx; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+1] = qy; +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + } + } + + const int ksx = threadIdx.x % scale_entries_per_row; + const int scale_block = ksx / scale_entries_per_block; + +#pragma unroll + for (int i0 = 0; i0 < I; i0 += nwarps) { + int i = i0 + threadIdx.y; + + if (fallback) { + i = min(i, i_max); + } + + const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + scale_block; + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_df[i*sram_stride + ksx] = bxi->d; +#else + x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + ksx] = bxi->d; +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + } +} + template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 7674c730e..47a15cb63 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -10,6 +10,9 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con case GGML_TYPE_Q1_0: mul_mat_q_case(ctx, args, stream); break; + case GGML_TYPE_Q2_0: + mul_mat_q_case(ctx, args, stream); + break; case GGML_TYPE_Q4_0: mul_mat_q_case(ctx, args, stream); break; @@ -264,6 +267,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t switch (type) { case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -298,6 +302,15 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return false; } + // MMQ tiles require at least 48 KiB per-block shared memory; fall back to BLAS otherwise. + { + const int id = ggml_cuda_get_device(); + const size_t smpbo = ggml_cuda_info().devices[id].smpbo; + if (smpbo < 48 * 1024) { + return false; + } + } + if (turing_mma_available(cc)) { return true; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 75de1711a..0c2a82ee1 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -61,6 +61,7 @@ static_assert(sizeof(block_fp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected b static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { switch (type_x) { case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -386,6 +387,7 @@ static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type, static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml_type type, int I) { switch (type) { case GGML_TYPE_Q1_0: return MMQ_DP4A_TXS_Q8_0; + case GGML_TYPE_Q2_0: return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_Q4_0: return MMQ_DP4A_TXS_Q4_0; case GGML_TYPE_Q4_1: return MMQ_DP4A_TXS_Q4_1; case GGML_TYPE_Q5_0: return MMQ_DP4A_TXS_Q8_0; @@ -543,6 +545,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func ggml_cuda_mmq_load_tiles_q1_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); + case GGML_TYPE_Q2_0: + return ggml_cuda_mmq_util_funcs( + VDR_Q2_0_Q8_1_MMQ, + ggml_cuda_mmq_load_tiles_q2_0, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( VDR_Q4_0_Q8_1_MMQ, @@ -701,6 +709,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func ggml_cuda_mmq_load_tiles_q1_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, ggml_cuda_mmq_write_back_mma); + case GGML_TYPE_Q2_0: + return ggml_cuda_mmq_util_funcs( + -1, + ggml_cuda_mmq_load_tiles_q2_0, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( -1, @@ -1551,6 +1565,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) \ extern DECL_MMQ_CASE(GGML_TYPE_Q1_0); +extern DECL_MMQ_CASE(GGML_TYPE_Q2_0); extern DECL_MMQ_CASE(GGML_TYPE_Q4_0); extern DECL_MMQ_CASE(GGML_TYPE_Q4_1); extern DECL_MMQ_CASE(GGML_TYPE_Q5_0); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index e18ada537..0589e65bd 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -10,6 +10,7 @@ typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: return vec_dot_q1_0_q8_1; + case GGML_TYPE_Q2_0: return vec_dot_q2_0_q8_1; case GGML_TYPE_Q4_0: return vec_dot_q4_0_q8_1; case GGML_TYPE_Q4_1: return vec_dot_q4_1_q8_1; case GGML_TYPE_Q5_0: return vec_dot_q5_0_q8_1; @@ -38,6 +39,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: return VDR_Q1_0_Q8_1_MMVQ; + case GGML_TYPE_Q2_0: return VDR_Q2_0_Q8_1_MMVQ; case GGML_TYPE_Q4_0: return VDR_Q4_0_Q8_1_MMVQ; case GGML_TYPE_Q4_1: return VDR_Q4_1_Q8_1_MMVQ; case GGML_TYPE_Q5_0: return VDR_Q5_0_Q8_1_MMVQ; @@ -1010,6 +1012,12 @@ static void mul_mat_vec_q_switch_type( nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); break; + case GGML_TYPE_Q2_0: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; case GGML_TYPE_Q4_0: mul_mat_vec_q_switch_ncols_dst (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, diff --git a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py index 614b1566c..d7cd27167 100755 --- a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py @@ -36,6 +36,7 @@ SOURCE_FATTN_MMA_CASE = "DECL_FATTN_MMA_F16_CASE({head_size_kq}, {head_size_v}, TYPES_MMQ = [ "GGML_TYPE_Q1_0", + "GGML_TYPE_Q2_0", "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0", "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K", "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S", diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0.cu new file mode 100644 index 000000000..750180e33 --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0.cu @@ -0,0 +1,5 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +DECL_MMQ_CASE(GGML_TYPE_Q2_0); diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index b9932bce9..0f039c735 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -109,6 +109,9 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) { #define VDR_Q1_0_Q8_1_MMVQ 1 // Process one 32-element chunk at a time for parallelism #define VDR_Q1_0_Q8_1_MMQ 4 // Q1_0 has 128 bits (4 ints) per block +#define VDR_Q2_0_Q8_1_MMVQ 1 // Process one 32-element chunk at a time for parallelism +#define VDR_Q2_0_Q8_1_MMQ 2 // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks + #define VDR_Q4_0_Q8_1_MMVQ 2 #define VDR_Q4_0_Q8_1_MMQ 4 @@ -722,6 +725,44 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1( return d1 * d8 * sumi; } +static __device__ __forceinline__ float vec_dot_q2_0_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_q2_0 * bq2_0 = (const block_q2_0 *) vbq + kbx; + + // Q2_0 (group 64): 64 elements with ONE scale, 2 bits per element (4 elements per byte) + // Q8_1: 32 elements per block with individual scales + // iqs selects which of the 2 chunks of 32 elements to process (0-1) + + const float d2 = bq2_0->d; + const int16_t * qs = (const int16_t *) bq2_0->qs + iqs * 4; + + // Process only the chunk specified by iqs + const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; + + int sumi = 0; +#pragma unroll + for (int j = 0; j < 4; ++j) { + const int q = qs[j]; + const int u = get_int_b4(bq8_1_chunk->qs, j*2+0); + const int v = get_int_b4(bq8_1_chunk->qs, j*2+1); + + // unpack even and odd crumbs into byte values + const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0); + const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2); + // unshuffle values + const int qx = __byte_perm(qe, qo, 0x5140); + const int qy = __byte_perm(qe, qo, 0x7362); + + sumi = ggml_cuda_dp4a(u, qx, sumi); + sumi = ggml_cuda_dp4a(v, qy, sumi); + } + + // Apply Q2_0's single scale and this chunk's Q8_1 scale + const float d8 = __low2float(bq8_1_chunk->ds); + return d2 * d8 * sumi; +} + static __device__ __forceinline__ float vec_dot_q4_0_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index d0956df50..91b841b67 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -213,7 +213,7 @@ typedef void * ggml_metal_rset_t; // a collection of residency sets (non-owning) typedef struct ggml_metal_rsets * ggml_metal_rsets_t; -ggml_metal_rsets_t ggml_metal_rsets_init(void); +ggml_metal_rsets_t ggml_metal_rsets_init(ggml_metal_device_t dev); void ggml_metal_rsets_free(ggml_metal_rsets_t rsets); // diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index a266f76b7..9f0fb6175 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -557,7 +557,32 @@ struct ggml_metal_rsets { dispatch_group_t d_group; }; -ggml_metal_rsets_t ggml_metal_rsets_init(void) { +#if defined(GGML_METAL_HAS_RESIDENCY_SETS) +static void ggml_metal_dummy_work(ggml_metal_device_t dev) { + if (dev->mtl_queue == nil) { + return; + } + + @autoreleasepool { + // perform a minimal dummy operation on the GPU + id buf = [dev->mtl_device newBufferWithLength:1 options:MTLResourceStorageModePrivate]; + id cmd_buf = [dev->mtl_queue commandBuffer]; + + { + id encoder = [cmd_buf blitCommandEncoder]; + + [encoder fillBuffer:buf range:NSMakeRange(0, 1) value:0]; + + [encoder endEncoding]; + } + + [cmd_buf commit]; + [buf release]; + } +} +#endif + +ggml_metal_rsets_t ggml_metal_rsets_init(ggml_metal_device_t dev) { ggml_metal_rsets_t res = calloc(1, sizeof(struct ggml_metal_rsets)); res->lock = [[NSLock alloc] init]; @@ -610,6 +635,15 @@ ggml_metal_rsets_t ggml_metal_rsets_init(void) { #endif }); +#if defined(GGML_METAL_HAS_RESIDENCY_SETS) + if (@available(macOS 15.0, iOS 18.0, tvOS 18.0, visionOS 2.0, *)) { + // workaround for residency set memory not being released if no GPU operation occurs + // https://developer.apple.com/forums/thread/839089 + // https://github.com/ggml-org/llama.cpp/issues/25937 + ggml_metal_dummy_work(dev); + } +#endif + return res; } @@ -870,7 +904,7 @@ ggml_metal_device_t ggml_metal_device_init(int device) { } if (dev->props.use_residency_sets) { - dev->rsets = ggml_metal_rsets_init(); + dev->rsets = ggml_metal_rsets_init(dev); } else { dev->rsets = nil; } @@ -1490,6 +1524,7 @@ static void ggml_metal_buffer_rset_free(ggml_metal_buffer_t buf) { if (buf->rset) { [buf->rset endResidency]; [buf->rset removeAllAllocations]; + [buf->rset commit]; [buf->rset release]; } } diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 5446df1b3..dab4401b4 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1487,6 +1487,11 @@ struct vk_op_binary_push_constants { float param1; float param2; int32_t param3; }; +// Distinct type with the same layout so concat can overload tensor offset initialization. +struct vk_op_concat_push_constants : vk_op_binary_push_constants {}; +static_assert(sizeof(vk_op_concat_push_constants) == sizeof(vk_op_binary_push_constants)); +static_assert(std::is_standard_layout_v); + struct vk_op_multi_add_push_constants { // shape for dst uint32_t ne20; uint32_t ne21; uint32_t ne22; uint32_t ne23; @@ -2252,6 +2257,40 @@ static uint32_t get_misalign_bytes(const ggml_backend_vk_context * ctx, const gg return ((vk_tensor_offset(t) + t->view_offs) & (ctx->device->properties.limits.minStorageBufferOffsetAlignment - 1));; } +static uint32_t ggml_vk_concat_unit_size(ggml_type type) { + const uint32_t type_size = ggml_type_size(type); + + if (!ggml_is_quantized(type)) { + return type_size; + } + + // Use the widest existing concat shader that evenly divides a quant block. + if (type_size % 8 == 0) { + return 8; + } + if (type_size % 4 == 0) { + return 4; + } + if (type_size % 2 == 0) { + return 2; + } + return 1; +} + +static bool ggml_vk_concat_supported(const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * dst) { + if (src0->type != src1->type || src0->type != dst->type) { + return false; + } + + if (!ggml_is_quantized(src0->type)) { + const size_t type_size = ggml_type_size(src0->type); + return type_size == 1 || type_size == 2 || type_size == 4 || type_size == 8; + } + + // Quantized tensor rows are block-aligned when created. + return ggml_is_contiguous_rows(src0) && ggml_is_contiguous_rows(src1) && ggml_is_contiguous_rows(dst); +} + template void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, T &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) { GGML_UNUSED(p); GGML_UNUSED(src0); @@ -10929,14 +10968,10 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const } return nullptr; case GGML_OP_CONCAT: { - if (src0->type != src1->type || src0->type != dst->type) { + if (!ggml_vk_concat_supported(src0, src1, dst)) { return nullptr; } - if (ggml_blck_size(src0->type) != 1) { - return nullptr; - } - const size_t type_size = ggml_type_size(src0->type); - switch (type_size) { + switch (ggml_vk_concat_unit_size(src0->type)) { case 1: return ctx->device->pipeline_concat_i8; case 2: @@ -11628,6 +11663,18 @@ template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk GGML_UNUSED(src3); } +template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk_op_concat_push_constants &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) { + const uint32_t unit_size = ggml_vk_concat_unit_size(dst->type); + const uint32_t a_offset = get_misalign_bytes(ctx, src0) / unit_size; + const uint32_t b_offset = get_misalign_bytes(ctx, src1) / unit_size; + const uint32_t d_offset = get_misalign_bytes(ctx, dst) / unit_size; + + p.misalign_offsets = (a_offset << 16) | (b_offset << 8) | d_offset; + + GGML_UNUSED(src2); + GGML_UNUSED(src3); +} + template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk_op_upscale_push_constants &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) { const uint32_t a_offset = get_misalign_bytes(ctx, src0) / ggml_type_size(src0->type); const uint32_t d_offset = get_misalign_bytes(ctx, dst) / ggml_type_size(dst->type); @@ -11663,7 +11710,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co } std::cerr << "), (" << dst << ", name=" << dst->name << ", type=" << dst->type << ", ne0=" << dst->ne[0] << ", ne1=" << dst->ne[1] << ", ne2=" << dst->ne[2] << ", ne3=" << dst->ne[3] << ", nb0=" << dst->nb[0] << ", nb1=" << dst->nb[1] << ", nb2=" << dst->nb[2] << ", nb3=" << dst->nb[3]; std::cerr << "), " << ggml_op_name(op) << ")"); - GGML_ASSERT(op == GGML_OP_GET_ROWS || op == GGML_OP_CPY || (!ggml_is_quantized(src0->type) && (src1 == nullptr || !ggml_is_quantized(src1->type)))); // NOLINT + GGML_ASSERT(op == GGML_OP_GET_ROWS || op == GGML_OP_CPY || op == GGML_OP_CONCAT || (!ggml_is_quantized(src0->type) && (src1 == nullptr || !ggml_is_quantized(src1->type)))); // NOLINT GGML_ASSERT(dst->buffer != nullptr); const uint64_t ne00 = src0->ne[0]; const uint64_t ne01 = src0->ne[1]; @@ -11918,6 +11965,9 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co ne *= ggml_type_size(src0->type) / 2; } } + if (op == GGML_OP_CONCAT && ggml_is_quantized(dst->type)) { + ne = ne / ggml_blck_size(dst->type) * ggml_type_size(dst->type) / ggml_vk_concat_unit_size(dst->type); + } // copy_to_quant has block size of 32, and each thread does QUANT_K elements. // Splitting into 512x512xZ wouldn't work well since each workgroup does 1024 elements. // So divide by block size here before splitting into 512x512 groups. @@ -12558,18 +12608,28 @@ static void ggml_vk_opt_step_sgd(ggml_backend_vk_context * ctx, vk_context& subc static void ggml_vk_concat(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { int * op_params = (int *)dst->op_params; - const uint32_t src0_type_size = ggml_type_size(src0->type); - const uint32_t src1_type_size = ggml_type_size(src1->type); - const uint32_t dst_type_size = ggml_type_size(dst->type); + const uint32_t unit_size = ggml_vk_concat_unit_size(dst->type); + const uint32_t units_per_block = ggml_type_size(dst->type) / unit_size; + const uint32_t block_size = ggml_blck_size(dst->type); + const bool quantized = ggml_is_quantized(dst->type); - ggml_vk_op_f32(ctx, subctx, src0, src1, nullptr, nullptr, dst, GGML_OP_CONCAT, { - (uint32_t)ggml_nelements(dst), - (uint32_t)src0->ne[0], (uint32_t)src0->ne[1], (uint32_t)src0->ne[2],(uint32_t)src0->ne[3], (uint32_t)src0->nb[0] / src0_type_size, (uint32_t)src0->nb[1] / src0_type_size, (uint32_t)src0->nb[2] / src0_type_size, (uint32_t)src0->nb[3] / src0_type_size, - (uint32_t)src1->ne[0], (uint32_t)src1->ne[1], (uint32_t)src1->ne[2],(uint32_t)src1->ne[3], (uint32_t)src1->nb[0] / src1_type_size, (uint32_t)src1->nb[1] / src1_type_size, (uint32_t)src1->nb[2] / src1_type_size, (uint32_t)src1->nb[3] / src1_type_size, - (uint32_t) dst->ne[0], (uint32_t) dst->ne[1], (uint32_t) dst->ne[2],(uint32_t) dst->ne[3], (uint32_t) dst->nb[0] / dst_type_size, (uint32_t) dst->nb[1] / dst_type_size, (uint32_t) dst->nb[2] / dst_type_size, (uint32_t) dst->nb[3] / dst_type_size, + // Address dimension 0 in packed storage units; higher strides may be noncontiguous. + const uint32_t ne00 = src0->ne[0] / block_size * units_per_block; + const uint32_t ne10 = src1->ne[0] / block_size * units_per_block; + const uint32_t ne20 = dst->ne[0] / block_size * units_per_block; + const uint32_t nb00 = quantized ? 1 : src0->nb[0] / unit_size; + const uint32_t nb10 = quantized ? 1 : src1->nb[0] / unit_size; + const uint32_t nb20 = quantized ? 1 : dst->nb[0] / unit_size; + + vk_op_concat_push_constants pc {{ + ne20 * (uint32_t)dst->ne[1] * (uint32_t)dst->ne[2] * (uint32_t)dst->ne[3], + ne00, (uint32_t)src0->ne[1], (uint32_t)src0->ne[2],(uint32_t)src0->ne[3], nb00, (uint32_t)src0->nb[1] / unit_size, (uint32_t)src0->nb[2] / unit_size, (uint32_t)src0->nb[3] / unit_size, + ne10, (uint32_t)src1->ne[1], (uint32_t)src1->ne[2],(uint32_t)src1->ne[3], nb10, (uint32_t)src1->nb[1] / unit_size, (uint32_t)src1->nb[2] / unit_size, (uint32_t)src1->nb[3] / unit_size, + ne20, (uint32_t) dst->ne[1], (uint32_t) dst->ne[2],(uint32_t) dst->ne[3], nb20, (uint32_t) dst->nb[1] / unit_size, (uint32_t) dst->nb[2] / unit_size, (uint32_t) dst->nb[3] / unit_size, 0, 0.0f, 0.0f, op_params[0], - }); + }}; + ggml_vk_op_f32(ctx, subctx, src0, src1, nullptr, nullptr, dst, GGML_OP_CONCAT, std::move(pc)); } static void ggml_vk_upscale(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) { @@ -17905,12 +17965,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm return op->src[0]->type == op->src[1]->type && op->src[0]->type == op->type && (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_I32); case GGML_OP_CONCAT: { - if (op->src[0]->type != op->src[1]->type || op->src[0]->type != op->type) { - return false; - } - const size_t type_size = ggml_type_size(op->type); - return ggml_blck_size(op->type) == 1 && - (type_size == 1 || type_size == 2 || type_size == 4 || type_size == 8); + return ggml_vk_concat_supported(op->src[0], op->src[1], op); } case GGML_OP_ADD1: return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32) @@ -18049,10 +18104,17 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_OP_CONV_2D: case GGML_OP_CONV_TRANSPOSE_2D: { + const bool transpose = op->op == GGML_OP_CONV_TRANSPOSE_2D; + const int64_t cout = !transpose ? op->src[0]->ne[3] : op->src[0]->ne[2]; + const int64_t cin = !transpose ? op->src[0]->ne[2] : op->src[0]->ne[3]; + // Channel-contiguous format is not supported yet. return ((op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) && + (op->src[0]->nb[0] == sizeof(float) || op->src[0]->nb[0] == sizeof(ggml_fp16_t) ) && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && + cout == op->ne[2] && + cin == op->src[1]->ne[2] && ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]) && ggml_is_contiguous(op)); diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 124ea28b0..650f1c8a5 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -4440,8 +4440,11 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = { MODEL_TENSOR.FFN_EXP_PROBS_B, MODEL_TENSOR.LAYER_OUT_NORM, MODEL_TENSOR.NEXTN_EH_PROJ, + MODEL_TENSOR.NEXTN_EMBED_TOKENS, MODEL_TENSOR.NEXTN_ENORM, MODEL_TENSOR.NEXTN_HNORM, + MODEL_TENSOR.NEXTN_SHARED_HEAD_HEAD, + MODEL_TENSOR.NEXTN_SHARED_HEAD_NORM, ], MODEL_ARCH.STEP35: [ MODEL_TENSOR.TOKEN_EMBD, diff --git a/include/llama.h b/include/llama.h index c4feba06f..f501b7752 100644 --- a/include/llama.h +++ b/include/llama.h @@ -1105,6 +1105,9 @@ extern "C" { LLAMA_API bool llama_vocab_get_add_eos(const struct llama_vocab * vocab); LLAMA_API bool llama_vocab_get_add_sep(const struct llama_vocab * vocab); + // model-specific suppress tokens (gguf key: tokenizer.ggml.suppress_tokens) + LLAMA_API const llama_token * llama_vocab_get_suppress_tokens(const struct llama_vocab * vocab, int32_t * n_suppress_tokens); + LLAMA_API llama_token llama_vocab_fim_pre(const struct llama_vocab * vocab); LLAMA_API llama_token llama_vocab_fim_suf(const struct llama_vocab * vocab); LLAMA_API llama_token llama_vocab_fim_mid(const struct llama_vocab * vocab); diff --git a/src/llama-context.cpp b/src/llama-context.cpp index a29dc5e7e..9bc4361c9 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -482,6 +482,9 @@ llama_context::llama_context( } llama_context::~llama_context() { + // wait for any pending asynchronous copies into the output buffers before they are freed + synchronize(); + if (!model.hparams.no_alloc) { for (size_t i = 0; i < backend_ptrs.size(); ++i) { ggml_backend_t backend = backend_ptrs[i]; @@ -1427,13 +1430,17 @@ int llama_context::encode(const llama_batch & batch_inp) { // micro-batching is not possible for non-causal encoding, so we process the batch in a single shot GGML_ASSERT(cparams.n_ubatch >= n_tokens && "encoder requires n_ubatch >= n_tokens"); + // TODO: this clear of the buffer can easily be forgotten - need something better + // sync first so any in-flight async copies into embd_seq complete before it is freed + if (!embd_seq.empty()) { + synchronize(); + } + embd_seq.clear(); + if (t_compute_start_us == 0) { t_compute_start_us = ggml_time_us(); } - // TODO: this clear of the buffer can easily be forgotten - need something better - embd_seq.clear(); - sched_reserve(); n_queued_tokens += n_tokens; @@ -1772,13 +1779,18 @@ int llama_context::decode(const llama_batch & batch_inp) { // GGML_ASSERT((cparams.causal_attn || cparams.n_ubatch >= n_tokens_all) && "non-causal attention requires n_ubatch >= n_tokens"); + // TODO: this clear of the buffer can easily be forgotten - need something better + // sync first so any in-flight async copies into embd_seq complete before it is freed + if (!embd_seq.empty()) { + synchronize(); + } + embd_seq.clear(); + if (t_compute_start_us == 0) { t_compute_start_us = ggml_time_us(); } n_queued_tokens += n_tokens_all; - // TODO: this clear of the buffer can easily be forgotten - need something better - embd_seq.clear(); output_swaps.clear(); sched_reserve(); @@ -3551,6 +3563,22 @@ llama_context * llama_init_from_model( } } + if ((model->hparams.is_mla() || model->arch == LLM_ARCH_DEEPSEEK4) && params.type_k != params.type_v) { + LLAMA_LOG_ERROR("%s: model does not support different K (%s) and V (%s) cache types\n", __func__, ggml_type_name(params.type_k), ggml_type_name(params.type_v)); + return nullptr; + } + + if (ggml_is_quantized(params.type_v) && params.flash_attn_type != LLAMA_FLASH_ATTN_TYPE_ENABLED) { + if (params.flash_attn_type == LLAMA_FLASH_ATTN_TYPE_AUTO) { + LLAMA_LOG_INFO("%s: enabling flash_attn since it is required for quantized V cache\n", __func__); + params.flash_attn_type = LLAMA_FLASH_ATTN_TYPE_ENABLED; + } + if (params.flash_attn_type == LLAMA_FLASH_ATTN_TYPE_DISABLED) { + LLAMA_LOG_ERROR("%s: quantized V cache requires flash_attn to be enabled\n", __func__); + return nullptr; + } + } + if (params.flash_attn_type != LLAMA_FLASH_ATTN_TYPE_DISABLED && ggml_is_quantized(params.type_k)) { const uint32_t blck_size = ggml_blck_size(params.type_k); for (uint32_t il = 0; il < model->hparams.n_layer(); ++il) { @@ -3573,11 +3601,6 @@ llama_context * llama_init_from_model( } } - if (ggml_is_quantized(params.type_v) && params.flash_attn_type == LLAMA_FLASH_ATTN_TYPE_DISABLED) { - LLAMA_LOG_ERROR("%s: V cache quantization requires flash_attn\n", __func__); - return nullptr; - } - if (params.pooling_type != LLAMA_POOLING_TYPE_UNSPECIFIED && params.pooling_type != model->hparams.pooling_type) { //user-specified pooling-type is different from the model default diff --git a/src/llama-model.cpp b/src/llama-model.cpp index dc54ac7a6..302de7017 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -2385,7 +2385,9 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, filter = [&](uint32_t il) { return il >= hparams.n_layer(); }; } - if ((arch == LLM_ARCH_STEP35 || arch == LLM_ARCH_HY_V3 || arch == LLM_ARCH_GLM_DSA) && hparams.n_layer_nextn > 0) { + if ((arch == LLM_ARCH_STEP35 || arch == LLM_ARCH_HY_V3 || arch == LLM_ARCH_GLM_DSA || + arch == LLM_ARCH_MIMO2) && + hparams.n_layer_nextn > 0) { if (params.ctx_type == LLAMA_CONTEXT_TYPE_MTP) { filter = [&](uint32_t il) { return il >= hparams.n_layer(); }; } else { diff --git a/src/llama-vocab.cpp b/src/llama-vocab.cpp index bbc5c6f73..bbcaf9ee7 100644 --- a/src/llama-vocab.cpp +++ b/src/llama-vocab.cpp @@ -2817,7 +2817,14 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) { if (suppress_idx != -1) { const int n = gguf_get_arr_n(ctx, suppress_idx); const int32_t * data = (const int32_t *) gguf_get_arr_data(ctx, suppress_idx); - suppress_tokens.assign(data, data + n); + // drop out-of-range ids + suppress_tokens.reserve(n); + for (int i = 0; i < n; ++i) { + const int32_t id = data[i]; + if (id >= 0 && id < (int) id_to_token.size()) { + suppress_tokens.push_back(id); + } + } } } @@ -4506,6 +4513,14 @@ bool llama_vocab_get_add_sep(const struct llama_vocab * vocab) { return vocab->get_add_sep(); } +const llama_token * llama_vocab_get_suppress_tokens(const struct llama_vocab * vocab, int32_t * n_suppress_tokens) { + const std::vector & tokens = vocab->get_suppress_tokens(); + if (n_suppress_tokens) { + *n_suppress_tokens = (int32_t) tokens.size(); + } + return tokens.data(); +} + llama_token llama_vocab_fim_pre(const struct llama_vocab * vocab) { return vocab->token_fim_pre(); } diff --git a/src/models/gemma4.cpp b/src/models/gemma4.cpp index 3511afd0d..c4f203380 100644 --- a/src/models/gemma4.cpp +++ b/src/models/gemma4.cpp @@ -142,33 +142,6 @@ std::unique_ptr llama_model_gemma4::build_arch_graph(const ll // idx * x->ne[0] * x->ne[1] * ggml_element_size(x)); // } -// TODO @ngxson : maybe improve this in the future -class llm_graph_input_logits_bias : public llm_graph_input_i { -public: - llm_graph_input_logits_bias(const llama_vocab & vocab) { - arr.resize(vocab.n_tokens(), 0.0f); - for (llama_token id : vocab.get_suppress_tokens()) { - if (0 <= id && id < (int32_t)vocab.n_tokens()) { - arr[id] = -INFINITY; - } - } - } - virtual ~llm_graph_input_logits_bias() = default; - - void set_input(const llama_ubatch * /*ubatch*/) override { - const int64_t n_vocab = arr.size(); - ggml_backend_tensor_set(logits_bias, arr.data(), 0, n_vocab*ggml_element_size(logits_bias)); - } - - bool can_reuse(const llm_graph_params & /*params*/) override { - return true; - } - - ggml_tensor * logits_bias = nullptr; // F32 [n_vocab] - - std::vector arr; -}; - llama_model_gemma4::graph::graph(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params), model(model), @@ -429,16 +402,6 @@ llama_model_gemma4::graph::graph(const llama_model & model, const llm_graph_para cur = ggml_scale(ctx0, cur, hparams.f_final_logit_softcapping); } - // apply logits bias if needed (e.g. for gemma4_unified patch) - // this is to mirror the suppress_tokens patch on transformers, to avoid model from outputing and tokens (which is a known issue related to the checkpoint) - // TODO: maybe handle this inside the sampling system in the future - if (!model.vocab.get_suppress_tokens().empty()) { - auto inp_bias = std::make_unique(model.vocab); - inp_bias->logits_bias = ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, inp_bias->arr.size()); - cur = ggml_add(ctx0, cur, inp_bias->logits_bias); - res->add_input(std::move(inp_bias)); - } - cb(cur, "result_output", -1); res->t_logits = cur; diff --git a/src/models/mimo2.cpp b/src/models/mimo2.cpp index 889891605..4080a934c 100644 --- a/src/models/mimo2.cpp +++ b/src/models/mimo2.cpp @@ -25,9 +25,13 @@ void llama_model_mimo2::load_arch_hparams(llama_model_loader & ml) { } } -void llama_model_mimo2::load_arch_tensors(llama_model_loader &) { +void llama_model_mimo2::load_arch_tensors(llama_model_loader & ml) { LLAMA_LOAD_LOCALS; + const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight"; + const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr); + const int mtp_flags = trunk_only ? TENSOR_NOT_REQUIRED : 0; + tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0); // output @@ -40,41 +44,46 @@ void llama_model_mimo2::load_arch_tensors(llama_model_loader &) { uint32_t n_embd_v_gqa = hparams.n_embd_v_gqa(i); uint32_t n_head = hparams.n_head(i); - // NextN/MTP layers (the last n_nextn blocks) are preserved but disabled pending support const bool is_nextn = i >= n_layer; - const int skip = is_nextn ? TENSOR_SKIP : 0; + const int flags = is_nextn ? mtp_flags : 0; - create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head, n_embd_k_gqa, n_embd_v_gqa, skip); - layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), { n_embd_head_v * n_head, n_embd }, skip); + create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head, n_embd_k_gqa, n_embd_v_gqa, flags); + layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), { n_embd_head_v * n_head, n_embd }, flags); - layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, skip); - layer.attn_sinks = create_tensor(tn(LLM_TENSOR_ATTN_SINKS, "weight", i), {n_head}, TENSOR_NOT_REQUIRED | skip); + layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, flags); + layer.attn_sinks = create_tensor(tn(LLM_TENSOR_ATTN_SINKS, "weight", i), {n_head}, TENSOR_NOT_REQUIRED | flags); - layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, skip); + layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, flags); // non-MoE branch - layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED | skip); - layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd}, TENSOR_NOT_REQUIRED | skip); - layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED | skip); + layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED | flags); + layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd}, TENSOR_NOT_REQUIRED | flags); + layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED | flags); // MoE branch int64_t n_ff_exp = hparams.n_ff_exp; - layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, TENSOR_NOT_REQUIRED | skip); - layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED | skip); - layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, TENSOR_NOT_REQUIRED | skip); - layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED | skip); - layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, TENSOR_NOT_REQUIRED | skip); + layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, TENSOR_NOT_REQUIRED | flags); + layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED | flags); + layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, TENSOR_NOT_REQUIRED | flags); + layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED | flags); + layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, TENSOR_NOT_REQUIRED | flags); if (is_nextn) { - layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), {2 * n_embd, n_embd}, skip); - layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), {n_embd}, skip); - layer.nextn.hnorm = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", i), {n_embd}, skip); - layer.layer_out_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, skip); + layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), {2 * n_embd, n_embd}, flags); + layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), {n_embd}, flags); + layer.nextn.hnorm = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", i), {n_embd}, flags); + layer.nextn.embed_tokens = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED | flags); + layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", i), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED | flags); + layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), {n_embd}, TENSOR_NOT_REQUIRED | flags); + layer.layer_out_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, TENSOR_NOT_REQUIRED | flags); } } } std::unique_ptr llama_model_mimo2::build_arch_graph(const llm_graph_params & params) const { + if (params.gtype == LLM_GRAPH_TYPE_DECODER_MTP) { + return std::make_unique(*this, params); + } return std::make_unique(*this, params); } @@ -89,6 +98,8 @@ llama_model_mimo2::graph::graph(const llama_model & model, const llm_graph_param ggml_tensor * inp_out_ids = build_inp_out_ids(); const float v_scale = hparams.f_attn_value_scale; + const bool emit_h_nextn = cparams.embeddings_nextn; + const bool crop_last_layer = inp_out_ids && (!emit_h_nextn || cparams.embeddings_nextn_masked); for (int il = 0; il < n_layer; ++il) { ggml_tensor * inpSA = inpL; @@ -168,7 +179,7 @@ llama_model_mimo2::graph::graph(const llama_model & model, const llm_graph_param } } - if (il == n_layer - 1 && inp_out_ids) { + if (il == n_layer - 1 && crop_last_layer) { cur = ggml_get_rows(ctx0, cur, inp_out_ids); inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids); } @@ -218,6 +229,15 @@ llama_model_mimo2::graph::graph(const llama_model & model, const llm_graph_param cur = inpL; + if (emit_h_nextn) { + cb(cur, "h_nextn", -1); + res->t_h_nextn = cur; + + if (!cparams.embeddings_nextn_masked && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + } + } + cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); @@ -233,3 +253,143 @@ llama_model_mimo2::graph::graph(const llama_model & model, const llm_graph_param ggml_build_forward_expand(gf, cur); } + +// Mirrors MiMo's appended NextN block: normalize and fuse token and hidden inputs, run the decoder block, +// expose its pre-head-norm state to the next draft step, then apply the shared output norm and LM head. +// Converted checkpoints may store that shared norm as layer_out_norm, so it remains in the fallback chain. +llama_model_mimo2::graph_mtp::graph_mtp(const llama_model & model, const llm_graph_params & params) + : llm_graph_context(params) { + GGML_ASSERT(hparams.n_layer_nextn > 0 && "MIMO2 MTP requires n_layer_nextn > 0"); + + const int il = hparams.n_layer() + cparams.nextn_layer_offset; + GGML_ASSERT(cparams.nextn_layer_offset >= 0 && + cparams.nextn_layer_offset < (int) hparams.n_layer_nextn && + "nextn_layer_offset out of range [0, n_layer_nextn)"); + + const auto & layer = model.layers[il]; + GGML_ASSERT(layer.nextn.eh_proj && "MIMO2 MTP block missing nextn.eh_proj"); + GGML_ASSERT(layer.nextn.enorm && "MIMO2 MTP block missing nextn.enorm"); + GGML_ASSERT(layer.nextn.hnorm && "MIMO2 MTP block missing nextn.hnorm"); + GGML_ASSERT(layer.wqkv && "MIMO2 MTP requires fused attn_qkv"); + + const uint32_t n_head_l = hparams.n_head(il); + const uint32_t n_head_kv_l = hparams.n_head_kv(il); + + const float freq_base_l = model.get_rope_freq_base(cparams, il); + const float freq_scale_l = model.get_rope_freq_scale(cparams, il); + const float v_scale = hparams.f_attn_value_scale; + + auto inp = std::make_unique(hparams.n_embd); + + inp->tokens = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens); + ggml_set_input(inp->tokens); + + inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, hparams.n_embd, n_tokens); + ggml_set_input(inp->embd); + ggml_set_name(inp->embd, "mtp_h_input"); + + ggml_tensor * tok_embd_w = layer.nextn.embed_tokens ? layer.nextn.embed_tokens : model.tok_embd; + ggml_tensor * h_input = inp->embd; + ggml_tensor * tok_embd = ggml_get_rows(ctx0, tok_embd_w, inp->tokens); + cb(tok_embd, "mtp_tok_embd", il); + + res->add_input(std::move(inp)); + + ggml_tensor * inp_pos = build_inp_pos(); + ggml_tensor * inp_out_ids = build_inp_out_ids(); + auto * inp_attn = build_attn_inp_kv_iswa(); + + ggml_tensor * h_norm = build_norm(h_input, layer.nextn.hnorm, nullptr, LLM_NORM_RMS, il); + cb(h_norm, "mtp_hnorm", il); + + ggml_tensor * e_norm = build_norm(tok_embd, layer.nextn.enorm, nullptr, LLM_NORM_RMS, il); + cb(e_norm, "mtp_enorm", il); + + ggml_tensor * concat = ggml_concat(ctx0, e_norm, h_norm, /*dim=*/ 0); + cb(concat, "mtp_concat", il); + + ggml_tensor * cur = build_lora_mm(layer.nextn.eh_proj, concat, layer.nextn.eh_proj_s); + cb(cur, "mtp_eh_proj", il); + + ggml_tensor * inpSA = cur; + + cur = build_norm(cur, layer.attn_norm, nullptr, LLM_NORM_RMS, il); + cb(cur, "mtp_attn_norm", il); + + ggml_tensor * qkv = build_lora_mm(layer.wqkv, cur, layer.wqkv_s); + cb(qkv, "mtp_wqkv", il); + + const size_t row_k = ggml_row_size(qkv->type, n_embd_head_k); + const size_t row_v = ggml_row_size(qkv->type, n_embd_head_v); + const size_t row_full = qkv->nb[1]; + const size_t k_off = row_k * n_head_l; + const size_t v_off = k_off + row_k * n_head_kv_l; + + ggml_tensor * Qcur = ggml_view_3d(ctx0, qkv, n_embd_head_k, n_head_l, n_tokens, row_k, row_full, 0); + ggml_tensor * Kcur = ggml_view_3d(ctx0, qkv, n_embd_head_k, n_head_kv_l, n_tokens, row_k, row_full, k_off); + ggml_tensor * Vcur = ggml_view_3d(ctx0, qkv, n_embd_head_v, n_head_kv_l, n_tokens, row_v, row_full, v_off); + + Qcur = ggml_rope_ext( + ctx0, Qcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base_l, freq_scale_l, + ext_factor, attn_factor, beta_fast, beta_slow); + + Kcur = ggml_rope_ext( + ctx0, Kcur, inp_pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base_l, freq_scale_l, + ext_factor, attn_factor, beta_fast, beta_slow); + + cb(Qcur, "mtp_Qcur", il); + cb(Kcur, "mtp_Kcur", il); + cb(Vcur, "mtp_Vcur", il); + + cur = build_attn(inp_attn, + layer.wo, nullptr, layer.wo_s, + Qcur, Kcur, Vcur, nullptr, layer.attn_sinks, nullptr, + 1.0f / sqrtf(float(n_embd_head_k)), il); + cb(cur, "mtp_attn_out", il); + + if (v_scale) { + cur = ggml_scale(ctx0, cur, v_scale); + cb(cur, "mtp_attn_out_scaled", il); + } + + ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA); + cb(ffn_inp, "mtp_ffn_inp", il); + + cur = build_norm(ffn_inp, layer.ffn_norm, nullptr, LLM_NORM_RMS, il); + cb(cur, "mtp_ffn_norm", il); + + GGML_ASSERT(layer.ffn_gate && layer.ffn_down && layer.ffn_up && "MIMO2 MTP requires dense FFN tensors"); + cur = build_ffn(cur, + layer.ffn_up, layer.ffn_up_b, nullptr, + layer.ffn_gate, layer.ffn_gate_b, nullptr, + layer.ffn_down, layer.ffn_down_b, nullptr, + nullptr, + LLM_FFN_SILU, LLM_FFN_PAR, il); + cb(cur, "mtp_ffn_out", il); + + cur = ggml_add(ctx0, cur, ffn_inp); + cb(cur, "mtp_post_ffn", il); + + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + + cb(cur, "h_nextn", -1); + res->t_h_nextn = cur; + + ggml_tensor * head_norm_w = layer.nextn.shared_head_norm + ? layer.nextn.shared_head_norm + : (layer.layer_out_norm ? layer.layer_out_norm : model.output_norm); + GGML_ASSERT(head_norm_w && "MIMO2 MTP missing head norm fallback"); + cur = build_norm(cur, head_norm_w, nullptr, LLM_NORM_RMS, -1); + cb(cur, "mtp_shared_head_norm", -1); + + ggml_tensor * head_w = layer.nextn.shared_head_head ? layer.nextn.shared_head_head : model.output; + ggml_tensor * head_s = layer.nextn.shared_head_head ? layer.nextn.shared_head_head_s : model.output_s; + GGML_ASSERT(head_w && "MIMO2 MTP missing LM head fallback"); + cur = build_lora_mm(head_w, cur, head_s); + cb(cur, "result_output", -1); + + res->t_logits = cur; + ggml_build_forward_expand(gf, cur); +} diff --git a/src/models/minimax-m3.cpp b/src/models/minimax-m3.cpp index 6068fc6b8..3e7bada64 100644 --- a/src/models/minimax-m3.cpp +++ b/src/models/minimax-m3.cpp @@ -2,7 +2,6 @@ #include "llama-kv-cache.h" #include #include -#include #include // MiniMax-M3: MiniMax-M2 style GQA (per-head QK-norm, partial rotary) with @@ -126,68 +125,6 @@ public: int64_t nblk; }; -// pooled score of a block with no visible token: -inf from the mask, or -FLT_MAX from the -// max-pool identity when every element of the block is -inf -static inline bool msa_score_masked(float x) { return x <= -1e30f; } - -// MSA block selection (batch regime) -// CPU custom op, the token-level expansion and the combination with the causal mask happen on the GPU. -static void msa_block_mask_op(struct ggml_tensor * dst, int ith, int nth, void * userdata) { - const struct ggml_tensor * bs = dst->src[0]; - const struct ggml_tensor * bias = dst->src[1]; - const msa_params * p = (const msa_params *) userdata; - - const int nblk = (int) bs->ne[0]; - const int Hd = (int) bs->ne[1]; - const int S = (int) bs->ne[2]; - - GGML_ASSERT(bs->type == GGML_TYPE_F32 && ggml_is_contiguous(bs)); - GGML_ASSERT(bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias)); - GGML_ASSERT(dst->type == GGML_TYPE_F16 && ggml_is_contiguous(dst)); - GGML_ASSERT(dst->ne[0] == nblk && dst->ne[1] == S && dst->ne[2] == Hd); - GGML_ASSERT(bias->ne[0] == nblk && bias->ne[1] == S); - - const int topk = p->topk_blocks < nblk ? p->topk_blocks : nblk; - - const ggml_fp16_t f16_zero = ggml_fp32_to_fp16(0.0f); - const ggml_fp16_t f16_ninf = ggml_fp32_to_fp16(-INFINITY); - - std::vector rank(nblk); - std::vector valid(nblk); - std::vector ord(nblk); - - ggml_fp16_t * out = (ggml_fp16_t *) dst->data; - - for (int i = ith; i < S; i += nth) { - const float * bias_col = (const float *) bias->data + (size_t) i * nblk; - for (int h = 0; h < Hd; ++h) { - const float * bs_col = (const float *) bs->data + ((size_t) i * Hd + h) * nblk; - - for (int bk = 0; bk < nblk; ++bk) { - // a block is selectable if it has a visible token or is locally forced - valid[bk] = !msa_score_masked(bs_col[bk]) || bias_col[bk] > 0.0f; - rank [bk] = bias_col[bk] > 0.0f ? bias_col[bk] : bs_col[bk]; - ord [bk] = bk; - } - - std::partial_sort(ord.begin(), ord.begin() + topk, ord.end(), - [&](int a, int b) { return rank[a] > rank[b]; }); - - ggml_fp16_t * dst_col = out + ((size_t) h * S + i) * nblk; - for (int bk = 0; bk < nblk; ++bk) { - dst_col[bk] = f16_ninf; - } - for (int t = 0; t < topk; ++t) { - const int bk = ord[t]; - if (!valid[bk]) { - break; // sorted desc: first invalid -> fewer than topk selectable blocks - } - dst_col[bk] = f16_zero; - } - } - } -} - // One FA call for all GQA groups (and at multi-stream decode, all streams) by mapping them onto the FA sequence dim (ne[3]) ggml_tensor * llama_model_minimax_m3::graph::build_attn_msa_fa( ggml_tensor * q_cur, // [D, HQ, T] @@ -433,8 +370,8 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_ msa_mf->nb[1], msa_mf->nb[1], st*msa_mf->nb[3]); ggml_tensor * km_s = ggml_view_3d(ctx0, msa_kqm, n_kv, n_tps, 1, msa_kqm->nb[1], msa_kqm->nb[3], st*msa_kqm->nb[3]); - ggml_tensor * bias_s = ggml_view_2d(ctx0, msa_loc->bias, nblk, n_tps, - msa_loc->bias->nb[1], st*n_tps*msa_loc->bias->nb[1]); + ggml_tensor * bias_s = ggml_view_3d(ctx0, msa_loc->bias, nblk, 1, n_tps, + msa_loc->bias->nb[1], msa_loc->bias->nb[1], st*n_tps*msa_loc->bias->nb[1]); ggml_tensor * q_s = ggml_view_3d(ctx0, Qcur, D, n_head, n_tps, Qcur->nb[1], Qcur->nb[2], st*n_tps*Qcur->nb[2]); ggml_tensor * k_s = ggml_view_4d(ctx0, k, D, HKV, n_kv, 1, @@ -453,15 +390,27 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_ ggml_tensor * bs = ggml_pool_2d(ctx0, sc, GGML_OP_POOL_MAX, blk, 1, blk, 1, 0, 0); cb(bs, "msa_bs", il); - // block-level 0/-inf keep mask on the CPU, tiny transfer - ggml_tensor * srcs[2] = { bs, bias_s }; - ggml_tensor * bm = ggml_custom_4d(ctx0, GGML_TYPE_F16, - nblk, n_tps, Hd, 1, - srcs, 2, msa_block_mask_op, GGML_N_TASKS_MAX, - const_cast(&mm.msa_p)); + // bias the scores so locally-forced blocks always rank first + ggml_tensor * bsf = ggml_add(ctx0, bs, bias_s); // [nblk, Hd, n_tps] + cb(bsf, "msa_bsf", il); + + ggml_tensor * idx = ggml_top_k(ctx0, bsf, K); // [K, Hd, n_tps] i32 + + ggml_tensor * ninf = ggml_cast(ctx0, + ggml_scale_bias(ctx0, bias_s, 0.0f, -1e30f), + GGML_TYPE_F16); // [nblk, 1, n_tps] + ninf = ggml_repeat_4d(ctx0, ninf, nblk, Hd, n_tps, 1); + ggml_tensor * zero = ggml_scale(ctx0, + ggml_cast(ctx0, idx, GGML_TYPE_F32), 0.0f); + ggml_tensor * bm = ggml_set_rows(ctx0, + ggml_reshape_3d(ctx0, ninf, 1, nblk, Hd*n_tps), + ggml_reshape_3d(ctx0, zero, 1, K, Hd*n_tps), + ggml_reshape_2d(ctx0, idx, K, Hd*n_tps)); + bm = ggml_reshape_3d(ctx0, bm, nblk, Hd, n_tps); + bm = ggml_cont(ctx0, ggml_permute(ctx0, bm, 0, 2, 1, 3)); // [nblk, n_tps, Hd] cb(bm, "msa_block_mask", il); - // expand block -> token granularity on the GPU (j = bk*blk + t), + // expand block -> token granularity (j = bk*blk + t), // then combine with the causal mask in place ggml_tensor * bmx = ggml_repeat_4d(ctx0, ggml_reshape_3d(ctx0, bm, 1, nblk, n_tps*Hd), diff --git a/src/models/models.h b/src/models/models.h index c73136f3b..bb372ece8 100644 --- a/src/models/models.h +++ b/src/models/models.h @@ -2127,6 +2127,10 @@ struct llama_model_mimo2 : public llama_model_base { graph(const llama_model & model, const llm_graph_params & params); }; + struct graph_mtp : public llm_graph_context { + graph_mtp(const llama_model & model, const llm_graph_params & params); + }; + std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; }; diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index 146eabce2..fec2b0180 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -33,7 +33,7 @@ enum resize_algo { RESIZE_ALGO_BILINEAR, // stretch to target resolution RESIZE_ALGO_BICUBIC, // center-crop when aspect ratio doesn't match RESIZE_ALGO_BICUBIC_PILLOW, - // RESIZE_ALGO_LANCZOS, // TODO + RESIZE_ALGO_LANCZOS, }; // Padding style for img_tool::resize diff --git a/tools/mtmd/mtmd-image.cpp b/tools/mtmd/mtmd-image.cpp index 36cd463b2..72d35fce6 100644 --- a/tools/mtmd/mtmd-image.cpp +++ b/tools/mtmd/mtmd-image.cpp @@ -68,6 +68,9 @@ struct img_tool { case RESIZE_ALGO_BICUBIC_PILLOW: resize_bicubic_pillow(src, dst, target_resolution.width, target_resolution.height); break; + case RESIZE_ALGO_LANCZOS: + resize_lanczos_pillow(src, dst, target_resolution.width, target_resolution.height); + break; default: throw std::runtime_error("Unsupported resize algorithm"); } @@ -97,6 +100,9 @@ struct img_tool { case RESIZE_ALGO_BICUBIC_PILLOW: resize_bicubic_pillow(src, resized_image, new_width, new_height); break; + case RESIZE_ALGO_LANCZOS: + resize_lanczos_pillow(src, resized_image, new_width, new_height); + break; default: throw std::runtime_error("Unsupported resize algorithm"); } @@ -337,22 +343,50 @@ private: } } - // Bicubic resize function using Pillow's ImagingResample algorithm + // Pillow-compatible separable resampling (Bicubic and Lanczos) // Adapted from https://github.com/python-pillow/Pillow/blob/main/src/libImaging/Resample.c // - // Key Difference with resize_bicubic: - // 1. Uses separable filtering: horizontal pass followed by vertical pass + // Key properties: + // 1. Separable filtering: horizontal pass followed by vertical pass // 2. Pre-computes normalized filter coefficients for each output pixel - // 3. Applies convolution using fixed-point integer arithmetic for performance + // 3. Fixed-point integer arithmetic (22 fractional bits) for speed and determinism static bool resize_bicubic_pillow(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) { + return resize_pillow(img, dst, target_width, target_height, /*use_lanczos=*/false); + } + + // Lanczos-3 (support radius 3), matches Pillow's Image.LANCZOS + static bool resize_lanczos_pillow(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) { + return resize_pillow(img, dst, target_width, target_height, /*use_lanczos=*/true); + } + + static bool resize_pillow( + const clip_image_u8 & img, + clip_image_u8 & dst, + int target_width, + int target_height, + bool use_lanczos) { // Fixed-point precision: 22 bits = 32 (int32_t) - 8 (uint8_t pixels) - 2 (headroom for accumulation) // This allows encoding fractional weights as integers: weight * 2^22 const int PRECISION_BITS = 32 - 8 - 2; - // Bicubic filter function with a = -0.5 (Note that GGML/PyTorch takes a = -0.75) + // Resample filter: Lanczos-3 (support [-3, 3]) or bicubic with a = -0.5 (support [-2, 2]) + // Note: GGML/PyTorch bicubic uses a = -0.75, Pillow uses a = -0.5 // Returns filter weight for distance x from pixel center - // Support: [-2, 2], meaning the filter influences pixels within 2 units of distance - auto bicubic_filter = [](double x) -> double { + auto resample_filter = [use_lanczos](double x) -> double { + if (use_lanczos) { + if (-3.0 <= x && x < 3.0) { + auto sinc = [](double v) { + if (v == 0.0) { + return 1.0; + } + const double pi_v = v * 3.141592653589793238462643383279502884; + return std::sin(pi_v) / pi_v; + }; + return sinc(x) * sinc(x / 3.0); + } + return 0.0; + } + constexpr double a = -0.5; if (x < 0.0) { x = -x; @@ -366,8 +400,8 @@ private: return 0.0; // Zero outside [-2, 2] }; - // Filter support radius: bicubic extends 2 pixels in each direction - constexpr double filter_support = 2.0; + // Filter support radius: 2 for bicubic, 3 for lanczos + const double filter_support = use_lanczos ? 3.0 : 2.0; // Clipping function for 8-bit values auto clip8 = [](int val) -> uint8_t { @@ -434,7 +468,7 @@ private: // Compute filter weights for each contributing input pixel for (x = 0; x < xmax; x++) { // Distance from input pixel center to output pixel center in input space - double w = bicubic_filter((x + xmin - center + 0.5) * ss); + double w = resample_filter((x + xmin - center + 0.5) * ss); pre_weights[xx * ksize + x] = w; ww += w; // Accumulate for normalization } @@ -463,6 +497,12 @@ private: const double fxp_scale = std::ldexp(1.0, PRECISION_BITS); // 1.0 * 2^PRECISION_BITS for (int i = 0; i < outSize * ksize; i++) { + if (use_lanczos) { + // Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice + const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5); + weights[i] = static_cast(rounded); + continue; + } double tmp_val = pre_weights[i] * fxp_scale; if (pre_weights[i] < 0) { tmp_val -= 0.5; diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 749bd9aac..4655b518e 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -78,31 +78,41 @@ struct server_batch { }; std::vector tokens; int32_t n_tokens_alloc = 0; + int32_t n_embd = 0; // track if given slot can be batched with slots already in the batch server_slot * slot_batched = nullptr; + // in embd mode, we temporarily swap out the tokens arr and restore it on clear() + bool has_embd = false; + llama_token * tokens_ptr = nullptr; + std::vector embd; + float alora_scale = -1.0f; size_t alora_disabled_id = 0; server_batch() { - batch.token = nullptr; // sentinel: uninitialized batch + batch.pos = nullptr; // sentinel: uninitialized batch } ~server_batch() { - if (batch.token != nullptr) { + if (batch.pos != nullptr) { + clear(); llama_batch_free(batch); } } - void init(int32_t n_tokens_alloc) { + void init(int32_t n_tokens_alloc, int32_t n_embd) { this->n_tokens_alloc = n_tokens_alloc; + this->n_embd = n_embd; batch = llama_batch_init(n_tokens_alloc, 0, 1); + tokens_ptr = batch.token; tokens.reserve(n_tokens_alloc); } bool add(int32_t id_slot, llama_token token, llama_pos pos, bool output) { - GGML_ASSERT(batch.token != nullptr); + GGML_ASSERT(!has_embd); // cannot mix tokens + embd in same batch + GGML_ASSERT(batch.pos != nullptr); if ((int32_t)tokens.size() >= n_tokens_alloc) { return false; } @@ -110,13 +120,30 @@ struct server_batch { return true; } + bool add(int32_t id_slot, const std::vector & embd_in, llama_pos pos, bool output) { + GGML_ASSERT(batch.pos != nullptr); + if ((int32_t)tokens.size() >= n_tokens_alloc) { + return false; + } + tokens.push_back({ id_slot, LLAMA_TOKEN_NULL, pos, output }); + has_embd = true; + embd.insert(embd.end(), embd_in.begin(), embd_in.end()); + return true; + } + void clear() { tokens.clear(); + embd.clear(); common_batch_clear(batch); slot_batched = nullptr; alora_scale = -1.0f; alora_disabled_id = 0; batch_rendered = false; + has_embd = false; + if (batch.token == nullptr) { + batch.token = tokens_ptr; + batch.embd = nullptr; + } } int32_t size() const { @@ -129,25 +156,33 @@ struct server_batch { } void render() { - GGML_ASSERT(batch.token != nullptr); + GGML_ASSERT(!batch_rendered); + GGML_ASSERT(batch.pos != nullptr); common_batch_clear(batch); for (int32_t i = 0; i < size(); i++) { const auto & t = tokens[i]; common_batch_add(batch, t.token, t.pos, { t.id_slot }, t.output); } + if (has_embd) { + batch.token = nullptr; // will be restored on clear() + batch.embd = embd.data(); + } batch_rendered = true; } llama_batch get_view(int32_t off, int32_t n_tokens) const { - GGML_ASSERT(batch.token != nullptr); + GGML_ASSERT(batch.pos != nullptr); GGML_ASSERT(batch_rendered); GGML_ASSERT(off >= 0 && off < size()); GGML_ASSERT(n_tokens > 0 && off + n_tokens <= size()); + auto * token = batch.token ? batch.token + off : nullptr; + auto * embd = batch.embd ? batch.embd + off * n_embd : nullptr; + llama_batch view = { n_tokens, - batch.token + off, - nullptr, + token, + embd, batch.pos + off, batch.n_seq_id + off, batch.seq_id + off, @@ -177,6 +212,7 @@ struct server_slot { llama_tokens spec_prompt; std::vector spec_i_batch; common_prompt_checkpoint spec_ckpt; + bool spec_is_replay = false; // TODO: move members that belong to the task (such as `generated_text`, `has_new_line`) to task_results_state // see https://github.com/ggml-org/llama.cpp/pull/18283#issuecomment-3710175837 @@ -270,6 +306,10 @@ struct server_slot { llama_token sampled; // in speculative mode, this is the last accepted token + // for TTS models, this is the embd generated from prev step, decode this to generate next hidden state + // corresponding to one token position (size = n_embd) + std::vector inp_embd; + // stats size_t n_sent_text = 0; // number of sent text character @@ -293,6 +333,8 @@ struct server_slot { void reset() { SLT_DBG(*this, "%s", "\n"); + spec_is_replay = false; + n_prompt_tokens_cache = 0; last_nl_pos = 0; @@ -378,7 +420,9 @@ struct server_slot { bool can_batch_with(server_slot & other_slot) const { GGML_ASSERT(task); - return task->type == other_slot.task->type && are_lora_equal(lora, other_slot.lora); + return task->type == other_slot.task->type + && inp_embd.size() == other_slot.inp_embd.size() + && are_lora_equal(lora, other_slot.lora); } bool has_budget(const common_params & global_params) { @@ -444,7 +488,11 @@ struct server_slot { // no speculative decoding i_batch = batch.size(); - add_ok &= batch.add(id, sampled, prompt.tokens.pos_next(), true); + if (!inp_embd.empty()) { + add_ok &= batch.add(id, inp_embd, prompt.tokens.pos_next(), true); + } else { + add_ok &= batch.add(id, sampled, prompt.tokens.pos_next(), true); + } SLT_DBG(*this, "slot decode token, id=%d, n_ctx = %d, n_tokens = %d, truncated = %d\n", sampled, n_ctx, prompt.n_tokens(), truncated); @@ -1334,7 +1382,8 @@ private: // note that n_batch can be > n_ctx (e.g. for non-causal attention models such as BERT where the KV cache is not used) { const int32_t n_batch = llama_n_batch(ctx_tgt); - batch.init(std::max(n_batch, params_base.n_parallel)); + const int32_t n_embd = llama_model_n_embd_inp(model_tgt); + batch.init(std::max(n_batch, params_base.n_parallel), n_embd); } if (params_base.cache_ram_mib != 0) { @@ -3578,6 +3627,15 @@ private: n_empty_consecutive = 0; } + // TODO @ngxson : dft model may have different n_embd than the tgt model, so we check & reject if that's the case + // this case is not currently used by any models, but may need to be supported in the future + if (spec && batch.has_embd) { + if (llama_model_n_embd_inp(model_dft) != llama_model_n_embd_inp(model_tgt)) { + SRV_ERR("%s", "unsupported batch.has_embd + spec case\n"); + throw std::runtime_error("unsupported batch.has_embd + spec case"); + } + } + const int ret = llama_decode(ctx_tgt, batch_view); metrics.on_decoded(slots); @@ -3820,6 +3878,7 @@ private: } // partial acceptance is not supported by the context -> truncate the draft and restore the state + slot.spec_is_replay = true; slot.spec_draft = std::move(accepted); const auto & ckpt = slot.spec_ckpt; @@ -3854,16 +3913,22 @@ private: const auto ids = std::move(slot.spec_draft); + size_t n_accepted = ids.size() - 1; + if (slot.spec_is_replay && n_accepted > 0) { + n_accepted--; + } + slot.spec_is_replay = false; + slot.t_token_generation = std::max(1, t_now - slot.t_start_generation) / 1e3; // update how many tokens out of those tested were accepted - slot.n_draft_accepted += ids.size() - 1; + slot.n_draft_accepted += n_accepted; slot.n_draft_verif_steps += 1; if (slot.n_accepted_per_pos.empty()) { slot.n_accepted_per_pos.resize(common_speculative_n_max(¶ms_base.speculative), 0); } - for (size_t i = 0; i < ids.size() - 1 && i < slot.n_accepted_per_pos.size(); ++i) { + for (size_t i = 0; i < n_accepted && i < slot.n_accepted_per_pos.size(); ++i) { slot.n_accepted_per_pos[i]++; } @@ -3899,7 +3964,7 @@ private: slot.print_timings_tg(); - SLT_DBG(slot, "accepted %d/%d draft tokens, new n_tokens = %d\n", (int) ids.size() - 1, (int) n_draft, slot.prompt.n_tokens()); + SLT_DBG(slot, "accepted %d/%d draft tokens, new n_tokens = %d\n", (int) n_accepted, (int) n_draft, slot.prompt.n_tokens()); }); } diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte index 905c2fe6f..f81dcf09c 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte @@ -48,6 +48,9 @@ }: Props = $props(); let dropdownOpen = $state(false); + // The system message action moves focus to the message editor, so the menu + // must not restore focus to the trigger on close + let suppressCloseAutoFocus = false; function handleMcpSettingsClick() { dropdownOpen = false; @@ -96,7 +99,16 @@ - + { + if (suppressCloseAutoFocus) { + suppressCloseAutoFocus = false; + e.preventDefault(); + } + }} + > @@ -148,7 +160,10 @@ { + suppressCloseAutoFocus = true; + onSystemPromptClick?.(); + }} > diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessage.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessage.svelte index 8e8a14ac3..b8068f790 100644 --- a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessage.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessage.svelte @@ -2,6 +2,7 @@ import { goto } from '$app/navigation'; import { getChatActionsContext, setMessageEditContext } from '$lib/contexts'; import { chatStore, pendingEditMessageId } from '$lib/stores/chat.svelte'; + import { isMobile } from '$lib/stores/viewport.svelte'; import { conversationsStore } from '$lib/stores/conversations.svelte'; import { DatabaseService } from '$lib/services/database.service'; import { SYSTEM_MESSAGE_PLACEHOLDER } from '$lib/constants'; @@ -46,7 +47,14 @@ assistantMessages: number; messageTypes: string[]; } | null>(null); - let editedContent = $derived(message.content); + // The system message placeholder must never surface as editable content; keeping + // it in the derived (not just in handleEdit) guards against prop invalidation + // reverting the override while editing + let editedContent = $derived( + message.role === MessageRole.SYSTEM && message.content === SYSTEM_MESSAGE_PLACEHOLDER + ? '' + : message.content + ); let rawEditContent = $derived.by(() => { if (message.role !== MessageRole.ASSISTANT) return undefined; @@ -265,6 +273,12 @@ chatActions.navigateToSibling(siblingId); } + // After the system message flow ends, hand focus to the main chat form + function focusMainChatForm() { + if (isMobile.current) return; + document.querySelector('.chat-screen-form-wrapper textarea')?.focus(); + } + async function handleSaveEdit() { if (message.role === MessageRole.SYSTEM) { // System messages: update in place without branching @@ -276,6 +290,8 @@ isEditing = false; if (conversationDeleted) { goto(ROUTES.START); + } else { + focusMainChatForm(); } return; } @@ -285,6 +301,7 @@ if (index !== -1) { conversationsStore.updateMessageAtIndex(index, { content: newContent }); } + focusMainChatForm(); } else if (message.role === MessageRole.USER) { const finalExtras = await getMergedExtras(); chatActions.editWithBranching(message, editedContent.trim(), finalExtras); diff --git a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenForm.svelte b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenForm.svelte index 600180742..8eb17eeae 100644 --- a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenForm.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenForm.svelte @@ -106,15 +106,23 @@ onFileRemove?.(fileId); } + // Auto-focus must not steal focus already claimed elsewhere (e.g. the system + // message editor opened just before a navigation) + function focusFormUnlessCaptured() { + const active = document.activeElement; + if (active instanceof HTMLTextAreaElement || active instanceof HTMLInputElement) return; + chatFormRef?.focus(); + } + onMount(() => { if (!isMobile.current) { - setTimeout(() => chatFormRef?.focus(), 100); + setTimeout(focusFormUnlessCaptured, 100); } }); afterNavigate((navigation) => { if (navigation?.from != null && !isMobile.current) { - setTimeout(() => chatFormRef?.focus(), 100); + setTimeout(focusFormUnlessCaptured, 100); } }); @@ -127,7 +135,7 @@ $effect(() => { if (previousIsLoading && !isLoading) { - setTimeout(() => chatFormRef?.focus(), 10); + setTimeout(focusFormUnlessCaptured, 10); } previousIsLoading = isLoading; diff --git a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenGreeting.svelte b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenGreeting.svelte index 018949aff..5b44bcf85 100644 --- a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenGreeting.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenGreeting.svelte @@ -11,7 +11,7 @@