BuildRiot
Sign in

Article

DeepSeek expands Ascend support, but migration still needs hardware-specific work

DeepSeek’s September 30, 2026, release expanded Ascend support across matrix kernels, communication libraries and TileLang. The repositories show where developers can reuse interfaces and where numerical layouts, compiler backends and deployment requirements still need work.

Blue modules cross a black bridge toward differently shaped sockets on an abstract circuit board, illustrating software migration that requires hardware-specific adapters.
Illustration: BuildRiot (AI-generated)

Key takeaway

Shared APIs and a common kernel language give developers more ways to move workloads to Ascend. Compatibility claims still need correctness tests, operation coverage checks and application benchmarks on obtainable hardware.

New libraries build on earlier Ascend support

DeepSeek’s September 30, 2026, infrastructure release expanded matrix kernel and communication support for Huawei Ascend hardware. Tom’s Hardware reported cooperation between DeepSeek and Huawei, including optimization on a 128-chip Ascend 950 supernode. The repositories document reusable APIs and a shared kernel language, with substantial Ascend-specific engineering underneath. They do not establish automatic migration from Nvidia’s CUDA platform.

TileLang’s Ascend support predates the announcement. Its upstream changelog records a public release on January 20, 2025, and external Ascend adapters on September 29, 2025. The external tilelang-ascend project later added automatic synchronization, vectorization and software pipelining. Upstream’s September 30, 2026, announcement introduced a native Ascend 950 backend, separate from the external A2/A3 projects.

TechRadar identified DeepGEMM-Ascend and DeepEP-Ascend as new repositories, alongside updates to existing projects. DeepGEMM-Ascend dates its initial release to September 30, 2026. The available evidence does not settle DeepEP-Ascend’s exact publication date. The announcement therefore combines new libraries with compiler support that already had a development history.

A shared language needs an Ascend compiler stack

TileLang illustrates how a common language can retain hardware-specific implementation choices. The upstream Ascend 950 guide describes a shared frontend with Ascend-specific compiler transformations, scheduling, synchronization and code generation. Its example combines matrix multiplication with an activation function, using bfloat16 inputs and 32-bit floating-point accumulation and output. Execution explicitly uses Ascend’s Cube and Vector facilities, including direct transfers between them.

The guide says developers can build this backend without a CUDA toolkit. The build still requires compatible Ascend drivers, Huawei’s software stack, Bisheng, a compatible linker, runtime libraries, PyTorch and torch_npu. Its default target architecture is dav-3510. Reusing the frontend therefore still requires changing the compilation and execution environment.

An external-adapter issue opened on September 22, 2026, reported an incorrect architecture target, wrong results after an override and a compiler route that hung on Ascend 950PR. These are the reporter’s findings and were not independently reproduced in the supplied research. The issue predates the release and does not establish defects in the new upstream backend. Migration reports need to identify the repository, commit, compiler route and hardware target.

DeepGEMM compatibility has a numerical boundary

DeepGEMM-Ascend claims full compatibility with the DeepGEMM API. Its documented coverage includes bfloat16, 8-bit and 4-bit floating-point matrix multiplication, along with MegaMoE. The implementation uses Ascend-specific primitives and pipelining. Requirements include Ascend 950, version 9.20 of Huawei’s compute software stack, torch_npu, Python 3.10 or later and C++20. The repository declares an MIT license.

The interface documentation identifies a scaling-factor layout that differs from Nvidia’s. It packs pairs of scaling factors along the matrix’s K dimension into 16-bit integers and supplies conversion helpers. Matching floating-point format labels therefore does not establish that existing buffers can transfer unchanged. Developers need to test conversion correctness and measure its overhead.

Issue 6 raises a separate question about argument behavior. A reporter examining commit 8491bbb alleged that separate accumulator c and output d tensors cause the kernel to use existing d values instead of c. The report says current tests use the same tensor for both arguments, potentially concealing the discrepancy. The issue remained open when retrieved, and the research did not independently reproduce it. A migration test should exercise both shared and separate tensors.

Kernel benchmarks have a limited scope

DeepGEMM-Ascend’s published dense benchmark reports 431 trillion floating-point operations per second for bfloat16 inputs, against a stated limit of 432 trillion. The repository reports 99.8% utilization for that case. It specifies Ascend 950DT, compute software version 9.20, bench_msprof, a cold second-level cache and matrix dimensions of M=4096, N=7168 and K=16384. Reported latency is 2,229.9 microseconds. These are publisher-reported measurements, not independently reproduced results.

For 4-bit floating-point inputs at the same dimensions, the table reports 1,701 trillion floating-point operations per second and 565.6 microseconds. Its stated utilization is 98.3%. A smaller 8-bit case, with dimensions of 128, 576 and 7168, reports 105 trillion floating-point operations per second and 10.1 microseconds. These results show why large-matrix utilization cannot represent every workload. They do not measure application speedups over Nvidia hardware.

Satyajit Ghana’s October 2, 2026, analysis distinguishes dense matrix multiplication’s 99.8% utilization from fused MegaMoE’s approximately 98% and DeepEP’s bandwidth utilization. Ghana examined repository code and published tables rather than independently benchmarking the hardware. The analysis attributes the dense-to-MegaMoE difference to communication cost. Comparing different workloads and precisions does not establish that cause; isolating it would require a controlled experiment.

Communication support depends on operations and networking

DeepEP-Ascend aligns its public buffer APIs with Nvidia DeepEP while using Ascend-specific communication components and runtime compilation. Its requirements include Ascend 950/A5, suitable interconnect channels and an even number of ranks. The repository lists missing Ascend kernels for all-reduce, reduce-scatter and expert-load-balancing operations. It also says graph capture is unsupported. Applications need to check implementation coverage beyond the available API names.

The communication benchmark specifies Ascend 950DT, compute software version 9.2.0, external Clos networking and capacity for 16,384 tokens per rank. It uses a hidden size of 7,168, routes tokens to six of 256 experts, and uses 8-bit floating-point dispatch with bfloat16 combine. Across eight expert-parallel ranks, reported dispatch bandwidth is 373–375 gigabytes per second and combine bandwidth is 345–347 gigabytes per second. At 128 ranks, those ranges fall to 313–320 and 272–278 gigabytes per second, respectively. The ranges reflect differences across ranks.

Benchmark timing includes issue and drain but excludes final epilogues. The results also rely on a manually configured, nonpublic proof-of-concept hardware development kit. The repository’s claimed 90–95% payload-bandwidth utilization applies to dispatch through 32 expert-parallel ranks; larger configurations and combine remain under optimization. Reproduction therefore depends on firmware, topology and timing boundaries as well as source availability.

Test migration against the application’s requirements

Start with a representative kernel and the communication operations the application actually uses. Record the backend and release commit, then track code changes, hardware-specific branches, tuning time and correctness failures. Test scaling-factor conversion and accumulator behavior before relying on API compatibility claims. Include small decode shapes and uneven expert loads, then measure complete application latency and throughput.

Release terms and hardware availability need separate checks. Ghana reported that DeepEP lacked a license file, and the retrieved first-party repository listing likewise showed no root license file. That observation leaves reuse terms unresolved without constituting a complete licensing audit. Developers should obtain explicit licensing terms and determine whether publicly obtainable firmware supports the required communication behavior.

Matched Ascend and Nvidia application benchmarks would provide stronger evidence for deployment decisions. Useful comparisons would report throughput, latency, power and cost under equivalent workloads. The supplied research established no independently reproduced comparison meeting those conditions. Before committing to a migration, request exact firmware, topology and rank mapping, then verify complete-operation timings on the system the team can obtain.

Attribution

Sources

Read the original source for full details and context.

Report a problem

Links, dates or what you expected. Do not include passwords or private data.