건축적 특징에서 생태계 구축까지 Muxi Dong Zhaohua는 국내 GPU에서 TVM의 적용 사례를 심층적으로 분석합니다.
7월 5일, HyperAI가 주최한 제7회 Meet AI Compiler Technology Salon이 성황리에 마무리되었습니다. GPU 아키텍처의 근본적인 혁신부터 크로스 하드웨어 컴파일 생태계의 최상위 설계, 단일 칩 연산자 최적화부터 다중 노드 분산 컴파일의 획기적인 발전까지... AI 컴파일 분야의 실무자와 학자들이 이 최고의 기술 축제에 모였습니다. 행사장은 사람들로 가득 찼고, 소통의 열기는 뜨거웠습니다.
위챗 공개 계정 "HyperAI Super Neuro"를 팔로우하고 키워드 "0705 AI Compiler"에 답글을 달면 공인 강사의 발표 PPT를 받으실 수 있습니다.
이 행사에서 AMD의 아키텍트 장닝은 AMD GPU 플랫폼에서 Triton 컴파일러의 성능 최적화 비법을 심층적으로 분석하여 Python 코드로 고성능 GPU 커널을 쉽게 제어하는 방법을 공개했습니다. Muxi Integrated Circuit의 이사인 동자오화는 국내 GPU에서 TVM 애플리케이션에 대한 실제 경험을 공유하며 독립 칩과 오픈 소스 컴파일 프레임워크 간의 충돌에 대한 사례를 보여주었습니다. ByteDance의 연구원 정 사이즈는 Triton 분산의 비밀을 밝히고 Python이 분산 통신의 성능 한계를 어떻게 극복하는지 공유했습니다. 베이징 대학의 왕레이 박사가 선보인 TileLang은 운영자 개발의 효율성 경계를 새롭게 정의했습니다.
기조연설 "Muxi GPU에서의 TVM 응용 실습"에서Muxi Integrated Circuit의 수석 이사인 동자오화는 자사의 GPU 제품의 기술적 특징, TVM 컴파일러 적응 솔루션, 실제 적용 사례 및 생태적 건설 비전을 소개했습니다.이는 고성능 컴퓨팅 및 AI 분야에서 국산 GPU의 기술적 혁신과 응용 잠재력을 입증했습니다.
HyperAI는 동자오화 교수의 연설을 원문의 의도를 훼손하지 않고 편집 및 요약했습니다. 다음은 연설 전문입니다.
Muxi GPU 소개
Muxi GPU는 현재 N 시리즈, C 시리즈, G 시리즈 등 다양한 제품군을 보유하고 있으며, AI 학습 및 추론부터 과학 컴퓨팅까지 광범위한 시나리오를 지원합니다. 다단계 소프트웨어 스택을 구축하여 주류 프레임워크와 완벽하게 통합됩니다. 소프트웨어 스택의 핵심 모듈인 컴파일러는 사용자 친화적인 프로그래밍 인터페이스를 제공하고, 상위 애플리케이션을 최적화하며, 다양한 머신 아키텍처에 맞는 머신 코드를 생성하여 GPU에 전달하여 실행합니다. 엔지니어의 정밀한 조정을 통해 Muxi GPU의 성능은 국제적으로 선진 수준에 도달했으며, 업계 주류 컴퓨팅 라이브러리와도 상호 호환되는 관계를 구축했습니다.
Muxi GPU는 풍부한 명령어 수준 함수 인터페이스를 갖추고 있습니다.자체 개발한 MACA C 인터페이스는 C 언어의 확장 기능을 기반으로 하며, 특정 분야의 문법적 요소를 통합하고 주요 제조업체의 기본 프로그래밍 인터페이스와 기능적으로 동등하여 개발자가 신속하게 마이그레이션 및 적응을 완료할 수 있도록 지원합니다. 동시에 Python, Triton, Fortran과 같은 다양한 프로그래밍 인터페이스를 제공하고, OpenACC 및 OpenCL과 같은 병렬 프로그래밍 표준을 지원하며, 자동 병렬화 코드 생성 효율성이 뛰어납니다.
또한,Muxi GPU는 GPGPU(범용 그래픽 처리 장치) 아키텍처를 채택했습니다.LLVM 기반 컴파일 시스템은 고급 언어에서 저수준 머신 코드까지 전체 프로세스 최적화를 지원하고, 개발 효율성과 하드웨어 성능을 모두 고려하여 고성능 소프트웨어 스택을 제공합니다.
Muxi GPU에서의 TVM 적응
오픈소스 딥러닝 컴파일러인 TVM은 딥러닝 모델을 다양한 하드웨어에서 효율적으로 실행 가능한 코드로 변환할 수 있습니다. Muxi 팀은 자체 GPU의 특성을 기반으로 완전한 TVM 적응 솔루션을 구축하여 모델 정의부터 하드웨어 실행까지 전체 프로세스 최적화를 달성했습니다.
컴파일러 아키텍처 관점에서 완전한 지원이 달성되었습니다.이론상으로는 4개의 핵심 수준에 연결할 수 있습니다.
C++ 인터페이스를 적용하기 위해 MACA 언어로 변환하여 솔루션을 구축하고자 합니다. 이 과정은 매우 어렵고, 도구 기반 자동 변환을 구현하는 데에는 몇 가지 어려움이 있습니다.
코드 추상화가 높을수록 크로스 레벨 적응을 달성하기가 더 쉽습니다. 또한, LLVM에 연결할 때는 버전 호환성 문제에 주의해야 합니다. LLVM 버전이 다양하고, 특정 버전의 적응은 해당 버전의 지원 여부에 따라 달라지며, 버전 불일치로 인해 비정상적인 컴파일 프로세스가 발생할 수 있습니다.
Muxi Arch의 GPU 적응 측면에서:
tvm.Target은 MACA Target을 추가하고 각 단계에 대한 지원을 추가합니다. 먼저, 일반 GPU 프로세스를 재사용하기 위해 transform/lower에 MACA Target의 파이프라인을 추가합니다. 그런 다음 튜닝 단계에 MACA Target의 스케줄링 규칙을 추가합니다. 마지막으로 CodeGenMACA에 대한 지원을 추가하고 MACAC 코드를 컴파일합니다.
또한, MACA 장치와 MACA 런타임 API의 사용이 tvm.Device에 추가되어 MACA 장치의 메모리 작업과 런타임에 커널을 시작합니다.
현재 당사 제품 중 다수가 TVM 수준에서 하위 장치를 지원하고 백엔드는 하위 장치를 기반으로 다양한 제품에 대한 최적화 작업을 수행합니다.다양한 제품의 경우 컴파일러는 장치 유형의 차이에 따라 자동으로 해당 적응 솔루션을 선택합니다.동시에, 일괄 컴파일 시나리오에서 우리는 다양한 아키텍처에 대해 고정된 선택 구성을 만들려고 노력하고 있습니다.일반 컴파일 단계에서 컴파일러는 특정 구성에 따라 다양한 아키텍처에 대한 함수 관련 컴파일 규칙을 동적으로 조정합니다.
운영자 적응 측면에서 우리는 주로 다음을 포함하는 상위 레벨 구성을 수행했습니다.
* 하위 일정에서 TIR 스케줄링 기본을 통해 PrimFunc에 일정을 추가합니다.
* Split Mixed Module에서 대상 및 기타 정보를 추가하고, MACA 내장을 주입하고, MACA에 대한 동기화 지침을 추가합니다.
* CodeGenMACA를 사용하는 경우 MACA가 의존하는 헤더 파일을 포함하고, tir.mma 관련 명령어에서 MACA WMMA API의 사용법을 생성하고, 다양한 유형의 변수를 선언하여 사용합니다.
적응 과정에서 더 나은 성능을 달성하기 위해 우리는 해당 특성에 따라 특별한 처리를 수행했습니다.
* 튜닝 중 텐서라이즈를 활성화할 수 없습니다. conv2d 연산자 매개변수 그룹이 1이 아니고 MACA 연산자 라이브러리 구현이 TOPI에서 직접 사용됩니다.
* onnx 모델 임포트 후 사용자 지정 연산자: Multi Head AttentionV1 연산자는 MACA onnx 런타임에서 고도로 최적화되었습니다. 연산자 호출은 contrib에 캡슐화되어 TVM이 onnx 모델을 임포트한 후 수동으로 최적화된 고성능 연산자 구현을 직접 사용할 수 있습니다.
우리에게는,공급업체 사용자 정의 최적화 연산자는 Python 및 MAC A로 구현될 가능성이 더 높습니다.
* Python 인터페이스에서 Relax.Function은 기본 연산자를 결합하여 정의됩니다. tir.PrimFunc는 tir을 사용하여 연산자 구현을 정의하고 필요에 따라 일정을 추가합니다.
* MACA C 인터페이스에서 tir.call_packed는 고성능 연산자 라이브러리의 구현을 캡슐화하여 사용합니다. tir.call_kernel은 MACA C로 구현된 커널 코드를 사용하여 TVM 스택을 통해 PackedFunc 호출로 컴파일합니다.
또한 Muxi GPU의 하드웨어 특성을 최대한 활용하기 위해이 팀은 TVM 스케줄링 알고리즘을 심층적으로 최적화했습니다.
* MACA 대상에 대한 WMMA float32 유형에 대한 지원이 추가되었습니다.
첫째, MACA에서 float32 유형 WMMA API를 지원하고 MACA ScheduleRules에 float32 유형 자동 텐서화 규칙을 추가하여 TVM이 자동으로 하드웨어 WMMA를 식별하고 사용할 수 있도록 하고, dlight 최적화 프레임워크에 해당 float32 텐서화 최적화를 추가하여 행렬 작업 효율성을 개선합니다.
* 스케줄링 알고리즘에 대한 비동기 복사의 영향을 평가합니다.
한 그룹을 로드하고 한 그룹을 계산하는 것에서 다음 데이터 그룹을 비동기로 로드하고 현재 데이터 그룹을 동기로 계산하는 것으로 여러 wmma 계산을 최적화하고, 파이프라인 효율성을 개선하고, MACA ScheduleRules에서 소프트웨어 파이프라인 최적화 논리를 활성화하고, MACA 대상에 대한 비동기 복사 명령 주입 및 코드 생성을 추가합니다.
또한 새로운 데이터 유형을 지원하기 위해 몇 가지 시도를 했습니다.DataType 시스템에서 MACA 대상 적응을 활성화합니다. MACA ScheduleRules에서 Float8 유형의 자동 텐서화 논리를 지원하고 Float8과 같은 사용자 정의 데이터 유형에 대한 TVM 지원을 확장합니다. CodeGenMACA에서 Float8 유형 변환 및 연산 코드 생성에 대한 지원을 구현하고 maca_half_t.h에서 관련 연산 정의를 보완합니다.
Muxi GPU에서의 TVM 애플리케이션
프레임워크 설계 측면에서 팀은 두 가지 액세스 방법을 구현했습니다.하나는 토치 모델을 직접 가져와서 Relay 프런트엔드에서 실행하는 것이고, 다른 하나는 torch.compile을 사용하여 TVM을 백엔드로 사용하는 것입니다.이는 상위 프레임워크와 기본 하드웨어 간의 효율적인 연결을 달성합니다.
성능 평가 단계에서는 ReseNet50과 Bert를 벤치마크 모델로 선정하고, 심층적인 최적화 없이 Torch와 TVM 컴파일 및 실행의 성능을 비교했습니다.실험 데이터에 따르면 TVM은 어떤 면에서는 상당한 이점이 있으며, 어떤 경우에는 토치보다 성능이 뛰어납니다.이는 TVM 중간 표현(IR)의 유연성과 하드웨어 특성에 대한 목표 최적화 덕분입니다.
TileLang은 TVM 생태계의 도메인 특정 언어(DSL)로, 텐서 컴퓨팅의 정교한 최적화에 중점을 두고 있습니다.저희 팀은 다음과 같은 측면에서 심층적인 기능적 적응을 수행했습니다.* MACA 대상 사용 지원 * MACA C 커널 코드 생성을 위해 CodeGenTileLangMACA 추가 * mcTVM을 종속성으로 교체 및 사용 * libgen, 어댑터 및 래퍼에 MACA 대상 처리 논리 추가 * tl_template에서 gemm을 사용하여 MACA 대상 정의 추가
최적화 측면에서 작업은 주로 tl_template에서 gemm을 구현하고 Muxi GPU의 특성에 맞게 조정되는 알고리즘을 구현하는 데 중점을 둡니다.
공급업체와 커뮤니티 간의 상호 작용 측면에서 TileLang의 디자인과 개발은 세 가지 주요 원칙의 균형을 이루어야 합니다.
언어 설계 측면에서,가장 먼저 해결해야 할 것은 추상화와 고성능 간의 균형을 맞추는 것입니다.특히, 여러 컴퓨팅 유닛을 다룰 때 컴파일러는 유연한 전략 선택 메커니즘을 제공해야 합니다. 즉, 기본 하드웨어 특성을 깊이 이해하는 개발자가 특정 코드 생성 경로를 지정할 수 있도록 하는 동시에, 일반 개발자가 기본 세부 사항에 신경 쓰지 않고도 추상 인터페이스를 통해 효율적인 프로그래밍을 수행할 수 있도록 지원해야 합니다. 이러한 균형은 DSL로서 TileLang이 설계 단계에서 집중해야 할 핵심 요소입니다.
둘째,공급업체의 맞춤형 구성과 DSL의 표준화를 고려해야 합니다.컴파일러는 다단계 최적화 옵션 구성을 지원해야 합니다. 예를 들어, 고급 사용자는 루프 언롤링 레벨 및 상속된 압력과 같은 기본 컴파일 전략을 매개변수를 통해 조정하여 더 나은 코드 생성 결과를 얻을 수 있습니다. 동시에, 컴파일러는 다양한 하드웨어 아키텍처에 대해 개발자가 하드웨어 특성에 따라 최적의 컴파일 경로를 선택하고 개발 효율성을 향상시킬 수 있도록 맞춤형 최적화 팁과 주석 메커니즘을 제공해야 합니다.
제삼,제품 세대 간의 인터페이스 연속성이 보장되어야 합니다.컴파일러와 언어 툴 체인의 반복 과정에서 인터페이스 설계의 하위 호환성을 보장해야 합니다. 현재 버전의 컴파일 로직과 생성된 코드는 차세대 하드웨어 제품에서도 효과적으로 실행될 수 있어야 아키텍처 반복으로 인한 코드 재구성 비용을 피할 수 있습니다. 이러한 연속성은 생태학적 구성의 기반이 되며, 비호환성으로 인한 "감소적" 손실이 아닌 컴파일러 툴 체인 기능의 "가산적" 축적을 달성할 수 있습니다. 동시에 사용자의 학습 및 마이그레이션 비용을 줄이고 사용상의 혼란을 방지할 수 있습니다.
도전과 기회
마지막으로, 현재 산업 발전에 따른 과제와 기회에 대해 이야기하고 싶습니다.이러한 과제는 주로 다음과 같은 측면에서 반영됩니다.
첫 번째는 프레임워크와 애플리케이션 기반 알고리즘이 빠르게 변화한다는 것입니다.딥러닝과 같은 분야의 급속한 발전으로 상위 프레임워크와 알고리즘의 업데이트 주기가 계속 짧아지고 있으며, 기능과 성능이 향상되면서 컴파일러 적응에 대한 압력이 커지고 있습니다. 컴파일러 커뮤니티는 새로운 운영자와 새로운 컴퓨팅 모델의 지원 요구에 신속하게 대응할 수 있는 효율적인 적응 메커니즘을 구축해야 합니다.
둘째, 하드웨어 아키텍처는 계속해서 발전하고 있습니다.현재 컴파일러는 일부 GPU 아키텍처의 기능을 지원할 수 있습니다. 향후 새로운 하드웨어 아키텍처 기능이 등장할 경우, 컴파일러는 이기종 하드웨어 기능도 수용할 수 있어야 합니다.
세 번째는 프로그래밍 패러다임이 계속해서 발전한다는 것입니다.기존의 C/C++부터 새로운 함수형 프로그래밍과 이기종 프로그래밍 모델까지, Python과 관련된 생태적 사슬을 정의하는 방법은 큰 과제입니다.
마지막으로, 정확성, 성능, 전력 소비 간의 균형입니다.실제 애플리케이션에서 컴파일러는 코드 성능뿐만 아니라 하드웨어 전력 소비에도 주의를 기울여야 하며, 이는 마찬가지로 중요합니다. 이러한 요소들은 후속 명령어 선택 및 아키텍처 설계와 관련이 있습니다.
미래,우리는 지역 사회와 협력하여 다음을 구축하고 싶습니다.
오픈 소스 전략 측면에서는 FlashMLA와 같은 핵심 컴퓨팅 모듈을 포함하여 프레임워크와 연산자 라이브러리의 핵심 구성 요소를 공개할 계획입니다. 오픈 소스 모델을 통해 컴파일러 툴 체인의 반복적 최적화를 촉진하고 생태계적 규모 효과를 형성할 것입니다.
둘째, 업계의 애플리케이션, 프레임워크, 연산자 라이브러리, 컴파일러, 그리고 하드웨어 아키텍처 간의 협력 기회가 더욱 확대되기를 바랍니다. 정기적인 기술 교류(예: 업계 포럼)를 통해 컴파일 최적화의 어려움, 연산자 스케줄링 전략과 같은 핵심 이슈에 집중하고, 도메인 간 협력을 증진하며, 기술 혁신을 모색해 나갈 것입니다.
Muxi는 생태계 공동 구축에도 중점을 두고 있습니다. Muxi의 구축 계획에는 컴파일러 툴 체인에 대한 개발자들의 피드백과 문제 보고서를 받을 수 있는 기술 커뮤니티 포럼 구축, 운영자 및 프레임워크 테마 경진대회 개최, 벤치마크 제공, 그리고 커뮤니티와 함께 분야별 테스트 스위트 및 벤치마크 공동 구축 등이 포함됩니다.