인공지능
DeepSeek이 화웨이 Ascend용 라이브러리 여섯 개를 오픈소스로 공개했다. 수출 통제가 한 번도 건드리지 않은 것이 바로 이 소프트웨어 계층이다
DeepSeek이 9월 30일 화웨이 Ascend 950용 TileLang 백엔드와 커널·통신 라이브러리 다섯 개를 공개했다. 화웨이와 함께 개발했고, CUDA 백엔드 옆으로 언어 본체에 업스트림 병합됐다. 수출 정책은 칩을 막았다. CUDA를 막은 적은 없고, 바꾸기 어려운 쪽은 언제나 CUDA였다.
MAI
DeepSeek이 9월 30일 화웨이의 Ascend AI 가속기용 오픈소스 소프트웨어 구성 요소 여섯 개를 공개했다. 화웨이와 공동으로 개발했고, 발표는 회사 자체 WeChat 채널을 통해 나왔다. 중심에 있는 것은 연산 커널을 작성하기 위한 언어 TileLang이며, 여기에 공식 Ascend 950 백엔드가 붙었다. 나머지 다섯 개는 커널과 통신 라이브러리로, 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 패스가 의존성을 고려한 명령 파이프라이닝, 멀티 버퍼링, 그리고 코어 내·코어 간 동기화 배리어 삽입을 담당한다. 디바이스 코드는 화웨이 CANN 툴킷의 bisheng 컴파일러를 거치고, 텐서와 스트림은 PyTorch의 NPU 백엔드와 상호 운용된다. DeepSeek과 화웨이는 또한 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년치의 커널, 라이브러리, 프레임워크 통합, 그리고 엔지니어의 손에 붙은 감각이 그 프로그래밍 모델을 전제로 한다는 데 있다. Ascend 부품을 창고째 받은 중국 연구소는 돌리고 싶은 소프트웨어가 다른 것을 위해 쓰여 있다는 문제를 여전히 마주한다.
이번 공개가 메우려는 것이 그 틈이며, 그것을 메울 수 있는 유일한 위치에서 메운다. 화웨이는 한동안 CANN을 오픈소스로 풀고 PyTorch, Triton, vLLM에 연결해왔다. 그러나 칩 벤더가 자기 도구를 쓰는 것은 벤더가 영업을 하는 것일 뿐이다. DeepSeek이 쓴다는 것은 그 나라에서 가장 유능한 모델 연구소가 이 도구가 프런티어 훈련에 쓸 만하다고 보증하는 일이다. 회사 설명에 따르면 V4 계열 훈련에 쓰이는 TileLang 연산자는 이제 모두 고성능 Ascend 구현을 갖췄다. DeepSeek은 4월에 V4를 화웨이 하드웨어로 이식했다. 네이멍구에 Ascend 가속기 약 16만 기를 담는 데이터센터를 짓고 있다는 보도도 있다. 오늘 공개된 라이브러리들은 그 숫자에 의미를 부여하는 쪽이다.
이것이 아닌 것
라이브러리 여섯 개와 컴파일러 백엔드 하나는 CUDA가 아니다. 백엔드가 겨누는 것은 Ascend 950에 한정되며, 그 이전의 A2와 A3는 여전히 별도의 커뮤니티 프로젝트가 떠받친다. PTO 코드 생성 경로와 가상 머신 내장 명령은 업스트림 과정에서 빠졌다. 커널들이 덮는 범위는 DeepSeek 자신의 워크로드가 필요한 것——어텐션, 행렬곱, 전문가 혼합 모델의 통신, 데이터 선택——이며, 이는 좁고 대단히 잘 고른 단면이지 범용 생태계가 아니다. 자기 모델이 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