DeepSeek, Huawei Ascend 950용 DeepGEMM 커널 공개

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 패키지 내보내기 항목 에는 밀집 및 그룹 GEMM 함수, MQA 함수, 스케일링 인자 배치 도구, JIT 컴파일 및 AI core 사용 제어 기능이 있습니다. 저장소는 MIT 라이선스입니다.

DeepSeek은 Ascend 포트가 상위 DeepGEMM과 API 수준에서 완전히 호환된다고 말합니다. 동일한 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로 보고했습니다. 이는 지정된 형식과 형태에서 측정한 커널 성능이지 모델 처리량 측정치가 아닙니다. 표에는 각기 다른 설정의 그룹 GEMM, MQA, MegaMoE, prenorm 결과도 있습니다. MegaMoE 수치는 8개 rank, 8-way expert parallelism, top-k 라우팅 6, 공유 expert 1개 설정의 평균입니다.

해당 테스트 코드는 GEMM 출력을 생성된 기준값과 비교하고 지원되는 밀집 사례에서 ACLNN 교차 검사를 사용합니다. 일부 형태와 레시피에서는 이 추가 ACLNN 비교를 명시적으로 건너뜁니다. 공개된 코드와 테스트를 통해 주장을 살펴볼 수 있지만 BIG CHANGE가 README 수치를 독립적으로 재현하지는 않았습니다. 다른 가속기 스택과 동일 조건으로 비교한 벤치마크, 비용 비교, 종단 간 서비스 결과는 공개되지 않았습니다.

주요 변화

DeepSeek은 Ascend 950 시리즈용 DeepGEMM Ascend 커널 코드를 공개했습니다.

개발자는 해당 하드웨어에서 익숙한 GEMM 및 관련 모델 연산을 살펴보고 빌드할 수 있습니다. 공개 자료에는 CANN 9.20과 Ascend 전용 스케일링 인자 배치가 설명되어 있습니다.

Ascend 950을 사용하는 팀은 이제 자체 모델 형태에 맞춰 평가할 커널 코드를 확보했습니다. DeepSeek은 한 Ascend 950DT 테스트 구성에서 하드웨어 한도에 가까운 결과를 보고했으며, BIG CHANGE는 커널을 실행하지 않았습니다.

전체 애플리케이션 성능, 다른 Ascend 기기 지원, 운영 비용은 아직 확인되지 않았습니다.

출처 및 추가 읽을거리

• DeepGEMM Ascend 저장소와 README: 9월 30일 출시 기록, 사전 요구 사항, 인터페이스 관련 주장, 프로젝트 자체 벤치마크 표입니다. README는 프로젝트 유지 관리자가 제공한 1차 자료이지 독립 검증 결과가 아닙니다.
• 패키지 내보내기 항목 및 빌드 설정: 최초 커밋에 기록된 버전, 공개 Python 함수, 종속성 선언, 소스 빌드 연결 방식입니다.
• GEMM 테스트 하네스: 저장소가 출력을 검사하고 프로파일러를 호출하는 방식, ACLNN 비교를 건너뛰는 조건을 설명합니다. 이 파일은 공개 수치를 독립적으로 재현하지 않습니다.