DeepSeek publie les kernels DeepGEMM pour Huawei Ascend 950

DeepSeek a publié une version Ascend de sa bibliothèque de kernels DeepGEMM pour les développeurs qui utilisent les NPU de la série 950 de Huawei. Le README du dépôt date la première publication du 30 septembre 2026 ; le paquet Python s’identifie comme version 0.1.0. Il offre aux équipes une base open source pour la multiplication de matrices et d’autres opérations de modèles sur ce matériel. À ce jour, les chiffres de performance proviennent des propres tests de DeepSeek.

Cette publication intéresse les ingénieurs qui évaluent une pile logicielle Ascend : elle met à disposition le code source des kernels, des points d’entrée DeepGEMM familiers et une méthode documentée pour compiler l’extension. Les performances d’un modèle complet et le coût d’exploitation doivent encore être mesurés avec les charges de travail prévues sur le système cible.

Ce que comprend la publication

DeepGEMM est une bibliothèque pour les opérations matricielles récurrentes dans les modèles d’IA. Le README Ascend répertorie des kernels GEMM BF16, FP8 et FP4, des opérations groupées pour les modèles à mélange d’experts, des logits MQA, une opération MegaMoE fusionnée et un kernel mHC prenorm. Ses exports du paquet Python présentent des fonctions GEMM denses et groupées, des fonctions MQA, des outils de disposition des facteurs d’échelle et des commandes pour la compilation JIT et l’utilisation d’AI core. Le dépôt est sous licence MIT.

DeepSeek indique que le port Ascend est entièrement compatible au niveau de l’API avec son DeepGEMM amont. Le même nom de paquet, deep_gemm deep_gemm, et une grande partie de l’interface publique peuvent faciliter la migration du code, mais le README signale une différence importante de disposition des données : Ascend regroupe par paires les facteurs d’échelle UE8M0 le long de K dans une valeur de 16 bits int16 et stocke les valeurs regroupées selon l’ordre MN-major. Le paquet exporte des fonctions qui convertissent les facteurs d’échelle vers cette disposition requise. Une application utilisant des entrées quantifiées doit toujours fournir des données adaptées et valider ses propres résultats ; la compatibilité est une affirmation des mainteneurs, pas un test de migration réalisé par BIG CHANGE.

Ce qu’il faut pour essayer la bibliothèque

Les prérequis documentés sont un NPU Ascend, CANN 9.20 avec bisheng et ld.lld comme outils, torch_npu, Python 3.10 ou une version ultérieure et un compilateur et une bibliothèque standard compatibles avec C++20 <format>. Les mainteneurs indiquent que le développement et la validation ont eu lieu sur la série Ascend 950. tilelang est déclaré comme dépendance du paquet pour le kernel HC prenorm ; tree-sitter et tree-sitter-cpp sont nécessaires pour générer les stubs de types lors d’une compilation depuis les sources. Le README ne fixe aucune version de torch_npu et n’établit pas de validation sur d’autres appareils Ascend.

Le README indique de cloner le dépôt avec ses sous-modules puis de compiler l’extension C++ avec develop.sh, ou de l’installer depuis les sources récupérées avec pip install . --no-build-isolation. Sa configuration de compilation importe torch et torch_npu, trouve la boîte à outils Ascend via ASCEND_HOME_PATH ou ASCEND_TOOLKIT_HOME, et lie les bibliothèques Ascend et Torch NPU. L’évaluation des kernels exige donc le matériel et la chaîne d’outils appropriés. Le dépôt comprend des tests des opérations. BIG CHANGE a examiné la documentation et le code, mais n’a exécuté ni la compilation ni les kernels.

Jusqu’où vont les benchmarks

DeepSeek rapporte des mesures sur un Ascend 950DT avec CANN 9.20, en utilisant bench_msprof avec un cache L2 froid. Pour une forme de multiplication matricielle BF16 dense — M=4096, N=7168 et K=16384 — son tableau de performances indique 2 229,9 microsecondes, 431 TFLOPS et 99,8 % de la limite matérielle indiquée dans ce tableau. Pour la même forme, il rapporte 1 117,6 microsecondes et 861 TFLOPS en FP8×FP8. Il s’agit de mesures de kernels pour des types et des formes précis, pas du débit mesuré d’un modèle. Le tableau donne aussi des résultats GEMM groupés, MQA, MegaMoE et prenorm avec leurs propres configurations ; les valeurs MegaMoE sont une moyenne sur huit ranks, un parallélisme d’experts de huit, un routage top-k de six et un expert partagé.

Le code de test compare les sorties GEMM à des références générées et effectue une vérification croisée ACLNN pour les cas denses pris en charge. Il omet explicitement cette comparaison ACLNN supplémentaire pour certaines formes et recettes. Le code et les tests publiés permettent d’examiner les affirmations, mais BIG CHANGE n’a pas reproduit indépendamment les chiffres du README. La publication ne fournit ni benchmark comparable avec une autre pile d’accélérateurs, ni comparaison de coûts, ni résultat de service de bout en bout.

Le principal changement

DeepSeek a publié le code des kernels DeepGEMM Ascend pour la série Ascend 950.

Les développeurs peuvent examiner et compiler des opérations GEMM familières et des opérations de modèle connexes pour ce matériel. La publication documente CANN 9.20 et une disposition des facteurs d’échelle propre à Ascend.

Les équipes équipées d’Ascend 950 disposent ainsi de code de kernels à évaluer sur les formes de leurs propres modèles. DeepSeek rapporte des résultats proches de la limite dans une configuration de test Ascend 950DT ; BIG CHANGE n’a pas exécuté ces kernels.

Les performances d’applications complètes, la prise en charge d’autres appareils Ascend et le coût d’exploitation restent à établir.

Sources et lectures complémentaires

• Dépôt et README DeepGEMM Ascend : note de publication du 30 septembre, prérequis, affirmations sur les interfaces et tableaux de benchmarks produits par le projet. Le README est une source primaire des mainteneurs du projet, pas une validation indépendante.
• Exports du paquet et configuration de compilation : version, fonctions Python publiques, dépendances déclarées et liaison de compilation depuis les sources au commit initial.
• Banc de test GEMM : fonctionnement des vérifications de sortie et appel au profileur du dépôt, y compris les cas où la comparaison ACLNN est omise. Ces fichiers ne reproduisent pas indépendamment les chiffres publiés.