diff --git a/docs/design/typedpython.md b/docs/design/typedpython.md index 689458fa..c0869d44 100644 --- a/docs/design/typedpython.md +++ b/docs/design/typedpython.md @@ -410,9 +410,46 @@ Mojo 의 절반은 CPU SIMD 다. 등급 1 `@compiled` 의 C 백엔드(§4.3)가 - **자동 융합은 뒤로 미룬다.** torchnative 의 그래프 캡처로 `torch.compile` 을 대신하는 길(캡처한 그래프의 융합 구간을 커널 컴파일러로 넘김)은 가치가 크지만, 그 캡처는 동적 shape · in-place 변경 · 제어 흐름을 거부하고 LLM 의 KV 캐시 갱신이 in-place 라 정면으로 걸린다. 그전까지는 리졸버를 통한 - 레이어 교체로 착지한다. + 레이어 교체로 착지한다 (§5.9). -### 5.7 선례, Taichi +### 5.7 빌드 타임 오토튜닝 + +같은 커널도 타일 크기, 파이프라인 단계 수, 워크그룹 크기에 따라 몇 배씩 차이가 난다. Triton · +TileLang · MLC-LLM 은 이것을 자동 탐색한다. 모바일에서는 실행 중에 탐색할 수 없으므로(§5.1) 탐색을 +**빌드 타임으로 옮긴다.** + +- 커널마다 조정 가능한 매개변수(타일 크기, 단계 수, 워크그룹 크기, 행렬 유닛 사용 여부)를 선언한다. +- **GPU 계열별로 실기기에서 탐색한다:** Adreno 세대, Mali, Apple GPU 세대, 데스크톱. 실기기 팜이나 + CI 의 기기 러너가 필요하다. +- 결과는 **설정표로 번들에 실리고**, 런타임은 기기를 식별해 표에서 고른다. 표에 없는 기기는 보수적인 + 기본 설정과 SIMT 폴백을 쓴다. +- 탐색 결과는 커널 소스, 컴파일러 버전, 드라이버 계열을 키로 캐시한다. 소스가 바뀌지 않으면 다시 + 탐색하지 않는다. + +탐색 없이는 첫 커널이 손으로 튜닝한 코드보다 느릴 가능성이 크다. 그래서 오토튜닝은 단계 5(타일 +행렬곱)의 일부다. + +### 5.8 성능 기준선 + +이 설계는 성능 상한을 막는 장애물을 치운 것이지, 속도를 보장하지 않는다. 경쟁력은 측정으로만 +주장한다. + +- **비교 대상:** 같은 기기에서 llama.cpp(Vulkan, Metal 백엔드)와 Apple MLX 의 softmax, matmul, + attention. 모바일에서 실제 경쟁 상대는 Mojo 가 아니라 이들과 MLC-LLM, ExecuTorch 다. 데스크톱 CUDA + 에서는 Triton · TileLang 이 기준이다. +- **첫 조각부터** 측정 테스트로 박는다(AGENTS.md §14). 기록하는 값은 커널 시간과 기준 대비 비율이고, + 수치 오차는 §5.6 의 세 쪽 대조로 따로 본다. +- **목표(안):** 모바일에서 llama.cpp · MLX 의 손 튜닝 커널에 근접하는 것. 구체적인 비율은 첫 측정 뒤 + 메인테이너가 정한다. +- 측정은 AGENTS.md §8 을 따른다: 단독 실행, 부하 확인, 필터 없이. + +### 5.9 자동 융합 (torchnative #68) + +LLM 디코딩은 작은 연산이 많아 연산 사이의 디스패치와 메모리 왕복이 커널만큼 시간을 먹는다. +`torch.compile` 이득의 대부분이 융합이다. §5.6 이 미룬 그래프 캡처 기반 융합(동적 shape, in-place KV +갱신 문제)은 torchnative #68 에서 다룬다. 그전까지 모델 전체 수준의 속도 경쟁력은 제한된다. + +### 5.10 선례, Taichi 파이썬 문법으로 Vulkan / Metal 커널을 AOT 컴파일해 모바일에 실었던 가장 가까운 선례다. 지금은 활동이 크게 줄었다. 기술보다는 **독자 런타임과 독자 생태계로 고립된 것**이 원인으로 보인다. @@ -426,9 +463,9 @@ TypedPython 은 그래서 자체 텐서 런타임을 만들지 않고 **torch op | 1b | 빌드 기본 단계로 연결 | PythonMultiplatform Gradle 플러그인이 사용자 파이썬 소스에 게이트를 항상 돌리고, 에러면 빌드 실패 | 1 | | 2 | typed IR → C 백엔드 (#41) | 프론트엔드·검증기·C 생성·런타임, 데스크톱에서 끝까지 | 없음 (pypackpack 슬롯은 3 의 의존) | | 3 | Kotlin 직접 호출 | 스텁 타입 기반 C ABI 호출 생성 | 스텁 배선 완료 | -| 4 | 첫 커널 조각 | f16 원소별 연산 또는 softmax 하나를 SPIR-V 로 컴파일해 torchnative `vulkan.rs` 로 실행, 기기에서 일반 파이썬 실행 · torchnative CPU 와 수치 일치 (추정 2~4 주) | torchnative 의 임의 SPIR-V 실행 진입점 | +| 4 | 첫 커널 조각 | f16 원소별 연산 또는 softmax 하나를 SPIR-V 로 컴파일해 torchnative `vulkan.rs` 로 실행, 기기에서 일반 파이썬 실행 · torchnative CPU 와 수치 일치, llama.cpp 대비 시간 기록 (§5.8) (추정 2~4 주) | torchnative 의 임의 SPIR-V 실행 진입점 | | 4b | Metal 조각 | 같은 커널을 MSL 로 직접 내 iOS · macOS 에서 수치 일치 | 4, candle-metal 연결 방식 (§8) | -| 5 | 타일 행렬곱 | 행렬 유닛(cooperative matrix, simdgroup_matrix)을 쓰는 타일 matmul, 없는 기기는 SIMT 폴백 (추정 2~4 개월, MLIR 패스 사용 시 1~3 개월) | 4, 4b | +| 5 | 타일 행렬곱 | 행렬 유닛(cooperative matrix, simdgroup_matrix)을 쓰는 타일 matmul, 없는 기기는 SIMT 폴백, 빌드 타임 오토튜닝 (§5.7) (추정 2~4 개월, MLIR 패스 사용 시 1~3 개월) | 4, 4b | | 6 | 착지 | 커널을 `torch.library` op · `kernels` 번들 항목으로 등록, `transformers` 무수정 사용 | 5 | | 7 | CPU SIMD | `@compiled` C 백엔드의 벡터 확장 · NEON 경로 (§5.3) | 2 | @@ -470,6 +507,10 @@ upstream 대비 수치 오차와 커널 시간을 기준선으로 박아 둔다. 단위로 삼아 범위를 제한한다. - **행렬 유닛의 실제 지원 범위.** Adreno 는 신형 칩에서 확장 두 개가 필요하고 Mali 는 대부분 없다. SIMT 폴백이 기본 경로가 되는 기기가 많을 수 있다. 기기별 측정 전에는 성능을 약속하지 않는다. +- **오토튜닝에 필요한 실기기.** GPU 계열별 탐색(§5.7)은 기기 팜이나 기기 러너 없이는 못 한다. + 확보 전에는 계열 하나(개발 기기)의 설정표로 시작한다. +- **속도 경쟁력은 아직 주장할 수 없다.** 설계는 상한만 열었다. §5.8 의 기준선 측정 전에는 성능을 + 약속하지 않고, 자동 융합(§5.9)이 없으면 모델 전체 속도는 제한된다. - **MLIR 휠의 변동.** nightly · alpha 이고 3.15t 휠이 없다. 선택적 보조로만 두고 날짜 고정 휠을 쓴다. 빠져도 SPIR-V 를 직접 내는 경로가 남아야 한다 (§5.4). - **기본 검사의 오탐.** 검사가 기본값이므로, 오탐 하나가 모든 사용자의 빌드를 막는다. 1b 를 켜기 전에