DeepSeekがHuawei向けAI計算基盤を公開 — 「脱CUDA」が実は「依存の乗り換え」である理由

この記事は時点のニュースです。現在の提供状況とは異なる場合があります。誤りの訂正は随時行います。


共通化されるのは「書く層」だけ TileLang(カーネル記述の共通言語) CUDA SM70〜SM120 Huawei固有の通信スタック HCCL/HCOMM・UBMEM・URMA NVIDIA GPU Ascend 950 NPU dav-3510 CUDA向けの資産はそのまま 通信層はHuawei固有のまま ※本体リポジトリが対応するのはAscend 950(dav-3510)のみ。 A2・A3はコミュニティ管理の別プロジェクト(v0.1.15時点)
共通になるのはカーネルを書く層だけで、実行基盤と通信スタックはベンダー固有のまま残る
ひよこ ひよこ
ペンギン先生、DeepSeekがNvidiaの「CUDA」の代わりになる道具を無料で公開したって聞いたんだけど、そんなことできるの?
ペンギン先生 ペンギン先生
何が公開されたのかを先に見てみようか。DeepSeekのGitHubには、Huawei Ascend NPU向けのリポジトリが並んでいるよ。行列計算の「DeepGEMM-Ascend」、通信用の「DeepEP-Ascend」、コード補完用の「clangd-ascend」の3つで、どれも2026年9月末の更新だね。そしてカーネルを書くための言語「TileLang」のv0.1.15が2026年9月30日に出て、Ascend 950向けのバックエンドが本体に入った。取り込んだPRがマージされたのは前日の9月29日だよ。ただね、「CUDAの代わり」という言い方は、実際の資料とは少しずれているんだ。
ひよこ ひよこ
えっ、ずれてるの? CUDAを置き換えるための道具じゃないの?
ペンギン先生 ペンギン先生
TileLangのREADMEを見ると、対応バックエンドの一覧にNVIDIA CUDA(SM70〜SM120)がちゃんと並んでいるんだよ。AMDのROCm、Huawei Ascend 950、Apple Metal、実験的なLLVM CPUと一緒にね。つまりNvidiaのGPUで動かすときは今までどおりCUDAを通る。共通化されるのは「カーネルを書く層」で、1回書けば複数のハード向けにコードを生成できるようにする、という話なんだ。
ひよこ ひよこ
カーネルって何なの? そこを共通にすると、どうして大事なことになるの?
ペンギン先生 ペンギン先生
行列のかけ算や、チップ同士のデータのやりとりみたいに、GPUやNPUの上で実際に走る小さなプログラムのことだよ。AIの学習や推論の速さは、ほぼここで決まる。だからチップだけ別の会社のものに買い替えても、このカーネル群を書き直さないと性能が出ないんだ。Nvidiaの本当の強みは「チップが速いこと」よりも、CUDA向けに何年も積み上がってきたカーネルとツールの蓄積のほうだ、と言われてきたのはそういう理由だね。
ひよこ ひよこ
じゃあ今回、その積み上がった差が埋まったってことなんだね?
ペンギン先生 ペンギン先生
そこは範囲を確かめたほうがいいよ。TileLangのPR説明には「他のAscendバックエンドの実装はこの統合には含まれない」とはっきり書かれていて、対象はAscend 950(開発コード名 dav-3510)だけ。A2・A3についてはコミュニティが管理する別プロジェクトに案内されている。DeepGEMM-Ascend側も「Ascend 950シリーズで開発・検証した」という書き方なんだ。だから「Ascendなら全部使える」ではなく、「最新世代の1機種向けに整った」が今わかる範囲だね。
ひよこ ひよこ
最新のチップ専用なんだね。でも動くようになったなら、Nvidiaへの依存はなくなったんじゃないの?
ペンギン先生 ペンギン先生
ここが今回いちばん面白いところでね。DeepEP-AscendのREADMEには「HuaweiのHCCL/HCOMM、UBMEM、URMAという通信スタックの上に構築されている」と書かれているんだ。チップ同士をつなぐ一番下の層は、Huawei独自の仕組みをそのまま使っている。つまりNvidia固有の層から離れた代わりに、Huawei固有の層に乗った。依存が消えたんじゃなくて、依存先が移った、と読むのが正確だよ。
ひよこ ひよこ
乗り換えただけなんだ…。それでも性能が同じくらい出るなら、乗り換える価値はあるのかな?
ペンギン先生 ペンギン先生
性能の数字も、条件のほうを見てほしいんだ。DeepEP-AscendのREADMEには、EP(エキスパート並列)のサイズが32までなら、持続的なdispatchが物理的なペイロード帯域の上限の90〜95%程度に達する、と書かれている。性能表自体はEPサイズ8から128まで載っているけれど、「90〜95%」という水準が明記されているのは32までなんだ。規模を大きくしてもそのまま同じ効率で伸びる、とは書かれていない。こういう数字は「どこまでの条件で測ったか」とセットで覚えるのが安全だよ。
ひよこ ひよこ
条件つきの数字なんだね。ところでそのTileLangって、DeepSeekが自分で作った言語なの?
ペンギン先生 ペンギン先生
それも正確には違ってね。TileLangはGitHubの「tile-ai」という組織のプロジェクトで、READMEの謝辞には北京大学の研究者による開発と、Microsoft Researchの研究者によるメンターシップが挙げられている。今回のAscend 950バックエンドのPRにDeepSeekのメンバーが大きく貢献した、という関係なんだ。それにDeepGEMM-Ascend自体は主にC++20で書かれていて、TileLangは一部のカーネルの依存として使われているだけ。「全部がTileLangに置き換わった」わけでもないんだよ。
ひよこ ひよこ
思っていたより細かいんだね…。ぼくみたいに普段AIを使うだけの人には、何か変わることがあるのかな?
ペンギン先生 ペンギン先生
すぐに体感で変わるものではないよ。ただ「同じモデルを別の会社のチップでも動かせる」余地が広がることは、長い目で見ればAIを使う値段と、チップが手に入るかどうかの話につながっていく。そしてもう一つ、今の時点で確認できるのはリポジトリとリリースノートに書かれていることだけ、というのも覚えておきたいな。DeepGEMM-AscendはMITライセンスと明記されている一方で、DeepEP-AscendのREADMEにはライセンスの記載が見当たらないんだ。こういう技術の話は、「書いてあること」と「まだ書かれていないこと」を分けて読むのが、いちばん確実な付き合い方だね。

出典・確認した情報

  • tile-ai/tilelang — Release v0.1.15

    発表日: 2026-09-30 / 確認日: 2026-10-02

    • リリースノートにHuawei Ascend 950(dav-3510)向けバックエンドの追加が記載され、tilelang.ascend.language と target="ascend" が加わったとされている。
    • Cube GEMMとVector計算を1つのカーネル内で組み合わせられ、T.SimdVF・T.SimtVF のリージョンを使うと記載されている。
    • UB・L1・L0の明示的な記憶領域、タイル化したコピー、コア間転送、MXFP8/MXFP4のblock-scaled GEMMに対応すると記載されている。
    • 自動のスケジューリング・パイプライニング・マルチバッファリング・同期挿入が含まれると記載されている。
    • 同じリリースにCUDA向けの自動warp specialization、ROCm対応の拡張など他プラットフォーム向けの変更も含まれている。
  • tile-ai/tilelang — Pull Request #3308「Introduce Ascend 950 backend」

    発表日: 2026-09-29 / 確認日: 2026-10-02

    • 対象は Ascend 950(dav-3510)であり、「other Ascend backend implementations are not included in this integration」と記載され、A2・A3についてはコミュニティが管理する別プロジェクトへ案内されている。
    • Cube(AIC)とVector(AIV)の演算を1つのカーネル内で構成でき、SIMD/SIMTのプログラミングを専用リージョンで行えると記載されている。
    • 2026年9月29日にマージされ、ascend-950-0930 ブランチの134コミットが本体に取り込まれたとされている。
    • 貢献者としてDeepSeekのメンバーが挙げられており、Huawei所属の人物は貢献者一覧に記載されていない。
  • tile-ai/tilelang — リポジトリREADME(2026年10月2日時点の掲載内容)

    発表日: 2026-09-30 / 確認日: 2026-10-02

    • TileLang自身の説明は「A concise domain-specific language designed to streamline the development of high-performance GPU/CPU/NPU kernels」と記載されている。
    • 対応バックエンドとしてNVIDIA CUDA(SM70〜SM120)、AMD ROCm/HIP、Huawei Ascend 950 NPU、Apple Metal、LLVM CPU(実験的)が並記されている。つまりCUDAを排除する構成ではない。
    • Ascend A2/A3、MetaX MACA、Moore Threads MUSAなどは本体とは別のエコシステム用リポジトリとして挙げられている。
    • 謝辞では北京大学の研究者による開発とMicrosoft Researchの研究者によるメンターシップが挙げられており、プロジェクトはGitHubのtile-ai組織で管理されている。
  • deepseek-ai/DeepGEMM-Ascend — リポジトリREADME

    発表日: 2026-09-30 / 確認日: 2026-10-02

    • 「DeepGEMM Ascend is a port of DeepGEMM to the HUAWEI Ascend platform」と記載され、Ascend の MAD(matrix multiply-add)プリミティブへの軽量な抽象を提供するとされている。
    • 「developed and validated on the Ascend 950 series」と記載されており、検証範囲はAscend 950シリーズとされている。
    • 実装は主にC++20で、TileLangは「used by the HC prenorm kernel」として依存の1つに挙げられている。全体がTileLangで書かれているわけではない。
    • ライセンスはMIT Licenseと記載されている。
    • 謝辞で「gratefully acknowledge Huawei for its technical support and engineering expertise throughout the development of DeepGEMM-Ascend」とHuaweiの技術支援に言及している。
  • deepseek-ai/DeepEP-Ascend — リポジトリREADME

    発表日: 2026-09-30 / 確認日: 2026-10-02

    • 「a high-performance communication library for machine learning training and inference on Huawei Ascend NPUs」と記載され、MoEのdispatch/combine向けのエキスパート並列all-to-allなどを提供するとされている。
    • 「DeepEP-Ascend is built on top of Huawei's HCCL/HCOMM, UBMEM and URMA communication stack」と記載されており、Huawei固有の通信スタックの上に構築されている。
    • 対応プロセッサはAscend 950(A5)とされ、性能測定はAscend 950DT NPUで行われたと記載されている。
    • 性能表はEPサイズ8から128まで掲載されているが、「sustained dispatch reaches roughly 90–95% of the physical payload bandwidth limit for EP sizes up to 32」と、90〜95%という水準はEPサイズ32までの記述になっている。
    • 2026年10月2日に確認した時点で、READMEにライセンスの記載は見当たらない。