← 記事一覧

人工知能

DeepSeek が Huawei Ascend 向けにライブラリ 6 本をオープンソース化した。輸出規制が一度も触れなかったのは、このソフトウェア層である

DeepSeek は 9 月 30 日、Huawei の Ascend 950 向け TileLang バックエンドと、カーネル/通信ライブラリ 5 本を公開した。Huawei との共同開発であり、CUDA バックエンドの隣、言語本体へ上流マージされている。輸出政策はチップを止めた。CUDA を止めたことは一度もなく、そして置き換えが難しいのは常に CUDA の側だった。

MAI
deepseek-ai/DeepGEMM-Ascend リポジトリの GitHub ソーシャルカード。プロジェクト名と、Huawei Ascend NPU 向け行列積カーネルライブラリという説明が示されている。

DeepSeek は 9 月 30 日、Huawei の Ascend AI アクセラレータ向けに、オープンソースのソフトウェア構成要素 6 本を公開した。Huawei との共同開発であり、発表は同社自身の WeChat で行われた。中心にあるのは、計算カーネルを記述するための言語 TileLang であり、これに公式の Ascend 950 バックエンドが加わった。残る 5 本はカーネルおよび通信のライブラリで、DeepSeek がすでに Nvidia GPU 向けに公開してきた一式と、ほぼ一対一で対応している。

この対応関係こそが要点である。これはモデルの発表ではない。自社のフロンティアモデルを訓練するために使っている道具立てが、中国製シリコンを一級の対象として扱うようになった、という宣言だ。

何が公開されたのか

構成要素役割
TileLang(Ascend 950 バックエンド)Ascend 向けのネイティブコード生成、自動スケジューリング、同期を備えたカーネル記述言語
DeepGEMM-Ascend行列積カーネル
TileKernelsベクトル演算およびメモリ演算子
FlashMLA長文脈推論のためのスパースアテンション
DeepEPデバイス間のエキスパート並列通信
DeepSelect効率的なデータ選択

Ascend 950 バックエンドは、tile-ai/tilelang プロジェクトへプルリクエスト #3308 として上流マージされた。マージは 9 月 29 日、134 コミット、貢献者として SiriusNEO と LeiWang1999 が記されている。この一点は見た目より重い。これはインターネットの片隅に置かれたベンダーのフォークではない。言語本体のメインブランチに、CUDA バックエンドの隣へ着地した。しかも CUDA のスレッドモデルを借りるのではなく、カーネル起動のための独自ダイアレクトを定義している。

このバックエンドが露出させているのは、Ascend という機械のありのままの姿である。カーネルは AIC ユニット上の Cube 行列演算と AIV ユニット上の Vector 演算を一つの中で組み合わせられる。メモリは UB、L1、L0 の各階層に明示的に配置できる。ブロックスケーリングされた MXFP8 および MXFP4 の演算には、それぞれの経路が用意されている。AutoSchedule パスが依存関係を見た命令パイプライニング、マルチバッファリング、そしてコア内およびコア間の同期バリアの挿入を受け持つ。デバイスコードは Huawei の CANN ツールキットに含まれる bisheng コンパイラを通り、テンソルとストリームは PyTorch の NPU バックエンドと相互運用される。DeepSeek と Huawei はさらに、Ascend 950 を 128 個つないだスーパーノード構成を、計算と通信をまとめて最適化したものとして説明している。

性能に関する主張は DeepSeek 自身のものであり、DeepGEMM-Ascend リポジトリに掲載されたもので、外部の誰かが検証したわけではない。README は、密行列積がハードウェア上限の最大 99.8% に達し、FP8 推論が最大 861 TFLOPS、MegaMoE 演算が最大 846.3 TFLOPS、MQA logits が FIX パイプを 99% の利用率で飽和させたと報告している。いずれも CANN 9.20 を用いた Ascend 950 シリーズ上での数字だ。誰かが再現するまでは、ベンダー数値として扱うのが妥当である。プルリクエスト側の評価はより控えめで、そしてより有用だ。bfloat16 の行列積は PyTorch の既存 NPU 経路と同等かそれ以上、FP8 のキャストは帯域幅の上限近くに達している。

一度も課されなかった規制

米国の輸出政策は 3 年をかけてチップの販売を制限してきた。CUDA を制限したことは一度もない。そして置き換えが難しいのは CUDA の側である。Nvidia の優位は、アクセラレータが速いことだけではない。20 年分のカーネル、ライブラリ、フレームワーク統合、そしてエンジニアの手の記憶が、その programming model を前提にしていることだ。Ascend の部品を倉庫いっぱいに渡された中国のラボは、動かしたいソフトウェアが別物のために書かれている、という問題に依然として直面する。

今回の公開が埋めようとしているのはこの隙間であり、それを埋められる唯一の立場から埋めている。Huawei はしばらく前から CANN をオープンソース化し、PyTorch、Triton、vLLM へつなぎ込んできた。しかしチップベンダーが自社の道具立てを書くのは、ベンダーが売り込みをしているにすぎない。DeepSeek が書くというのは、その国で最も能力のあるモデルラボが、この道具立てはフロンティア訓練に耐えると保証することだ。同社によれば、V4 系列の訓練で使われる TileLang のオペレータはすべて、高性能な Ascend 実装を持つようになったという。DeepSeek は 4 月に V4 を Huawei のハードウェアへ移植している。内モンゴルに Ascend アクセラレータ約 16 万基を収めるデータセンターを建設中とも報じられている。今日公開されたライブラリは、その数字に意味を持たせるためのものだ。

これが何ではないか

ライブラリ 6 本とコンパイラバックエンド 1 つは、CUDA ではない。バックエンドが狙うのは Ascend 950 に限られ、それ以前の A2 および A3 は依然として別のコミュニティプロジェクトが支えている。PTO のコード生成経路と仮想マシン組み込み命令は、上流化の過程で落とされた。カーネル群が覆うのは DeepSeek 自身のワークロードが必要とするもの——アテンション、行列積、Mixture-of-Experts の通信、データ選択——であり、これは狭く、そして極めてよく選ばれた切り口であって、汎用のエコシステムではない。自分のモデルが DeepSeek のそれに似ていない者は、おおむね自力でやることになる。

とはいえ、注目すべきは方向である。チップはいずれ手に入るものだった。歩留まりは上がり、設計は反復し、在庫は積まれる。ソフトウェア層は、そうする動機を特に持たない人々による、忍耐のいる地味な作業を 10 年分は要するように見えた部分だった。DeepSeek には動機があり、エンジニアがおり、そしてその作業を、中国の内外の誰でも拾い上げられる寛容なライセンスの下で公開しながら進めている。輸出規制は出荷を止められる。マージされたプルリクエストを取り消すことはできない。

Sources: Bloomberg: DeepSeek Unveils Huawei AI Chip Tools That May Replace Nvidia's, South China Morning Post: China's DeepSeek open-sources tools to help Huawei chips supplant Nvidia in AI, deepseek-ai/DeepGEMM-Ascend on GitHub, tile-ai/tilelang pull request #3308: Introduce Ascend 950 backend, tile-ai/tilelang-ascend on GitHub

続けて読む