Skip to content
Merged
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
49 changes: 45 additions & 4 deletions docs/design/typedpython.md
Original file line number Diff line number Diff line change
Expand Up @@ -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 컴파일해 모바일에 실었던 가장 가까운 선례다. 지금은
활동이 크게 줄었다. 기술보다는 **독자 런타임과 독자 생태계로 고립된 것**이 원인으로 보인다.
Expand All @@ -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 |

Expand Down Expand Up @@ -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 를 켜기 전에
Expand Down
Loading