DeepSeekは、Huaweiの950シリーズNPUを使う開発者向けに、DeepGEMMカーネルライブラリのAscend版を公開しました。リポジトリのREADMEによると、初回リリース日は2026年9月30日で、Pythonパッケージのバージョンは0.1.0です。このライブラリは、同ハードウェアで行列乗算や関連するモデル演算を行うためのオープンソースの出発点をチームに提供します。現時点での性能データはDeepSeek自身のテストによるものです。

この公開は、Ascendのソフトウェアスタックを評価するエンジニアにとって重要です。カーネルのソースコード、使い慣れたDeepGEMMのエントリーポイント、拡張機能のビルド手順が文書化されているためです。モデル全体の性能と運用コストは、対象システムで実際のワークロードを試して確認する必要があります。

リリースに含まれるもの

DeepGEMMは、AIモデルで繰り返し使われる行列演算のライブラリです。Ascend READMEには、BF16、FP8、FP4のGEMMカーネル、混合エキスパートモデル向けのグループ演算、MQA logits、融合型MegaMoE演算、mHC prenormカーネルが列挙されています。 Pythonパッケージのエクスポートには、denseおよびgrouped GEMM関数、MQA関数、スケーリング係数の配置を扱うヘルパー、JITコンパイルとAI coreの使用制御が含まれます。リポジトリのライセンスはMITです。

DeepSeekによると、Ascend移植版はAPIレベルで上流版DeepGEMMと完全に互換性があります。同じdeep_gemmパッケージ名と公開インターフェースの多くはコード移行を容易にする可能性がありますが、READMEは重要なデータ配置の違いを示しています。Ascendでは、K方向に並ぶUE8M0スケーリング係数の各ペアを16ビット値にまとめ、int16に格納し、パックした値をMN-major順で保存します。パッケージはスケーリング係数を必要な配置に変換する関数を公開しています。量子化入力を使うアプリケーションは、適切なデータを用意し、自ら結果を検証する必要があります。互換性は保守担当者の主張であり、BIG CHANGEによる移行テストの結果ではありません。

開発者が試すために必要なもの

文書に記載された要件は、Ascend NPU、CANN 9.20とbishengおよびld.lldの各ツール、torch_npu、Python 3.10以降、C++20に対応するコンパイラーと標準ライブラリ<format>です。保守担当者によると、開発と検証はAscend 950シリーズで行われました。tilelangはHC prenormカーネルのパッケージ依存関係として宣言されています。tree-sitterとtree-sitter-cppは、ソースからビルドする際に型スタブを生成するために必要です。READMEはtorch_npuのバージョンを固定しておらず、ほかのAscendデバイスでの検証も示していません。

READMEは、サブモジュールを含めてリポジトリをクローンし、develop.shでC++拡張をビルドするよう案内しています。または、チェックアウトしたソースからpip install . --no-build-isolationでインストールできます。そのビルド設定はtorchとtorch_npuをインポートし、ASCEND_HOME_PATHまたはASCEND_TOOLKIT_HOMEを使ってAscendツールキットを探し、AscendとTorch NPUのライブラリにリンクします。したがって、カーネルの評価には適切なハードウェアとツールチェーンが必要です。リポジトリには演算テストも含まれます。BIG CHANGEは文書とコードを確認しましたが、ビルドやカーネルの実行は行っていません。

ベンチマークで分かる範囲

DeepSeekは、CANN 9.20を搭載したAscend 950DTで、bench_msprofを使い、L2キャッシュが冷えた状態で測定したと報告しています。密なBF16行列乗算の形状M=4096、N=7168、K=16384では、性能表に2,229.9マイクロ秒、431 TFLOPS、同表に示されたハードウェア上限の99.8%と記載されています。同じ形状のFP8×FP8では、1,117.6マイクロ秒、861 TFLOPSと報告されています。これらは指定された型と形状におけるカーネルの測定値であり、モデルのスループットではありません。表には、それぞれ異なる構成でのgrouped GEMM、MQA、MegaMoE、prenormの結果もあります。MegaMoEの数値は、8ランク、8-wayのexpert parallelism、top-kルーティング6、共有expert 1つの条件での平均です。

そのテストコードはGEMMの出力を生成した参照値と照合し、対応するdenseケースではACLNNによるクロスチェックを行います。一部の形状やレシピでは、この補足的なACLNN比較を明示的に省略します。公開コードとテストにより主張を調べることはできますが、BIG CHANGEはREADMEの数値を独自に再現していません。この公開には、別のアクセラレータスタックとの同条件ベンチマーク、コスト比較、エンドツーエンドのサービス結果はありません。

主な変化

DeepSeekは、Ascend 950シリーズ向けにDeepGEMM Ascendのカーネルコードを公開しました。

開発者は、このハードウェア向けの使い慣れたGEMMや関連するモデル演算を確認し、ビルドできます。公開資料にはCANN 9.20とAscend固有のスケーリング係数配置が記載されています。

これにより、Ascend 950を使うチームは、自社モデルの形状に合わせて評価できるカーネルコードを入手できます。DeepSeekは、1つのAscend 950DTテスト構成で上限に近い結果を報告していますが、BIG CHANGEはカーネルを実行していません。

アプリケーション全体の性能、ほかのAscendデバイスへの対応、運用コストは今後の確認が必要です。

出典と関連資料

  • DeepGEMM AscendリポジトリとREADME：9月30日のリリースノート、前提条件、インターフェースに関する主張、プロジェクト作成のベンチマーク表。READMEはプロジェクト保守担当者による一次資料であり、独立検証ではありません。
  • パッケージのエクスポートとビルド設定：初回コミット時点のバージョン、公開Python関数、依存関係の宣言、ソースビルドのリンク設定。
  • GEMMテストハーネス：リポジトリが出力を検証し、プロファイラーを呼び出す方法と、ACLNN比較を省略する条件を説明しています。これらのファイルは公表値を独立して再現したものではありません。