deepseek-ai/DeepGEMM-Ascend

DeepGEMM-Ascend: clean and efficient matrix multiplication kernel library for Huawei Ascend NPUs

DeepGEMM‑Ascend란 무엇인가?

DeepGEMM‑Ascend는 오픈소스 DeepGEMM 행렬 곱셈 라이브러리의 하드웨어 특정 포트로, Huawei의 Ascend NPU 제품군(예: Ascend 950)을 대상으로 합니다. DeepGEMM과 동일한 Python 수준 API를 구현하므로 다른 플랫폼(CUDA, ROCm 등)용으로 작성된 코드는 이 패키지를 설치하기만 하면 Ascend에서 실행할 수 있습니다.

이 라이브러리는 Ascend MAD(행렬 덧셈 곱셈) 프리미티브를 얇은 C++/Python 확장으로 래핑하여 프랙탈 레이아웃, 정렬, 주소 계산과 같은 저수준 세부 사항을 처리합니다. 또한 희소 데이터 로딩 및 코루틴 기반 파이프라이닝과 같은 Ascend 전용 최적화를 추가하여 다양한 행렬 형태에서 하드웨어의 이론적 최대 성능에 도달하는 것을 목표로 합니다.

누가 사용할까?

대상	중요성
AI 연구자／엔지니어가 Ascend NPU에서 실행되는 대규모 언어 모델(LLM)을 구축하는 경우 – 빠른 GEMM, MQA, MoE 커널이 필요합니다.
프레임워크 개발자가 PyTorch‑NPU(`torch_npu`)를 사용자 정의 고성능 연산으로 확장하는 경우.
성능 엔지니어가 극한으로 최적화된 Ascend 커널의 참조 구현을 찾는 경우.
하드웨어 인식 ML 실무자가 Ascend에서 BF16, FP8, FP4 및 혼합 정밀도 행렬 연산을 벤치마킹하려는 경우.

핵심 기능(README에 설명된 대로)

기능	세부 사항
API 호환성	업스트림 DeepGEMM 프로젝트와 동일한 패키지 이름과 함수 시그니처를 사용합니다.
정밀도 지원	BF16, FP8, FP4 GEMM; MQA logits 및 MegaMoE용 특수 커널도 제공합니다.
경량 추상화	Ascend 특정 레이아웃 및 정렬 복잡성을 숨기면서 커널 코드를 간결하게 유지합니다.
Ascend 특정 최적화	희소 데이터 로딩, 코루틴 기반 파이프라이닝, 사용자 정의 스케일링 팩터 패킹(UE8M0 쌍 → `int16`).
유틸리티 헬퍼	AI 코어 수 설정, NPU 아키텍처 선택, 결정적 모드 제어, 스케일링 팩터 레이아웃 변환, 이론적 정렬 한계 조회 기능.
JIT 런타임	DeepJIT 런타임에 의존; 커널은 첫 사용 시 컴파일되고 캐시됩니다(`DG_JIT_*` 환경 변수로 구성 가능).
환경 변수 튜닝	디버그, 캐시 위치, 실행 시간 초과 및 기타 설정을 `DG_JIT_*` 및 표준 Ascend 변수로 노출합니다.
성능	벤치마크는 BF16/FP8/FP4에서 고밀도 GEMM에 대해 하드웨어 한계의 98‑99%를 달성하며, 그룹화된 GEMM, MQA, MegaMoE 및 사용자 정의 "HC Prenorm" 커널에 대한 자세한 수치도 제공합니다.

시작 방법(빠른 시작 요약)

하드웨어 및 소프트웨어 사전 요구 사항
- Huawei Ascend NPU(950 시리즈에서 검증됨)
- CANN 9.20 툴킷(bisheng, ld.lld)
- torch_npu Python 패키지
- Python 3.10+, C++20 호환 컴파일러, tilelang, tree‑sitter, tree‑sitter‑cpp

서브모듈과 함께 클론

git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git
cd DeepGEMM-Ascend

C++ 확장 빌드(리포지토리에는 include를 링크하고 컴파일러를 실행하는 develop.sh 스크립트가 제공됩니다).
설치
```
pip install . --no-build-isolation
```

Python 코드에서 동일한 DeepGEMM API 사용, 예:

import deep_gemm as dg
dg.set_num_sms(8)          # 사용할 AI 코어 수 선택
out = dg.aclnn_fp8_fp4_gemm_nt(A, B, scaling_factors)

언제 빛을 발할까?

README에는 라이브러리가 Ascend 하드웨어 한계에 근접하는 여러 벤치마크 범주가 나열되어 있습니다:

고밀도 GEMM – BF16, FP8, FP4에서 이론적 TFLOPS의 최대 99.8%.
M‑Grouped GEMM – Mixture‑of‑Experts(MoE) 전문가에게 유용하며, 실제 전문가‑토큰 패턴에서 > 800 TFLOPS 달성.
MQA Logits – 멀티 쿼리 어텐션을 위한 빠른 스코어링, 지연 시간 100‑300 µs 범위.
MegaMoE – 엔드투엔드 융합 커널(dispatch + 두 개의 GEMM + SwiGLU)로 계산 및 메모리 대역폭을 모두 높게 유지.
HC Prenorm GEMM – "Manifold‑Constrained Hyper‑Connections" 모듈용 특수 커널로 HBM 쓰기 대역폭을 포화시킵니다.

이 수치는 라이브러리가 Ascend 하드웨어에서 프로덕션 규모의 LLM 추론／훈련 워크로드에 사용할 준비가 되었음을 시사합니다.

제한 사항／주의 사항

하드웨어 특정 – Huawei Ascend NPU에서만 작동; CPU 또는 GPU 폴백 없음.
툴체인 무거움 – CANN 9.20, bisheng 컴파일러, <format>을 지원하는 C++20 컴파일러 필요.
환경 변수 튜닝 – 성능은 DG_JIT_* 변수와 Ascend 실행 차단／디버그 플래그를 올바르게 설정하는 데 달려 있을 수 있습니다.
베타 상태 – 리포지토리는 초기 릴리스(2026‑09‑30)로 보고; 커뮤니티 지원 및 문서는 아직 제한적일 수 있습니다.

AI 생태계에서의 위치

DeepGEMM‑Ascend는 AI 스택의 저수준 선형 대수 계층에 위치합니다. 현대 LLM은 계산의 대부분을 대규모 행렬 곱셈(GEMM) 및 관련 어텐션 커널에 사용합니다. 널리 사용되는 DeepGEMM API를 미러링하는 고도로 최적화된 Ascend 네이티브 구현을 제공함으로써 이 프로젝트는 다음을 가능하게 합니다:

원활한 포팅 – 기존 PyTorch 기반 모델을 커널을 다시 작성하지 않고 Ascend로 포팅.
정확한 성능 기준 – Ascend를 NVIDIA/AMD GPU와 비교하는 연구자용.
참조 코드 – 새로운 Ascend 최적화 기술(희소 로딩, 코루틴 파이프라인, 사용자 정의 스케일링 팩터 패킹)을 탐구하는 하드웨어 엔지니어용.

라이선스 및 인용

라이선스: MIT – 상업 및 연구 목적으로 허용.
인용: 학술 논문용으로 복사 가능한 BibTeX 항목이 제공됩니다.

TL;DR

DeepGEMM‑Ascend는 Huawei Ascend NPU용 드롭인, 고성능 GEMM 라이브러리로, 원래 DeepGEMM 프로젝트와 동일한 Python API를 노출하면서 Ascend 특정 최적화를 추가합니다. LLM 워크로드(고밀도 GEMM, MoE, MQA)를 대상으로 하며 다양한 행렬 형태에서 거의 최대 하드웨어 활용률을 주장합니다. pip로 설치할 수 있으며 Ascend CANN 툴킷과 C++20 컴파일러가 필요합니다.