model: support Longcat-Flash (need testing)

[ngxson]

Overview

Support https://huggingface.co/meituan-longcat/LongCat-Flash-Chat

Next: impl the ngram model: #19167

LongCat-Flash-Chat uses MLA + "zero-computing experts" (see below). I tested with a slice of the bigger model (generated via this script, which extracts 2 layers and 64 experts from the original weight)

📊 NMSE Check for Model Comparison
NMSE:                         6.904346e-05
Max Absolute Error:           0.117372
NMSE (dB):                    -41.61 dB
✅ Excellent match
📚 NMSE BENCHMARKS
< 1e-4:  Excellent (typical for good conversions)
✅ RESULT: PASS (NMSE = 6.90e-05)

Additional information

This model is different from other MoE models, it use something called "zero-computing experts" (ref: link to paper)

Ideally #26631 is needed to make it correctly skip computation

The current impl support this by (provided that model has N+M expert, with N=real exp and M=zero exp):

• Inject a N+1 expert to up/down/gate; this exp is all zero, so output will be vector of all 0
• On cgraph, clamp real exp routing to N+1, so everything > N will be routed to the injected 0 exp in step 1
• For the rest of exp, make a mask of M "skip exp" by shift+clamp the router indexes

Old comment

To illustrate what's the problem, I will take an example of how a normal MoE FFN work:

• Calculate expert probs using the router
• Sort & get top_k experts
• Do FFN gate/up/down with the selected top_k experts, this is done via ggml_mul_mat_id
• Weighted sum the output

This means we spend the the same amount of computation for each token, proportionally to n_expert_used

However, with longcat-flash:

• After top_k expert, ONLY experts with ID < n_zero_experts will go through FFN; for the rest, they skip the FFN altogether
• This makes the amount of computation to be varied token-by-token. For example: one token can use n FFN expert, while another can use n-1, another can use 0 FFN expert (in other words, skipping the MoE altogether)

Apart from the weird MoE, the model has double block architecture, meaning there are 2 attentions and 2 FFNs per layer. Upon converting to GGUF, we convert it to a model of 2 * n_layer, which make the implementation much easier.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: Most of the code written by AI; core changes are impl by me (adapted from 1fa084e, which is human-coded)

[ggerganov]

Huh, interesting. Likely need to extend ggml_mul_mat_id with more general alpha_i*A^B + beta_i*C and implement special casing in the kernels for alpha_i=0 to skip the compute.

[hebangwen]

Huh, interesting. Likely need to extend ggml_mul_mat_id with more general alpha_i*A^B + beta_i*C and implement special casing in the kernels for alpha_i=0 to skip the compute.

Hello, I'm also paying attention to the adaptation of the loncat-flash model. Suppose B is the expert weight. If we change mul_mat_ids to alpha * A^B + beta * C, since the weights of zero-compute experts are not saved and the experts activated for each token are different, there might be an interleaved expert activation scenario. In this case, we need to determine for each token whether the activated expert is a zero-compute expert. Would this have an impact on performance?

If we bypass the requirement of ggml_mul_mat_ids to arrange in token order and add a new operator that rearranges the activations in expert order, like torch._grouped_gemm, would this provide better adaptability? In this case, we would only need to determine whether the currently computed expert ID is valid or whether it is a zero-compute expert. See #18369

[ggerganov]

@hebangwen Not sure I follow, but I think simply setting the coefficient like this should work:

# normal expert
alpha_i = 1.0f
beta_i  = 0.0f

# zero-compute expert
alpha_i = 0.0f
beta_i  = 1.0f

And the matrix C is just the MoE input (i.e. x_t from the paper).

[ngxson]

It can be simpler to explain the mul_mat_id in simple terms like this:

• A normal mul_mat(A, B) calculates A @ B = C, simple.
• Now, instead of having just a single B, we have B as a stack of multiple matrices: B = (B0, B1, B2, ..., Bn)
• For MoE, we want to mul_mat A with a subset of B, something like: A @ (B2, B8, ...)
• So mul_mat_id takes an extra param, the indexes of elements in B to be used: mul_mat_id(A, B, (2, 8)) --> (A @ B2, A @ B8)

As @ggerganov suggested, I imagine the mul_mat_id will now take an extra alpha, beta params:

idx = (2, 8)
alpha = (0.3, 0.0)
mul_mat_id(A, B, idx, alpha) --> (A @ B2 * 0.3, A @ B8 * 0.0)

In the example above, computation for A @ B8 will be skipped as its beta value is 0.0

However, one issue is that the router weight alpha is non-zero. Indeed, it depends on the top_k operation to sort out the activated experts. So, I think just adding a check like if (idx >= B->ne[2]) could be enough? So if B only has 4 experts and the experts ID=5 is accessed, then it's out-of-bound; we skip the mul_mat in such case.

However, yet another problem, even when the idea above is implemented: The output dim of mul_mat_id will not be the same as the input, and more importantly, there are also other ops like mul or gelu in-between FFN gate/up/down. We can resolve this by assuming that output of the skipped mul mat will be all 0.0, but I think that's not a generic solution.

For the calculation of the beta_i*C term, it will also be a bit tricky, as we need to filter out only the coefficients (aka router weights) that correspond to the zero-computing experts. Something like torch.where could be necessary, but that's another rabbit hole I think

[ngxson]

Seems like quite more works than I initially thought, so I think we should re-consider if this worth implementing. Currently, only longcat-flash family using this technique, so it can be quite risky to too many infrastructure to support it.

[ggerganov]

Yes, seems more complicated. Let's reconsider later in case this architecture shows any promise.

[ngxson]

This PR should be ready for testing now. I tested with a very small (8B params) sub-model extracted from the original one.

Appreciate if someone can test with the bigger model.

[ServeurpersoCom]

Oh, how I’d love to be able to test it :x

[ngxson]

A bit more digging, seems like CUDA doesn't support duplicated expert indexes, that need to be fixed first...

Testing to see if other backend has the same problem: #26657

[AesSedai]

I've uploaded an IQ1_S (phew, that PPL and KLD look horrible) and was able to get it to infer with this patch (warning, AI slop ahead): cuda-mmid-duplicate-expert-ids.patch

The repo includes the imatrix I made from the Q8_0, and the BF16 logits and the wiki.test.raw for others to run PPL / KLD if they want.

[CISC]

Maybe at least add conversion of the MTP layer?

[ngxson]

Ok some conversion fixes and NMSE is reduced from 6.9e-05 down to 6.9e-08, perfect numerical match now

NMSE:                         6.991135e-08
Max Absolute Error:           0.003572
Mean Absolute Error:          0.000606
✅ RESULT: PASS (NMSE = 6.99e-08)

@AesSedai appreciate if you can re-convert the GGUF, thanks :-)

[AesSedai]

@ngxson the imatrix crashes with this cuda error:

4.28.282.343 I compute_imatrix: tokenizing the input ..
4.28.781.446 I compute_imatrix: tokenization took 499.094 ms
4.28.782.153 I compute_imatrix: computing over 53 chunks, n_ctx=8192, batch_size=8192, n_seq=1
/home/jarvis/development/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:106: CUDA error
4.31.301.692 E CUDA error: an illegal memory access was encountered
4.31.301.702 E   current device: 0, in function ggml_backend_cuda_synchronize at /home/jarvis/development/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2499
4.31.301.703 E   cudaStreamSynchronize(cuda_ctx->stream())
/home/jarvis/development/llama.cpp/build/bin/libggml-base.so.0(+0x4825) [0x7fd5abb72825]
/home/jarvis/development/llama.cpp/build/bin/libggml-base.so.0(ggml_print_backtrace+0x1eb) [0x7fd5abb72bfb]
/home/jarvis/development/llama.cpp/build/bin/libggml-base.so.0(ggml_abort+0x11f) [0x7fd5abb72d7f]
/home/jarvis/development/llama.cpp/build/bin/libggml-cuda.so.0(+0x1f9183) [0x7fd5a95f9183]
/home/jarvis/development/llama.cpp/build/bin/libggml-cuda.so.0(+0x200708) [0x7fd5a9600708]
/home/jarvis/development/llama.cpp/build/bin/libggml-base.so.0(ggml_backend_sched_graph_compute_async+0x2bc) [0x7fd5abb8f91c]
/home/jarvis/development/llama.cpp/build/bin/libllama.so.0(_ZN13llama_context13graph_computeEP11ggml_cgraphb+0xa0) [0x7fd5abe47ab0]
/home/jarvis/development/llama.cpp/build/bin/libllama.so.0(_ZN13llama_context14process_ubatchERK12llama_ubatch14llm_graph_typeP22llama_memory_context_iR11ggml_status+0xeb) [0x7fd5abe4b75b]
/home/jarvis/development/llama.cpp/build/bin/libllama.so.0(_ZN13llama_context6decodeERK11llama_batch+0x398) [0x7fd5abe52178]
/home/jarvis/development/llama.cpp/build/bin/libllama.so.0(llama_decode+0xe) [0x7fd5abe53d4e]
./build/bin/llama-imatrix() [0x403993]
/lib64/libc.so.6(+0x35b5) [0x7fd5a2e105b5]
/lib64/libc.so.6(__libc_start_main+0x88) [0x7fd5a2e10668]
./build/bin/llama-imatrix() [0x404c45]

[ngxson]

@AesSedai you will need to apply #26294 on top of this PR

[AesSedai]

I updated the quant in the repo a few days ago, but it still needed the patch I posted earlier.

[ngxson]

@AesSedai thanks! I will go back to this a bit later this month

it's currently blocked by the issue of duplicated experts on CUDA, can be merged once that's fixed. in any cases, GGUF is fine and won't need to be reconverted