Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/mmq.cu (CUDA)

Code before:
    if (amd_mfma_available(cc)) {
        // As of ROCM 7.0 rocblas/tensile performs very poorly on CDNA3 and hipblaslt (via ROCBLAS_USE_HIPBLASLT)
        // performs better but is currently suffering from a crash on this architecture.
        // TODO: Revisit when hipblaslt is fixed on CDNA3
        if (GGML_CUDA_CC_IS_CDNA3(cc)) {
            return true;
        }
        if (n_experts > 64 || ne11 <= 128) {
            return true;
        }
        if (type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 || type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1) {
            return true;
        }
        if (ne11 <= 256 && (type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K)) {
            return true;
        }
        return false;
    }

    if (amd_wmma_available(cc)) {
        // High expert counts almost always better on MMQ
        // due to a large amount of graph splits
        // https://github.com/ggml-org/llama.cpp/pull/18202
        if (n_experts >= 64) {
            return true;
        }

        switch (type) {
            // These quants are really bad on MMQ
            case GGML_TYPE_Q2_K:
            case GGML_TYPE_Q6_K:
            // These quants are usually worse but not always
            case GGML_TYPE_IQ2_XS:
            case GGML_TYPE_IQ2_S:
                return ne11 <= 128;
            default:
                return true;
        }
    }

    return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;

}


Review comment (suggestion):
32 for GPT OSS 20B was also sort of a wash, but I'm assuming non-fp4 models will benefit more from the quant cases instead. There's not a lot of MOEs I can run fully in VRAM, so testing this was limited.

Code after:
    if (amd_mfma_available(cc)) {
        // As of ROCM 7.0 rocblas/tensile performs very poorly on CDNA3 and hipblaslt (via ROCBLAS_USE_HIPBLASLT)
        // performs better but is currently suffering from a crash on this architecture.
        // TODO: Revisit when hipblaslt is fixed on CDNA3
        if (GGML_CUDA_CC_IS_CDNA3(cc)) {
            return true;
        }
        if (n_experts > 64 || ne11 <= 128) {
            return true;
        }
        if (type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 || type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1) {
            return true;
        }
        if (ne11 <= 256 && (type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K)) {
            return true;
        }
        return false;
    }

    if (amd_wmma_available(cc)) {
        // RDNA 4 is consistently worse on rocblas
        // https://github.com/ggml-org/llama.cpp/pull/18537#issuecomment-3706422301
        if (GGML_CUDA_CC_IS_RDNA3(cc)) {
            // High expert counts almost always better on MMQ
            // due to a large amount of graph splits
            // https://github.com/ggml-org/llama.cpp/pull/18202
            if (n_experts >= 64) {
                return true;
            }

            switch (type) {
                // These quants are really bad on MMQ
                case GGML_TYPE_Q2_K:
                case GGML_TYPE_Q6_K:
                // These quants are usually worse but not always
                case GGML_TYPE_IQ2_XS:
                case GGML_TYPE_IQ2_S:
                    return ne11 <= 128;
                default:
                    return true;
            }
        }
        return true;
    }

    return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;

}