diff --git a/docs/INTENT.md b/docs/INTENT.md index 50fe642dc..cb707781c 100644 --- a/docs/INTENT.md +++ b/docs/INTENT.md @@ -40,9 +40,23 @@ CPython interpreter inside a Kotlin Multiplatform application so that: is **checked** by default; **compiling** it to native code is opt-in. The compiled code still runs on the embedded CPython, as an extension module (`docs/design/typedpython.md`). - > The maintainer also asked to *explore* GPU kernels written in Python without a separate language - > ("GPU 관련 최적화까지도 별도 언어 없이 지원할 방안을 모색해보자"). That is an exploration, not yet - > a commitment; it is recorded in the design note only. +5. **High-performance kernels are written in Python too** (TypedPython tier 2). The maintainer first + asked to explore it ("GPU 관련 최적화까지도 별도 언어 없이 지원할 방안을 모색해보자") and on + 2026-10-04 made it a commitment ("GPU 커널 등 탐색을 약속으로 올려"), so that TypedPython can stand + where Mojo stands without leaving Python syntax: + + - **GPU kernels** are valid Python functions marked `@kernel`, compiled **ahead of time** for Vulkan + (Android, desktop), Metal (iOS, macOS), WebGPU (web) and CUDA (desktop). iOS forbids runtime code + generation, so nothing is compiled on the device. + - **CPU SIMD** is part of the `@compiled` tier: vectorised native code for NEON and AVX from the + same valid Python. + - A kernel still runs unmodified on plain CPython, which is how it is debugged and how its numbers + are checked. + - Kernels land in the existing ecosystem as torch ops through torchnative, not in a separate tensor + runtime. + + The backend choice is in `docs/design/typedpython.md` §5; the evidence is in + `docs/investigations/typedpython-kernel-backend.md` and `typedpython-kernel-ir.md`. The CPython version moved on from 3.13 (the build pins 3.14.x today and offers opt-in free-threaded 3.14t builds on desktop only); the intent, embed **CPython itself**, not a re-implementation, is unchanged. diff --git a/docs/design/typedpython.md b/docs/design/typedpython.md index 05a9bf215..689458faf 100644 --- a/docs/design/typedpython.md +++ b/docs/design/typedpython.md @@ -31,7 +31,8 @@ TypedPython 의 목표는 **새 언어를 만들지 않고, 파이썬 문법 그 1. 타입이 확정된 코드는 **네이티브로 컴파일**해서 앱에 싣는다. PythonMultiplatform 의 배포 목표인 "파이썬 부분도 컴파일된 채로 들어간다" 와 같은 자리다. -2. GPU 커널을 **별도 언어(CUDA, Metal Shading Language, Mojo) 없이** 파이썬 서브셋으로 쓴다. +2. GPU 커널과 CPU SIMD 를 **별도 언어(CUDA, Metal Shading Language, Mojo) 없이** 파이썬 서브셋으로 + 쓴다. 2026-10-04 에 탐색에서 약속으로 올라갔다 (INTENT §1.5, §5). ## 2. 원칙 @@ -61,8 +62,8 @@ TypedPython 의 목표는 **새 언어를 만들지 않고, 파이썬 문법 그 |---|---|---|---| | 0 | 서드파티 패키지 (의존성) | 없음, 검사하지 않는다 | 바이트코드 (pypackpack 이 받은 그대로) | | 1 `checked` | **사용자 프로젝트의 파이썬 코드 전부, 기본값** | 빌드 타임 타입 오류, `Any` 유출 0 (§4.2) | 바이트코드 | -| 1 `compiled` | `@compiled` (내장 이름, import 없음) 또는 모듈 선언 `# typedpython: compiled` | `checked` + §3.1 금지 목록 | C 확장 모듈 | -| 2 `kernel` | `@kernel` | 객체·힙 할당 없음, 고정 dtype | SPIR-V / MSL / WGSL / PTX | +| 1 `compiled` | `@compiled` (내장 이름, import 없음) 또는 모듈 선언 `# typedpython: compiled` | `checked` + §3.1 금지 목록 | C 확장 모듈 (NEON · AVX 벡터화, §5.3) | +| 2 `kernel` | `@kernel` | 객체·힙 할당 없음, 고정 dtype | SPIR-V / MSL (직접) / WGSL / PTX, 빌드 타임 (§5) | - **사용자 코드는 표시하지 않아도 검사된다.** 타입 오류가 있으면 Kotlin 컴파일 에러처럼 빌드가 실패한다. TS 의 `strict` 기본값과 같은 위치다. @@ -322,43 +323,96 @@ PyPy·GraalVM 과 같은 방식이다. N-11 클래스의 객체를 PyObject 대 **선행 조건:** 스텁 생성기가 픽스처 대역이 아닌 실제 `pythonx-map.toml` 을 읽고, 생성된 스텁이 IDE·타입 체커가 찾는 자리(pythonx-compose 배포본)에 설치돼야 한다. 2026-10 현재 미완료. -## 5. 등급 2, 별도 언어 없는 GPU +## 5. 등급 2, 별도 언어 없는 커널 + +> **결정 (2026-10-04, 메인테이너):** GPU 커널과 CPU SIMD 는 탐색이 아니라 **약속**이다 (INTENT §1.5). +> 아래 백엔드 결정의 근거는 [`typedpython-kernel-backend.md`](../investigations/typedpython-kernel-backend.md) +> 와 [`typedpython-kernel-ir.md`](../investigations/typedpython-kernel-ir.md) 에 있다. 두 조사는 실제 기기 +> 성능을 재지 않았다. 아래 "측정" 표기가 붙은 항목은 첫 조각에서 확인한다. ### 5.1 하드 제약: AOT **iOS 는 런타임 코드 생성을 금지한다.** torchnative `docs/design/DESIGN.md` 는 이 이유로 Triton 을 기각했다. Android 에도 인앱 GPU 컴파일러 툴체인이 없고, autotuning 워밍업은 짧은 세션에서 상각되지 -않는다. 따라서 **GPU 커널은 빌드 타임에 컴파일된다.** JIT 은 데스크톱 개발 편의로만 허용한다. +않는다. 따라서 **커널은 빌드 타임에 컴파일된다.** JIT 은 데스크톱 개발 편의로만 허용한다. 컴파일러는 +빌드 도구에만 들어가고 앱에는 컴파일된 커널만 실린다. 그래서 컴파일러 쪽 의존성의 무게는 앱 크기가 +아니라 툴체인의 문제다. -### 5.2 커널 모델 +### 5.2 커널 모델: 타일 우선, SIMT 는 폴백 -Triton · NVIDIA Warp 처럼 파이썬 함수에 `@kernel` 을 붙인다. +파이썬 함수에 `@kernel` 을 붙인다. -- **처음에는 SIMT** (스레드 인덱스 기반). 컴퓨트 셰이더에 그대로 대응된다. -- 타일 추상화(Triton · cuTile 식)는 그다음. -- 커널 안: 고정 dtype 스칼라, 버퍼·텐서 뷰, 산술, 제어 흐름, 공유 메모리, 배리어. +- **기본은 타일 수준이다.** 타일 복사, 타일 행렬곱, 단계별 파이프라이닝, 공유 메모리·레지스터 조각 + 할당을 1급 연산으로 둔다 (TileLang 의 `T.copy` · `T.gemm` · `T.Pipelined` 와 같은 층위). AI 커널의 + 성능은 행렬 유닛(텐서 코어, Apple simdgroup_matrix · Metal 4 cooperative tensor, Vulkan + `VK_KHR_cooperative_matrix`)과 메모리 계층을 쓰는 데서 나오고, 그것을 표현하는 단위가 타일이다. +- **SIMT(스레드 인덱스 기반)는 폴백이자 저수준 탈출구다.** 행렬 유닛이 없는 기기가 실제로 많다 + (Mali 대부분, 구형 Adreno). 타일 연산은 행렬 유닛이 없으면 SIMT 코드로 낮춰진다. +- 커널 안: 고정 dtype 스칼라(f32, f16, bf16, int8 포함), 버퍼·텐서 뷰, 산술, 제어 흐름, 배리어. 객체 생성·힙 할당·예외·파이썬 호출은 없다. 위반은 등급 1 과 같은 게이트에서 컴파일 에러. -- 커널 함수도 유효한 파이썬이다. CPU 에서 참조 구현으로 실행해 수치를 대조할 수 있다. - -### 5.3 컴파일 경로 - - @kernel 서브셋 → 자체 소형 IR → SPIR-V ─┬─ Vulkan (Android, 데스크톱) - ├─ SPIRV-Cross → MSL (iOS, macOS) - └─ naga → WGSL (웹) - └──────────── PTX (데스크톱 CUDA, 별도) - -한 소스에서 모바일·데스크톱·웹 타깃이 나온다. - -### 5.4 착지점, torchnative - -컴파일된 커널을 **`torch.library` 커스텀 op 로 등록**하고 mps / vulkan 구현을 붙인다. 그러면 -`transformers` 코드가 수정 없이 그 커널을 쓴다. - -torchnative 의 설계는 이미 이 자리를 비워두고 있다. 융합 어텐션 같은 커널 계층을 "prefill 문제의 -예정된 해법" 으로 지목했고, 커널 소스 후보로 손으로 쓴 Metal/NEON, CubeCL, AOT Triton 을 올려 -두었다. TypedPython 등급 2 는 **같은 자리를 파이썬으로 채우는 네 번째 후보**다. - -### 5.5 선례, Taichi +- **커널 함수도 유효한 파이썬이다.** 타일 연산은 CPU 참조 구현을 가진 일반 파이썬 함수이고, 커널은 + CPython 에서 그대로 실행되어 수치를 대조할 수 있다. TileLang · Triton 과 갈리는 지점이다. + +### 5.3 CPU SIMD: `@compiled` 의 벡터화 + +Mojo 의 절반은 CPU SIMD 다. 등급 1 `@compiled` 의 C 백엔드(§4.3)가 이 자리를 맡는다. + +- 첫 단계는 **C 벡터 확장**(`__attribute__((vector_size))`)으로 NEON · AVX 를 내는 것이다. +- 행렬곱 같은 핫 커널은 손으로 쓴 NEON 경로를 둔다. +- 벡터 폭과 dtype 을 표현하는 타입은 등급 2 의 타일·dtype 표기와 같은 어휘를 쓴다. + +이 절은 조사가 얕다(investigations 의 [I] 표기). 첫 GPU 조각 다음에 별도로 설계한다. + +### 5.4 IR: 우리 것, MLIR 에 사상되게 + +프론트엔드(유효한 파이썬 커널을 읽는 부분)와 IR 은 우리가 만든다. "CPython 에서 그대로 실행" 원칙을 +지키는 길이 이것뿐이다. 다만 IR 은 **MLIR 의 gpu · vector · linalg 구조에 그대로 사상되게** 설계한다. +나중에 MLIR 로 옮기거나 둘 사이를 오가는 비용을 낮추기 위해서다. + +- **MLIR 파이썬 휠은 선택적 보조 도구다.** SPIR-V · PTX 직렬화, 검증, 교차 대조에 쓴다. 공식 휠은 + LLVM 의 nightly(eudsl, alpha)이고 PyPI 에 없으며, 3.14t 휠은 있으나 **3.15t 휠은 아직 없다**(조사 + 시점 2026-10-04). 그래서 3.14t 별도 프로세스에서 돌리고, 날짜를 고정한 휠을 잠금 파일에 둔다. 빠지면 + SPIR-V 를 직접 내는 쪽으로 돌아간다. +- MLIR 의 타일링 · 벡터 · transform 패스를 쓸지는 **첫 조각이 동작한 뒤** 정한다. +- 채택하지 않는 것(참고 자료로만): TileLang/TVM (Vulkan 백엔드 없음, 커널이 일반 파이썬으로 실행되지 + 않음), CubeCL (진입점이 Rust 매크로, JIT 전용), IREE (모델 단위 입력과 자체 출력 형식, SPIR-V 코드 + 생성만 떼어 쓸 공식 경로 없음), xDSL (spirv 방언·SPIR-V/MSL/PTX 출력기 없음). + +### 5.5 출력기: 대상별로 교체 가능 + + @kernel (유효한 파이썬) → 프론트엔드 → 우리 IR ─┬─ SPIR-V ──── Vulkan (Android, 데스크톱) + ├─ MSL 직접 ─ Metal (iOS, macOS) + ├─ WGSL ───── WebGPU (웹) + ├─ PTX ────── CUDA (데스크톱) + └─ C + 벡터 확장 ─ CPU (참조 · 폴백) + +- **Apple 은 MSL 을 직접 낸다.** SPIRV-Cross 를 거치면 행렬 유닛이 8x8 float simdgroup 까지만 가고, + Metal 4 텐서 연산(MSL 4 전용)에는 닿지 않는다. +- **MSL 과 WGSL 출력기는 어느 IR 을 택해도 우리가 쓴다.** MLIR 상류에 두 대상이 없다. +- SPIR-V 와 PTX 는 직접 내거나 MLIR 휠로 직렬화한다 (§5.4). + +### 5.6 런타임과 착지점: torchnative + +**별도 런타임을 만들지 않는다.** 커널은 torchnative 의 장치 계층에서 실행된다. + +- **Vulkan:** torchnative `torch_c/src/vulkan.rs` 는 이미 SPIR-V 를 싣고(`include_bytes!`), 파이프라인을 + 만들고, push constant 로 shape 을 넘겨 실행한다. 첫 조각은 glslc 로 만든 셰이더 자리에 우리 + 컴파일러가 만든 SPIR-V 를 넣는 일이다. +- **Metal:** candle-metal 경로. 우리 MSL 커널을 거기에 붙이는 방식은 torchnative 와 정한다 (§8). +- **wgpu 는 웹 전용이다.** torchnative 가 기기에서 두 방식을 나란히 재 보았다: ash 는 크레이트 4 개 + 719 KB, wgpu 는 60 개 6.77 MB 이고 torchnative 는 ash 를 권고했다(`docs/devices/VULKAN.md` §5.3). + wgpu 를 네이티브 런타임으로 쓰면 행렬 유닛이 실험 기능 뒤에 갇힌다. +- **착지:** 컴파일된 커널은 `torch.library` 커스텀 op, 또는 torchnative 가 빌드 타임으로 옮긴 HF + `kernels` 번들 리졸버의 항목으로 등록된다. 그러면 `transformers` 코드가 수정 없이 그 커널을 쓴다. + `kernels` 규격에 `vulkan` 슬롯이 없는 구멍을 이 커널이 메운다. +- **수치 기준:** torchnative CPU 경로(bf16/fp16 에서 비트 일치하는 `flash.rs` 등)가 정답지다. 커널의 + 일반 파이썬 실행, torchnative CPU, GPU 출력의 세 쪽을 대조한다. +- **자동 융합은 뒤로 미룬다.** torchnative 의 그래프 캡처로 `torch.compile` 을 대신하는 길(캡처한 + 그래프의 융합 구간을 커널 컴파일러로 넘김)은 가치가 크지만, 그 캡처는 동적 shape · in-place 변경 · + 제어 흐름을 거부하고 LLM 의 KV 캐시 갱신이 in-place 라 정면으로 걸린다. 그전까지는 리졸버를 통한 + 레이어 교체로 착지한다. + +### 5.7 선례, Taichi 파이썬 문법으로 Vulkan / Metal 커널을 AOT 컴파일해 모바일에 실었던 가장 가까운 선례다. 지금은 활동이 크게 줄었다. 기술보다는 **독자 런타임과 독자 생태계로 고립된 것**이 원인으로 보인다. @@ -372,11 +426,15 @@ TypedPython 은 그래서 자체 텐서 런타임을 만들지 않고 **torch op | 1b | 빌드 기본 단계로 연결 | PythonMultiplatform Gradle 플러그인이 사용자 파이썬 소스에 게이트를 항상 돌리고, 에러면 빌드 실패 | 1 | | 2 | typed IR → C 백엔드 (#41) | 프론트엔드·검증기·C 생성·런타임, 데스크톱에서 끝까지 | 없음 (pypackpack 슬롯은 3 의 의존) | | 3 | Kotlin 직접 호출 | 스텁 타입 기반 C ABI 호출 생성 | 스텁 배선 완료 | -| 4 | 커널 하나 | 융합 softmax 를 Vulkan · Metal 양쪽에서 upstream 과 수치 일치 | torchnative 착지 승인 | -| 5 | 커널 컴파일러 일반화 | SIMT 서브셋 전체, 이후 타일 | 4 | +| 4 | 첫 커널 조각 | f16 원소별 연산 또는 softmax 하나를 SPIR-V 로 컴파일해 torchnative `vulkan.rs` 로 실행, 기기에서 일반 파이썬 실행 · torchnative CPU 와 수치 일치 (추정 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 | +| 6 | 착지 | 커널을 `torch.library` op · `kernels` 번들 항목으로 등록, `transformers` 무수정 사용 | 5 | +| 7 | CPU SIMD | `@compiled` C 백엔드의 벡터 확장 · NEON 경로 (§5.3) | 2 | -각 단계는 AGENTS.md §5·§14 의 TDD 와 측정 테스트 규정을 따른다. 특히 3 은 경계 비용 감소를, 4 는 -upstream 대비 수치 오차를 기준선으로 박아 둔다. +각 단계는 AGENTS.md §5·§14 의 TDD 와 측정 테스트 규정을 따른다. 특히 3 은 경계 비용 감소를, 4 와 5 는 +upstream 대비 수치 오차와 커널 시간을 기준선으로 박아 둔다. 4 와 5 의 기간은 조사의 추정이지 측정이 +아니다. ### 6.1 단계 1 의 현재 상태 (2026-10-03, #16) @@ -410,6 +468,10 @@ upstream 대비 수치 오차를 기준선으로 박아 둔다. 그 목록이 2단계 IR 의 요구사항이 된다. - **등급 2 컴파일러의 규모.** 가장 큰 작업이다. 커널 하나를 끝까지 하는 것(§6 의 4)을 첫 검증 단위로 삼아 범위를 제한한다. +- **행렬 유닛의 실제 지원 범위.** Adreno 는 신형 칩에서 확장 두 개가 필요하고 Mali 는 대부분 없다. + SIMT 폴백이 기본 경로가 되는 기기가 많을 수 있다. 기기별 측정 전에는 성능을 약속하지 않는다. +- **MLIR 휠의 변동.** nightly · alpha 이고 3.15t 휠이 없다. 선택적 보조로만 두고 날짜 고정 휠을 쓴다. + 빠져도 SPIR-V 를 직접 내는 경로가 남아야 한다 (§5.4). - **기본 검사의 오탐.** 검사가 기본값이므로, 오탐 하나가 모든 사용자의 빌드를 막는다. 1b 를 켜기 전에 이 저장소의 실제 사용자 측 파이썬 코드(`sample/`, pythonx-compose 의 소비 예시)에 게이트를 돌려 에러 수와 그 원인을 기준선으로 박아 둔다. @@ -430,7 +492,10 @@ upstream 대비 수치 오차를 기준선으로 박아 둔다. 모듈 전체를 지정할 수 있다. 2. ~~고정폭 정수의 오버플로 의미~~, **결정됨:** `int` 는 i64 로 컴파일하고 오버플로 시 인터프리터 재실행으로 큰 정수로 승격(의미 보존). 3. ~~pypackpack Cython 슬롯 인터페이스 소유~~, **결정됨:** pypackpack 이 정한다. TypedPython 은 생성한 C 소스, 빌드 플래그, C 를 낸 모듈 목록을 넘긴다. 확장이 원본 소스를 품으므로 별도 `.py` 폴백은 없다(Cython 폐기 후 갱신). -4. GPU 착지를 torchnative 커스텀 op 로 할 것인가 (다른 레포 변경). +4. ~~GPU 착지를 torchnative 커스텀 op 로 할 것인가~~, **결정됨 (2026-10-04, 메인테이너):** 커널은 + torchnative 장치 계층에서 실행되고 torch op · `kernels` 번들로 착지한다 (§5.6). torchnative 와 정할 + 것은 남는다: Metal 커널을 candle-metal 에 붙이는 방식, 등록 경로를 ATen 디스패처와 `kernels` + 리졸버 중 무엇으로 정식화할지, Vulkan f16 · 행렬 유닛을 장치 계층에 넣을지. 5. ~~TypedPython 을 별도 레포로 분리할 것인가~~, **결정됨 (2026-10-02):** PythonMultiplatform 의 기본 동작. 검사는 기본, 컴파일은 opt-in. 6. 기본 검사에서 빠져나가는 방법을 둘 것인가, 예: 동적 기능이 꼭 필요한 모듈 하나를 `checked` 에서 diff --git a/docs/investigations/typedpython-kernel-backend.md b/docs/investigations/typedpython-kernel-backend.md new file mode 100644 index 000000000..b3fc7f75d --- /dev/null +++ b/docs/investigations/typedpython-kernel-backend.md @@ -0,0 +1,206 @@ +# TypedPython 커널 백엔드 조사 (2026-10-04) + +표기: **[V]** = verified (출처 URL 병기, 이번 조사에서 직접 fetch/search 결과로 확인). **[I]** = unverified / my inference. +출처 중 WebSearch 요약은 검색 엔진 요약문이므로, 숫자나 버전은 채택 전에 원문 재확인이 필요하다. + +결론의 반영: `docs/design/typedpython.md` §5 (2026-10-04 결정). IR 선택은 후속 조사 [`typedpython-kernel-ir.md`](typedpython-kernel-ir.md) 가 다듬었다. + +## 0. 조사 당시의 계획 (읽은 문서, 이후 §5 가 바뀌었다) + +- typedpython.md §2.1: 유효한 파이썬만. 커널도 수정 없이 CPython 에서 실행돼 참조 구현이 된다. +- §5: iOS 는 런타임 코드 생성 금지이므로 AOT 가 하드 제약. 경로는 `@kernel` 서브셋 -> 자체 IR -> SPIR-V -> (Vulkan | SPIRV-Cross->MSL | naga->WGSL), PTX 별도. 착지점은 torchnative 의 `torch.library` 커스텀 op. SIMT 먼저, 타일은 나중. +- torchnative DESIGN.md: Triton 기각(JIT), 후보는 손으로 쓴 Metal/NEON, CubeCL("성숙도 확인 필요"), AOT Triton. 안드로이드 GPU 는 "Vulkan/wgpu 백엔드 도입 또는 작성" 으로 비어 있다. + +## 1. TileLang (tile-ai/tilelang) + +- 백엔드 [V github.com/tile-ai/tilelang]: CUDA(주), ROCm/HIP, Ascend, Apple Metal 은 supported. LLVM CPU, CuTe DSL, WebGPU 는 README 상 experimental. v0.1.13 (2026-08-03 출시라는 검색 요약), 문서 사이트는 0.1.14. +- 불일치 [V]: 공식 docs(tilelang.com/get_started/targets.html)는 Metal, LLVM, WebGPU, C 를 모두 "Fully supported" 로 적지만 README 는 WebGPU 를 experimental 로 적는다. 문서 쪽이 낙관적이라고 보는 게 안전하다 [I]. +- Vulkan/SPIR-V: 타깃 문서에 항목 없음. 로드맵 이슈 #56 (2025-01-26 개설) 에 "planned" 로만 있고 진행 표시 없음 [V github.com/tile-ai/tilelang/issues/56]. 즉 Android Vulkan 경로는 **없다** (Mali/Adreno 불가) [V 부재 확인, 완료 여부는 이슈 본문 한계로 100% 는 아님]. +- 라이선스: MIT [V raw.githubusercontent.com/tile-ai/tilelang/main/LICENSE]. 2024-12~2025-03 기간 Microsoft 협업 부가 조항 언급이 있다. TVM 은 Apache-2.0 [I, 기억 기반]. +- AOT: README 가 "AOT export" 를 기능으로 언급 [V]. 단 무엇이 export 되는지(소스 vs 바이너리, iOS 에 싣는 형태)는 확인 못함. 타깃 문서에 `C` 타깃(소스만 출력)이 있어 소스 생성은 가능 [V]. +- 순수 파이썬 CPU 실행: LLVM CPU 타깃이 있으나 "빌드 시 USE_LLVM=ON, LLVM 15+" 필요, 검색 요약상 experimental [V 검색요약]. 파이썬 인터프리터가 커널 본문을 그대로 돌리는 참조 실행(우리 §2.1 요건)은 TileLang 의 `T.prim_func` 가 TVM script 빌더라서 **그대로는 불가** [I: 커널이 파이썬 AST 를 직접 실행하지 않고 DSL 빌더로 구성된다는 구조적 추론, 실행해 보지 않음]. +- TVM 의 단점(빌드 복잡도, 바이너리 크기, 컴파일 시간, 포크 변경, free-threading 지원): **이번 조사에서 확인하지 못함**. README 가 TVM 을 핵심 의존성으로 명시한다는 것만 [V]. 나머지는 [I] 이며 측정이 필요하다 (§7). +- 우리 요건과 맞는 점: 타일 추상화, 텐서 코어/simdgroup 접근이 1급. 어긋나는 점: 새 DSL 문법(`T.` API), "유효한 파이썬만" 은 만족하나 CPython 에서 의미 있게 실행되지는 않음, iOS/Android Vulkan 부재. + +## 2. wgpu + naga + +### 2.1 현재 상태 +- wgpu 문서 v30.0.1 [V docs.rs/wgpu/latest/wgpu/struct.Features.html]: + - SUBGROUP: Vulkan, DX12, Metal 에서 지원, "WebGPU 표준으로 곧 이동 예상". + - SHADER_F16: Vulkan, Metal, DX12, WebGPU 모두. + - EXPERIMENTAL_COOPERATIVE_MATRIX: 문서상 8x8 f32 로 제한, **native 전용** (Metal, Vulkan). + - PASSTHROUGH_SHADERS: Vulkan, DX12, Metal, WebGPU 에서 셰이더를 파싱 없이 직접 제출 가능. 즉 우리가 SPIR-V/MSL 을 직접 만들어 wgpu 런타임에만 얹는 길이 열려 있다. + - 해당 문서에 PUSH_CONSTANTS, int8 dot 항목은 보이지 않음 (이 문서 요약 기준. 부재의 증거가 아님) [V 한계 있음]. +- 협력 행렬 PR #8251: 2025-12-22 merge, Vulkan(KHR_cooperative_matrix)과 Metal(simdgroup) 대상, 후속 PR 로 정수(i32/u32/i8/u8)와 비정방 타일 [V github.com/gfx-rs/wgpu/pull/8251, /pull/9490, /pull/10414]. 문서의 "8x8 f32 만" 과 후속 PR 이 어긋나므로 어느 릴리스에 어디까지 들어갔는지는 미확정. +- WebGPU 표준: subgroups 는 Chrome 134 에서 출시 [V developer.chrome.com/blog/new-in-webgpu-134], Chrome 144 에서 subgroup_id/num_subgroups 추가 [V 검색요약]. subgroup matrix 제안은 2026 에도 활발히 수정 중(2026-09 D3D12 반영 PR) 이므로 **표준 미확정, 브라우저 출시 전** [V github.com/gpuweb/gpuweb/issues/4195, PR #10964; 출시 여부는 확인 못함]. +- naga: subgroupElect 외 subgroup 연산 지원 [V github.com/gfx-rs/wgpu/issues/7396, 이슈 시점 기준]. + +### 2.2 성능 증거 +- Burn 블로그 요약: Vulkan(SPIR-V) 컴파일러가 WGSL 보다 matmul 에서 크게 낫다. 이유는 SPIR-V 쪽에서 텐서 코어와 f16 을 쓸 수 있기 때문 [V 검색요약, burn.dev/blog/sota-multiplatform-matmul]. 같은 맥락에서 "CUDA/Vulkan 이 WebGPU 보다 prefill 최대 10x, decode 2.5x 빠르고 Metal 이 WebGPU 보다 prefill 2x 이상" 이라는 문장이 나왔으나, 어느 모델/논문의 수치인지 출처 귀속이 불분명하다 (Llamas on the Web arXiv 2605.20706 로 추정되나 PDF 본문에서 수치를 직접 확인하지 못함) [V 검색요약, 귀속 미확인]. +- Meganeura (arXiv 2608.01563): wgpu 를 거치지 않고 Vulkan/Metal 로 직접 컴파일하며, wgpu 층의 오버헤드를 이유로 든다고 요약됨 [V 요약, 수치 미확인]. +- 결론 [I]: ML 커널에서 "wgpu 의 WGSL 경로" 는 협력 행렬/int8/f16 접근에서 native 에 뒤진다. wgpu 를 **런타임** 으로만 쓰고 셰이더는 PASSTHROUGH 로 넣는 것이 중간 길이다. + +### 2.3 모바일, torchnative 공존 +- Android Vulkan, iOS Metal 은 wgpu 가 지원하는 백엔드 [I: 일반 지식, 이번에 문서로 재확인하지 않음]. +- torchnative 의 `mps`/`vulkan` 은 torch 디바이스이고, wgpu 는 별도 디바이스/버퍼 소유권을 가진다. 같은 MTLBuffer/VkBuffer 를 공유하려면 interop 이 필요하다 [I]. 커스텀 op 가 자체 Vulkan/Metal 호출을 쓰는 편이 버퍼 공유가 쉽다 [I]. + +## 3. CubeCL (tracel-ai/cubecl) + +- 백엔드: CUDA, ROCm(HIP), wgpu 를 통한 Vulkan(SPIR-V)/WebGPU(WGSL)/Metal, CPU(SIMD). 라이선스 MIT/Apache-2.0. 상태 "Alpha, Burn 이 프로덕션 사용" [V github.com/tracel-ai/cubecl]. +- 네이티브 Metal: cubecl-metal 이 0.11.0-pre.4 에 추가되고 MSL 경로가 WGSL 폴백과 병존. 관련 이슈들(컴파일 실패, 폴백 혼선)이 최근 열려 있어 **이 경로는 아직 불안정** [V 검색요약, github.com/tracel-ai/cubecl/pull/1734, /issues/1722, /issues/1730]. +- 타일/매트릭스: cubek/cubecl 의 matmul 이 텐서 코어, double buffering, 벡터화로 cuBLAS/CUTLASS 와 경쟁(RTX 4080 에서 Simple 변형이 대체로 앞섬). Apple M2 Pro 에서는 Specialization 변형이 Metal 의 plane 제어 부족으로 뒤짐 [V burn.dev/blog/sota-multiplatform-matmul]. +- AOT: 공식 입장이 "JIT, 실제 실행되는 변형만 생성" [V github 요약]. 즉 **iOS 에서 그대로는 불가**. 다만 IR 에서 소스(WGSL/MSL/SPIR-V)를 빌드 타임에 뽑아 저장하는 길은 코드 구조상 가능해 보인다 [I, 확인 안 함]. +- 파이썬 프런트엔드 -> CubeCL IR: CubeCL 의 입구는 Rust `#[cube]` 매크로가 IR 빌더 호출로 바꾸는 것 [V]. IR 은 Rust 크레이트 내부 API 라 파이썬에서 직접 만들려면 Rust 쪽에 IR 을 직렬화하는 얇은 계층을 우리가 써야 한다 [I]. 또 comptime 특수화는 Rust 제네릭에 묶여 있어 파이썬 쪽 표현이 어렵다 [I]. +- CPU 참조 실행: CubeCL 자체 CPU 런타임이 있으나 이는 Rust 쪽이고 파이썬 인터프리터 실행이 아니다 [I]. + +## 4. MLIR 계열 (Triton 스택, IREE, Mojo) + +- Triton: JIT 이 전제. Apple 용 외부 백엔드(triton-metal, triton-msl, triton-ext PR #126/#127)가 존재하며 MSL 을 생성하고 upstream test_core 를 통과한다고 주장하나, 개인 프로젝트이며 fp8/mxfp 행렬 유닛 없음, 1024 스레드/32KB threadgroup 상한 등 하드웨어 한계가 있다 [V 검색요약, github.com/bledden/triton-msl]. Vulkan/SPIR-V 백엔드는 이번 검색에서 찾지 못함 [V 부재는 검색 한계]. 이전 조사(FlexLA, ICLR 2026)는 Triton 위 AOT 디스패처 선례 [V torchnative DESIGN.md 인용]. +- IREE: AOT 중심, Vulkan/Metal/CUDA/CPU, "Android, iOS, 베어메탈 지원", 런타임은 임베디드에서 약 30KB 까지 작아질 수 있음 [V 검색요약, iree.dev]. 그러나 이는 최소 구성이고 GPU 백엔드 포함 모바일 앱의 실제 크기와 컴파일러 빌드 부담은 확인 못함 [I]. 입력 단위가 모델(그래프) 이라, 단일 커널 DSL 을 임베드하려면 linalg/TOSA/StableHLO 로 내려야 한다 [I]. +- Mojo: Apple GPU 는 nightly 에서 초기 지원, 경로가 Mojo -> LLVM IR -> AIR 비트코드 -> metallib. MAX 그래프/모델은 아직 Apple GPU 에서 미지원(조사 시점 검색 요약) [V forum.modular.com/t/apple-silicon-gpu-support-in-mojo/2295]. Modular 25.6 이 NVIDIA/AMD/Apple 통합을 표방 [V modular.com 블로그 제목]. 모조의 SIMD 구현 세부(SIMD[dtype,width] 1급 타입 등)는 이번에 재확인하지 않음 [I, 일반 지식]. +- 이 계열의 공통 한계: 컴파일러 빌드/배포가 무겁고, 우리 "파이썬 AST -> IR" 프런트엔드를 MLIR 로 올리려면 MLIR 바인딩을 끌어와야 한다 [I]. + +## 5. 자체 소형 IR -> SPIR-V (현재 계획) + +만들어야 하는 것 [I, 다른 시스템 경험에 근거한 추정. 규모는 측정 아님]: +1. 프런트엔드: `@kernel` 파이썬 AST 서브셋 -> 타입된 IR. (등급 1 의 Pyrefly 게이트와 공유 가능) +2. SIMT 코드 생성: SPIR-V 를 직접 emit (rspirv 같은 크레이트나 직접 binary writer). 구조화된 제어 흐름 규칙, 메모리 모델, 용량 한정자 처리가 필요. +3. 공유 메모리, 배리어, 원자 연산, 서브그룹 연산. +4. 행렬 유닛: SPIR-V `OpCooperativeMatrix*KHR` 직접 emit (Vulkan 경로). Metal 은 SPIRV-Cross 가 simdgroup_matrix 로 내려주는지 불확실. SPIRV-Cross 의 협력 행렬 지원은 이번에 확인 못함 [I]. 한편 naga 는 협력 행렬 입력 WGSL, 출력 SPIR-V/MSL/WGSL 이 구현 중 [V wgpu PR #8251]. +5. 타일 추상화를 올리려면 레이아웃 추론, 소프트웨어 파이프라이닝(cp.async 대응), 레지스터/공유 메모리 swizzle, autotuning 대체물(AOT 이므로 빌드 타임 탐색)이 필요. 이것이 Triton/TileLang 이 수년 투자한 부분. +6. Metal 4 `matmul2d`/cooperative_tensor 는 SPIR-V 경로로 접근 불가 (MSL 4 전용, MPP 라이브러리) [V WWDC26 330, developer.apple.com]. Apple Neural Accelerator(M5) 를 쓰려면 SPIR-V -> MSL 변환이 아니라 MSL 직접 emit 이 필요하다 [I]. +7. Vulkan 쪽 협력 행렬 하드웨어: Adreno 는 VK_KHR_cooperative_matrix + VK_QCOM_cooperative_matrix_conversion 이 있는 신형에서만 llama.cpp 가 활성화 중, Mali 는 일부(PanVK) 외에는 없음으로 보고됨 [V llama.cpp PR #29328 검색요약]. 따라서 **Android 에서는 협력 행렬이 기본값이 아니라 선택 경로**이고 SIMT/서브그룹 폴백이 반드시 필요하다. + +## 6. CPU SIMD 쪽 + +- 우리 `@compiled` 는 자체 typed IR -> C (typedpython.md §4.3). 선택지 [I, 이번에 재조사 안 함]: + a. C 벡터 확장 (`__attribute__((vector_size))`, clang/GCC 공통, MSVC 불가). 소스 생성이 단순하고 NEON/AVX 로 내려감. 폭이 고정이라 SVE 류는 못 씀. + b. Highway: 런타임 디스패치와 가변 폭 지원. C++ 템플릿이라 C 생성기와 결합하려면 C++ 로 emit. 라이선스 Apache-2.0/BSD-3 [I, 기억 기반]. + c. LLVM 직접 emit: 가장 강력하나 의존성 최대. iOS 앱에 LLVM 을 싣는 것은 비현실적이고 AOT 빌드 머신에서만 쓸 수 있음. 빌드 타임 도구로서는 가능. + d. 모조 방식: `SIMD[dtype,N]` 을 1급 타입으로 두고 LLVM 에 내림 [I 일반 지식]. +- 권고(추정): a 를 기본, 핵심 BLAS 류는 손으로 쓴 NEON (torchnative 가 이미 후보로 둠), 성능이 필요한 곳에만 b. + +## 7. 비교표 + +범례: O 가능(검증), o 가능 추정, X 불가, ? 미확인. + +| 항목 | TileLang | wgpu+naga | CubeCL | MLIR(IREE/Triton/Mojo) | 자체 IR->SPIR-V | +|---|---|---|---|---|---| +| iOS AOT | ? (export 있으나 iOS 형태 미확인) | o (PASSTHROUGH 로 MSL/SPIR-V 제출 가능; 런타임은 Rust 크레이트) | X 기본(JIT). 소스 추출 개조 필요 [I] | IREE o (AOT 전제). Triton X (JIT). Mojo ? | O 설계 목표 | +| Android Vulkan | X (SPIR-V 미구현) | O (wgpu Vulkan 백엔드) | O (SPIR-V 컴파일러, JIT) | IREE O, 그 외 ? | o (직접 emit) | +| Metal | O 소스 생성(macOS, iOS ?) | O (naga MSL) | O 네이티브 cubecl-metal 신규/불안정 | IREE o, Triton 외부 포크, Mojo nightly | o (SPIRV-Cross 또는 MSL 직접) | +| WebGPU | 실험 | O (WGSL) | O (WGSL, SPIR-V 대비 느림) | IREE ? | o (naga) | +| CUDA | O (주력) | X (wgpu 에 CUDA 없음) | O | Triton/Mojo O | PTX 별도 작업 | +| 행렬 유닛 | O (CUDA/HIP/Metal) | 실험(협력 행렬, native 전용, WebGPU 표준 미확정) | O (SPIR-V/CUDA 경로, WGSL 경로 약함) | O (Triton/Mojo) | 직접 구현 필요 | +| 타일 추상화 | O | X | 부분(cubek 라이브러리 수준) | Triton O | X (뒤로 미룸) | +| 순수 파이썬 CPU 참조 실행 | X [I] | 해당 없음 | 해당 없음 | Triton 은 interpreter 모드 있음 [I 기억] | O (§2.1 설계) | +| 의존성 무게 | 큼(TVM) [I] | 중(Rust) | 중(Rust) | 큼 | 작음(우리 코드) | +| 라이선스 | MIT (TVM Apache-2.0 [I]) | MIT/Apache-2.0 [I] | MIT/Apache-2.0 [V] | IREE Apache-2.0 [I], Mojo 폐쇄적 부분 있음 [I] | 해당 없음 | +| 성숙도 | 활발, 릴리스 0.1.x | 활발, 협력 행렬 신기능 | Alpha, 변화 빠름 | IREE 성숙, Mojo Apple 초기 | 0 | + +## 8. 각 선택지가 **못 하는** 것 + +- TileLang: Android Vulkan 불가, 파이썬 인터프리터 참조 실행 불가(추정), TVM 빌드 체인을 우리 Gradle/uv 배포에 넣는 부담, iOS AOT 형태 미확인. +- wgpu/naga: WebGPU 표준에 협력 행렬 아직 없음, WGSL 경로는 int8 dot / 협력 행렬 / 큰 공유 메모리에서 native 에 뒤짐(Burn 사례), Metal 4 tensor ops 접근 불가, CUDA 불가. 런타임이 Rust 라 torchnative 의 버퍼와 별도 소유. +- CubeCL: JIT 전제로 iOS 불가(개조 전), Rust 가 소스 언어라 파이썬 프런트엔드 입구가 없음, Alpha, Metal 네이티브 경로 불안정. +- MLIR: 무거운 빌드, 커널 단위 DSL 아님(IREE), Triton 은 JIT, Mojo 는 우리가 소유할 수 없는 폐쇄 도구. +- 자체 IR: 타일 수준 최적화(파이프라인, 레이아웃)를 전부 직접 만들어야 하고, 벤더 드라이버 버그와 기기별 튜닝은 어떤 선택지도 대신해 주지 않는다. Apple Neural Accelerator 는 SPIR-V 를 거치면 못 쓴다. + +## 9. 권고 (불확실성 명시) + +1. **프런트엔드와 IR 은 자체 소유를 유지한다.** 이유: §2.1(참조 CPU 실행 = 파이썬 그대로)을 만족하는 후보는 이것뿐이고 [I], 등급 1 의 게이트와 공유된다. 외부 시스템은 모두 자체 DSL 이 입구라서 §2.1 과 충돌한다 [V 각 DSL 의 입구 확인, 충돌 판단은 I]. +2. **백엔드는 한 개로 고정하지 말고 emitter 를 교체 가능한 계층으로 둔다.** 1단계는 SIMT 로 SPIR-V(Vulkan, Android)와 MSL(Apple) 두 개를 직접 emit. SPIRV-Cross 경유보다 MSL 직접 emit 을 권한다(Metal 4 tensor ops, simdgroup_matrix 에 닿으려면 필요) [I]. +3. **런타임은 먼저 wgpu 를 쓰지 않는다.** torchnative 의 mps/vulkan 디바이스와 버퍼를 공유해야 하는 커스텀 op 라는 위치 때문에, 자체 얇은 Vulkan/Metal 디스패처가 맞을 가능성이 높다 [I]. 웹 타깃에서만 wgpu/WebGPU 를 쓰고, 이때 WGSL 의 성능 한계를 감수한다. 협력 행렬이 필요한 웹 경로는 표준 확정 이후. +4. **CubeCL 은 "라이브러리 수준 참조" 로 쓰고 의존하지 않는다.** cubek matmul 의 타일 알고리즘(double buffering, line size)과 벤치 결과는 우리 커널 설계의 비교 기준으로 가치가 크다 [V]. 의존하려면 AOT 개조와 파이썬 IR 직렬화가 먼저다 [I]. +5. **TileLang 은 데스크톱 CUDA/Metal(macOS) 에서 비교 기준(baseline)으로 쓴다.** 모바일 경로가 없으므로 주 백엔드로는 부적합 [V Vulkan 부재 / I 판단]. +6. 타일 추상화는 SIMT 가 돌아간 뒤. 처음부터 Triton 급을 목표로 하면 §5 의 항목 5 가 프로젝트를 잡아먹는다 [I]. +7. CPU SIMD 는 C 벡터 확장으로 시작 [I]. + +불확실성: (a) TVM/TileLang 의 실제 비용을 측정하지 않았음, (b) wgpu 협력 행렬의 릴리스별 현황과 성능, (c) Adreno/Mali 협력 행렬 보급률, (d) 직접 emit 한 SPIR-V 가 기기별 드라이버에서 얼마나 안 깨지는지, (e) SPIR-V -> MSL 에서 협력 행렬이 내려가는지(SPIRV-Cross 미확인). + +## 10. 결정을 위해 필요한 측정 + +1. TileLang 을 `pip`/`uv` 로 설치한 크기, 소스 빌드 시간, Python 3.13/3.14t 지원 여부, 소스 export 로 얻는 MSL/CUDA 가 iOS 앱에 링크 가능한지(실제로 export 후 Xcode 에서 컴파일). +2. 동일 커널(GEMM 1024^3 f16, softmax, 융합 어텐션) 을 (i) 손 작성 MSL (ii) 우리 SPIR-V->SPIRV-Cross MSL (iii) naga MSL (iv) CubeCL Metal 로 M-시리즈/A-시리즈에서 GFLOPS 비교. Apple 은 simdgroup_matrix 와 Metal 4 matmul2d 둘 다. +3. Android: Adreno 7xx/8xx, Mali G7xx 에서 VK_KHR_cooperative_matrix 노출 여부와 SIMT 폴백 대비 속도(실기 3종 이상). +4. wgpu PASSTHROUGH + 우리 SPIR-V 와 직접 Vulkan 디스패치의 호출당 오버헤드(소형 커널 반복 10만 회), 그리고 torchnative 텐서 버퍼 공유 가능 여부. +5. 직접 emit SPIR-V 의 spirv-val, 드라이버 크래시율(가능한 많은 기기/Android 버전, CI 팜). +6. CPU: C 벡터 확장 vs 손 NEON vs Highway 로 dot/GEMV/q4 matmul 비교(torchnative 의 기존 커널 기준). 측정은 규칙 8 에 따라 단독, 비필터로. +7. 자체 IR 구현 비용 추정을 위한 프로토타입: SIMT vector-add + reduction + 공유 메모리 GEMM 을 SPIR-V 와 MSL 로 emit 하는 데 걸리는 기간. + +## 11. 이 조사 방법이 찾지 못하는 것 + +웹 검색 요약과 문서 페이지 요약(작은 모델이 변환한 결과)에 의존했으므로, 실제 성능 수치, 실기 드라이버 버그, 빌드해 보아야 보이는 의존성 문제, 요약기가 누락하거나 잘못 적은 세부(특히 PDF 본문 수치)는 찾지 못한다. + +## 12. torchnative 를 커널 컴파일러의 일부로 보기 + +표기는 위와 같다. 이 절의 **[R]** = verified (torchnative 저장소의 파일, 이번에 직접 읽음. 경로 병기, 실행하지는 않았으므로 문서가 주장하는 측정값을 재현한 것은 아님). **[V]** 는 웹 출처, **[I]** 는 추론. + +### 12.0 읽고 나서 바뀌는 사실 (§2, §5 의 전제 수정) + +- torchnative 에는 **이미 자체 Vulkan 경로가 있다.** `torchnative/rust/torch_c/src/vulkan.rs` (3801줄), `ash`(loaded, dlopen) 기반, GLSL `.comp` 를 빌드 머신에서 `glslc` 로 컴파일해 `.spv` 를 **저장소에 체크인**하고 `include_bytes!` 로 싣는다. `torch_c/shaders/` 에 add, matmul, bmm, softmax, layer_norm, gelu, embedding, gather 등 f32 셰이더가 있다 [R: torch_c/shaders/compile.sh 와 목록]. +- 셰이더는 shape 마다 재생성하지 않는다. 모양은 push constant 로 넘기고 셰이더 하나가 모든 크기를 처리한다 [R: docs/devices/VULKAN.md §3]. +- `ash` vs `wgpu` 를 **기기(에뮬레이터)에서 둘 다 돌려 비교**했다. 둘 다 통과, ash 는 크레이트 4개 719 KB / 클린 빌드 8.5 s, wgpu 는 60 크레이트 6.77 MB / 24 s. 권고는 ash. 단 "Apple GPU 를 어떻게 채우나" 가 걸려 있어, candle 의 metal 을 켜면 ash, 커널을 직접 소유하면 wgpu 가 맞다고 명시하고 결정은 조율 세션에 남겼다 [R: docs/devices/VULKAN.md §4.1, §5.1-5.3]. 이 조사의 질문(커널을 우리가 소유)은 바로 그 두 번째 칸이다. +- 단 wgpu 측정은 `Limits::downlevel_defaults()` 로 돌렸고 이는 기기 한계가 아니라 wgpu 하한, 버퍼는 최대 4 MB 였다 [R: VULKAN.md 약 381행]. 큰 텐서·협력 행렬·f16 에서의 비교는 아직 없다. +- Apple 쪽은 `mps` 디바이스가 candle Metal 위에 있고, candle Metal 에는 조용한 CPU 폴백이 없으며 이 크레이트의 커널 중 54개 op 이 호스트 되읽기라 `NotImplementedError` 로 거절한다 [R: docs/devices/MPS.md]. `Cargo.toml` 은 Apple 에서 `accelerate` 만 켜고 `metal` 은 기본값으로 끈다는 주석이 있다 [R: torch_c/Cargo.toml, DESIGN.md 1150-1160]. +- 즉 §5 의 "SPIRV-Cross->MSL, naga->WGSL" 같은 **번역 계층 전략**은 torchnative 입장에서는 이미 한 번 검토되어 "가장 가벼운 것 = 직접 ash + 체크인된 SPIR-V" 로 기운 상태다. 그 방향과 정합하도록 §9 권고 3("wgpu 를 먼저 쓰지 않는다")은 근거가 강화된다. + +### 12.1 (a) 빌드 타임 그래프 캡처 -> 융합 구간 -> 커널 컴파일러 + +**상태 [R]**: `torch_c/src/capture.rs` + `docs/graph/CAPTURE.md`. `_aten_dispatch` 단일 문(door) 끝 한 줄에서 op 을 기록한다. 기록 = (op 이름 `aten..`, 피연산자 shape/dtype/device, 결과), 가드 = 입력의 **정확한** shape/dtype/device, 재생 = eager 와 대조. FX/ExportedProgram 과 1:1 대응이라 ExecuTorch 로 갈 수 있는 모양이라는 판단이 문서에 있다 (단 "부족한 네 가지" 가 명시됨, 이번에는 열람하지 않음). +**거절하는 것 [R: CAPTURE.md §4]**: 텐서 값 기반 제어 흐름(`_local_scalar_dense`), 인플레이스/별칭, **동적 shape**, 난수, 비텐서 결과, 컨테이너. 상수(가중치)는 참조로 잡아 가드하지 않음. 배치 차원은 자유롭지 않음. +**torch.compile 은 불가 [R]**: `docs/graph/COMPILE.md` 는 abi3 에서 PEP 523 프레임 평가 훅이 막힌다고 진단하고, 대안으로 `torch.export` 경로를 본다. `docs/graph/EXPORT6.md` 는 실제 transformers 40개 아키텍처 중 export 성공이 **0/40** (벽 8개 중 상당수 해결, 그래도 헤드라인은 0/40) 이라고 적는다. + +이 역할이 필요로 하는 것과 막는 것: + +1. **융합 패턴 선정기.** 기록에서 직선 구간을 찾아 "gated MLP", "RMSNorm+residual", "SDPA", "dequant+matmul" 같은 패턴을 추출. `NPUFUSE.md` 가 Intel NPU 에서 gated MLP 를 OpenVINO 그래프 하나로 낮춘 선례다 (단 NPU 에서는 아직 실행 안 됨, 문서가 스스로 명시) [R]. 이것은 새 컴파일러가 아니라 **패턴 매처**로 시작하는 게 현실적이다 [I]. +2. **동적 shape (LLM 디코드).** 막힘의 핵심. 디코드는 KV 길이가 매 스텝 변하고, 현재 캡처 가드는 정확한 shape 를 요구한다 [R]. 외부 사례: ExecuTorch 는 LLM 에서 sequence 차원을 dynamic 으로 표시하며, 정적 shape 만 되는 백엔드(예: QNN)에는 prefill(최대 길이로 패딩)과 단일 토큰 decode 모델을 따로 export 한다 [V docs.pytorch.org/executorch/1.0/llm/export-custom-llm.html, 검색요약]. 즉 업계도 "shape 별 그래프 2~몇 개 + 패딩" 으로 해결한다. torchnative 도 같은 길(prefill 그래프, decode 그래프를 각각 캡처, KV 는 최대 길이 사전 할당 + 길이는 스칼라 입력)을 택하는 것이 가능해 보인다 [I]. AOTInductor 가 동적 shape 를 어떻게 다루는지는 이번에 원문으로 확인하지 못했다 (ExecuTorch 의 Metal/CUDA 가 AOTInductor 를 활용 중이라는 검색요약만 [V]). +3. **커널 쪽 계약.** 융합 구간을 TypedPython `@kernel` 로 쓴 커널에 매핑하려면, 커널이 **shape 를 데이터(push constant/uniform)로 받는** 형태여야 한다. torchnative Vulkan 셰이더가 이미 그렇게 한다 [R]. TypedPython 커널 시그니처에 "shape/stride 를 런타임 인자로, tile 크기는 컴파일 타임 상수" 를 요구사항으로 넣는 것이 자연스럽다 [I]. +4. **자동 코드 생성은 이 단계에서 불필요.** 융합 영역을 즉석에서 Inductor 식으로 생성하려면 pointwise/reduction 스케줄러가 필요한데 규모가 크다 [I]. 1단계는 "사람이 `@kernel` 로 쓴 융합 커널 + 캡처 기록에서 패턴 매칭으로 치환" 이면 된다. 이것은 DESIGN.md §8 의 `use_kernel_forward_from_hub` 레이어 치환과 같은 자리다 [R]. + +못 하는 것: torch.compile 의 자동 융합 대체, 데이터 의존 제어 흐름(MoE 라우팅, 조기 종료), 인플레이스 KV 갱신(변이는 거절됨, 단 KV 캐시 갱신이 정확히 그 경로 [R: DESIGN.md §4 "주의"]). **KV 캐시 인플레이스가 캡처 거절과 정면 충돌**한다 점이 가장 큰 구체적 장애다 [I: 두 문서의 결합에서 도출]. + +### 12.2 (b) torchnative 디바이스 층을 커널 런타임으로 공유 + +**상태 [R]**: Vulkan(ash) 디바이스 존재(f32 위주 셰이더, 1훈련 스텝까지 도는 단계 VULKAN10), Metal 은 candle 경유 mps, CPU 는 candle + 자체 커널. 각 디바이스에서 Metal/Vulkan 의 저장소(storage)는 torchnative 가 소유한다. + +장점 [I]: TypedPython 커널이 별도 런타임을 갖지 않고 같은 버퍼·큐·동기화를 쓰면 (1) 호스트 복사 없음, (2) 커널 소유 문제(수명, 풀링)가 한 곳, (3) 앞서 §0 이 제기한 "Taichi 식 고립" 을 피한다. +필요한 것: +- torchnative `vulkan.rs` 에 **외부 SPIR-V 를 받는 "임의 커널 디스패치" 진입점**: 파이프라인 캐시, 디스크립터 레이아웃 선언(버퍼 개수, push constant 크기), 워크그룹 크기. 현재 프로브는 디스패치마다 파이프라인·풀을 만들고 부수며 이는 "프로브라서 그럼" 이라고 문서가 명시, 커널이 몇 개를 넘으면 얇은 내부 계층이 필요하다고 스스로 적었다 [R: VULKAN.md §5.2]. +- Metal 쪽: 현재 candle-metal 이 MTLDevice/큐를 소유한다. 커스텀 MSL 을 끼우려면 candle 의 command buffer 에 인코딩하거나 같은 MTLBuffer 를 받아 별도 인코딩해야 한다. candle-metal-kernels 가 `rust/vendor/` 에 벤더링되어 있어 [R] 수정은 가능하나 버퍼 공유 방식은 코드를 읽지 않아 모른다 [I]. +- 행렬 유닛 접근: 현재 셰이더에 cooperative matrix 사용 흔적은 확인하지 못함(.comp 이름만 열람, 내용 미열람) [R 한계]. +막힘: Vulkan 셰이더가 f32 뿐이면 AI 커널 목적(f16/bf16/int4 matmul)에 부족. 이는 TypedPython 커널의 첫 번째 실제 수요처다 [I]. + +**wgpu 선택지와의 상호작용 [R + I]**: torchnative 는 wgpu 를 이미 기기에서 통과시켜 둔 프로브(`vulkan_probe` 의 `--features wgpu-route`)를 보존하고 있다 [R]. 그래서 두 런타임이 공존하면 버퍼가 둘로 갈라진다. 권고: 네이티브(Android Vulkan, Apple Metal)는 torchnative 디바이스 층(ash/candle-metal), **웹 타깃만 wgpu(WebGPU)** 로 분리하고, 두 경로는 같은 IR 에서 SPIR-V/MSL/WGSL 을 각각 emit 하는 emitter 로만 만난다. 웹에서 torchnative 를 쓰는지는 `wasm_probe` 가 있다는 것만 확인 [R]. + +### 12.3 (c) torchnative CPU 경로를 수치 오라클로 + +**상태 [R]**: `torch_c/src/flash.rs` 는 ATen CPU flash-attention 을 bf16/fp16 에서 **bit-identical** (226136개 원소 중 차이 0) 로 재현한 "참조 구현" 이고, f32/f64 는 upstream 이 BLAS 를 쓰므로 비트 일치 불가(차이 4.17e-07 수준)라고 문서가 구분한다. `vulkan_probe/src/check.rs` 는 GPU 와 CPU 참조를 같은 하니스로 비교하는 패턴이다. +역할 적합성: +- 수치 검증의 좋은 오라클이 된다. 단 **무엇과 일치해야 하는가를 먼저 정해야 한다**: TypedPython 의 순수 파이썬 참조 실행(§2.1) vs torchnative CPU(candle) vs 상류 PyTorch. 세 개의 오라클이 서로 다른 reduction 순서를 쓰므로 f32 에서는 허용오차 기반(ULP 한도 선언) 비교가 필요하다 [I, flash.rs 의 문서가 이 현상을 직접 보여줌 R]. +- 기기에서 CPU 오라클을 쓰는 것(디바이스 내 자기 대조)은 이미 프로브가 한다. 같은 방식으로 `@kernel` 하나당 (입력 시드, CPU 참조, GPU 결과) 를 출력하는 테스트 형식을 재사용할 수 있다 [I]. +- 한계: 오라클이 candle 이라면 candle 의 의미론 차이(복사 지향, view/in-place 임피던스 [R]) 가 섞인다. 커널 단위 테스트는 raw 버퍼 입력이라 영향이 작다 [I]. + +### 12.4 (d) kernels 번들 리졸버를 착지 슬롯으로 + +**상태 [R: DESIGN.md §8]**: HF `kernels` 계약(백엔드 변형 선언, transformers 무수정 교체)을 채택하고, 해석 시점을 빌드 타임으로 옮겨 Gradle 플러그인이 변형을 골라 AOT 컴파일·번들, 런타임 리졸버가 번들을 조회. 타깃 표에서 Android GPU 는 `kernels` 에 `vulkan` 백엔드가 없어 **구멍**이고, 권고는 "1) Android 는 CPU 만으로 시작 + 2) 상류에 vulkan 백엔드 제안 병행". +TypedPython 입장 [I]: +- `@kernel` 하나가 `metal` / `vulkan`(비표준) / `cpu` / `cuda` 변형 번들을 산출하는 **`kernel-builder` 대체 빌더**가 되는 구도. 이것이 §5.4 의 "`torch.library` 커스텀 op 등록" 보다 구체적인 착지 형태다. 두 개를 동시에 만족시키려면: op 등록은 torchnative 의 ATen 디스패처(`aten..`), 변형 선택은 리졸버로 하는 이중 구조가 필요할 수 있다. 어느 쪽이 정본인지는 결정 필요. +- 업스트림 `vulkan` 백엔드 슬롯 제안과 TypedPython 의 Vulkan emitter 는 한 묶음으로 움직이는 게 맞다. 이것이 DESIGN.md §8 의 선택지 2 를 실체화한다. +- 막힘: 빌드 타임 AOT 컴파일을 Gradle 플러그인이 호출해야 하므로 TypedPython 컴파일러(파이썬 프로세스) 호출 비용과 증분 컴파일(SPEC N-10)의 연계가 필요 [I]. + +### 12.5 종합 (이 절의 권고, 모두 [I]) + +1. 가장 얇은 수직 슬라이스: **`@kernel` 1개 (f16 GEMV 또는 RMSNorm) -> SPIR-V 직접 emit -> torchnative `vulkan.rs` 에 임의 SPIR-V 디스패치 진입점 추가 -> 기기에서 CPU 오라클과 대조.** 이미 있는 프로브·체크인 SPIR-V 흐름과 같은 모양이라 새 위험이 가장 적다. +2. wgpu 는 **웹 전용**으로 한정. torchnative 의 기존 측정(719 KB vs 6.77 MB)이 모바일 네이티브에서의 기각 근거다 [R]. +3. 그래프 캡처(a)는 커널 컴파일러의 선행 조건이 아니다. 먼저 레이어 치환(`use_kernel_forward_from_hub` 식)으로 가고, 캡처 기반 융합은 KV 인플레이스와 동적 shape 문제가 풀린 뒤로 미룬다. torch.export 0/40 이 그 이유다 [R]. +4. 결정이 필요한 것: (i) Apple 은 candle-metal 에 끼울지, TypedPython 이 MSL 을 emit 하고 자체 Metal 디스패처를 둘지, (ii) 정본 등록 경로(ATen 디스패처 vs kernels 리졸버), (iii) Vulkan f16/cooperative matrix 를 torchnative 디바이스 층의 범위로 볼지. + +### 12.6 추가 측정 항목 (§10 에 더해) + +8. torchnative Vulkan: 디스패치 1회당 오버헤드 (파이프라인 캐시 전/후), 동일 matmul 의 f32 GFLOPS (Adreno/Mali 실기), f16 셰이더 가능 여부. +9. candle-metal 에 외부 MSL 커널을 같은 command buffer 로 끼울 수 있는지(코드 열람 + 프로토타입). +10. 캡처: Llama 급 1개 모델의 decode 스텝을 "KV 최대 길이 고정 + 길이 스칼라" 로 쓰면 캡처 가드를 통과하는지 (인플레이스 KV 갱신 때문에 현재는 거절될 것으로 예상, 미측정). +11. 세 오라클 (파이썬 참조, candle CPU, 상류 PyTorch) 의 f32 softmax/matmul 편차 크기. + +이 절의 조사 방법이 찾지 못하는 것: 문서의 측정값을 재현하지 않았고(문서가 주장하는 숫자를 그대로 인용), `vulkan.rs` 3801줄과 candle-metal 연동 코드는 본문을 읽지 않았으며, 문서가 "해결됨" 이라 적은 항목이 현재 코드에서도 참인지는 확인하지 못한다. diff --git a/docs/investigations/typedpython-kernel-ir.md b/docs/investigations/typedpython-kernel-ir.md new file mode 100644 index 000000000..76d438456 --- /dev/null +++ b/docs/investigations/typedpython-kernel-ir.md @@ -0,0 +1,128 @@ +# TypedPython 커널 컴파일러 IR 선택 (2026-10-04) + +표기: **[V]** = verified (직접 실행한 명령 또는 fetch 한 URL), **[R]** = 로컬 저장소 파일을 직접 읽음, **[I]** = inference. +결론의 반영: `docs/design/typedpython.md` §5.4 (2026-10-04 결정). 조사 당시 작업 위치는 git 에 넣지 않은 `.tmp/research/` 였다(스크립트와 venv 는 커밋하지 않았다). venv 두 개(`venv` = py3.13 + xdsl, `venv314t` = py3.14 free-threaded + MLIR 바인딩 wheel). 스크립트 `t_spirv.py`, `t_bin.py`, `t_ptx.py`. + +## 0. 전제 (읽은 문서) + +- [R] typedpython.md §2.1: 유효한 파이썬만, 커널도 CPython 에서 그대로 실행(참조 구현). §5.1: iOS 때문에 AOT 가 하드 제약. §5.3 현재 경로: 자체 소형 IR -> SPIR-V -> (Vulkan | SPIRV-Cross->MSL | naga->WGSL), PTX 별도. 착지점 torchnative. +- [R] torchnative `vulkan.rs` (3801줄): `ash` 로 로더를 dlopen, GLSL 을 `glslc` 로 컴파일한 `.spv` 를 저장소에 체크인하고 `include_bytes!` 로 싣는다(71행). `create_shader_module`(692행) -> `create_compute_pipelines`(750행), 형상은 push constant(866행) 으로 전달, 디스패치는 `dispatch_kernel` / `dispatch_kernel_words`(788, 801행). 즉 **임의 SPIR-V 워드열을 받는 진입점이 이미 `dispatch_kernel_words` 모양으로 존재**하고, 첫 슬라이스의 도착 지점은 "`.spv` 를 컴파일러가 만든 것으로 바꾸는 것" 이다 (이전 조사 §12 의 "임의 SPIR-V 진입점이 필요" 는 이 파일의 현재 상태로 보면 부분적으로 이미 해결됐을 가능성, 단 파이프라인 캐시 상태는 읽지 않음) [R + I]. +- 이전 조사(kernel-backend.md)의 "자체 프런트엔드 + 자체 IR" 권고는 MLIR 을 내부 인프라로 평가하지 않았다. 이 문서가 그 공백을 채운다. + +프런트엔드(valid Python `@kernel` 서브셋 파싱, 등급 1 게이트와 공유)는 어느 선택지에서도 우리 몫이다. 선택지가 갈리는 곳은 **프런트엔드 뒤: IR, 변환(tiling/layout/pipelining), 코드 방출(emitter)** 이다. + +## 1. 직접 확인한 사실 (이번 조사에서 실행한 것) + +### 1.1 MLIR Python 바인딩 wheel + +- [V] PyPI 의 `mlir-python-bindings` 는 **이름 선점용 빈 placeholder** (0.0.1, "not the real distribution"). `mlir` 도 2023 년 빈 패키지. 즉 `uv add mlir-python-bindings` 는 진짜를 주지 않는다 (`curl pypi.org/pypi//json`). +- [V] 실제 wheel 은 GitHub release `llvm/eudsl` 태그 `mlir-python-bindings` 에 있다 (GitHub API 로 목록 조회: 658 assets, 최신 `mlir_python_bindings-20261004+758aa4a92-*`, 거의 매일 갱신). 설치 가이드: `pip install mlir-python-bindings -f https://llvm.github.io/eudsl` (eudsl README 요약). 이전 저장소 `makslevental/mlir-wheels` 는 2025-11-29 archive, eudsl 로 이전됨 [V WebFetch]. +- [V] 플랫폼/파이썬 매트릭스 (asset 이름 집계): cp310, 311, 312, 313, 314, **cp314t(free-threaded)** 각각 macOS arm64(`macosx_13_0_arm64`), Linux x86_64/aarch64(manylinux_2_34), Windows amd64. **cp315 / cp315t wheel 은 없음** (`cp315` 문자열 asset 0개). 크기 wheel 당 58 ~ 76 MB, 설치 후 `mlir/` 디렉터리 253 MB (macOS cp314t). +- [V] 실제로 설치해 실행: `uv venv venv314t --python 3.14t; uv pip install --python venv314t/bin/python ` 성공. `sys.version` = "3.14.7 free-threading build", `import mlir` 후 `sys._is_gil_enabled()` = **False** (GIL 이 다시 켜지지 않음). 등록된 dialect: spirv, gpu, nvvm, vector, linalg, transform, amdgpu, x86, nvgpu, emitc, llvm, scf, memref, tensor, tosa 등 (`pkgutil.iter_modules(mlir.dialects.__path__)`). +- [V] 공식 문서(mlir.llvm.org/docs/Bindings/Python/): free-threading 은 3.13+ 에서 "독립된 Context 라면 멀티스레드 안전". 안전하지 않다고 명시한 것: IR printing 과 PassManager 동시 사용, location error emit, module dump, transform interpreter, **GPU dialect 관련 작업**. 우리는 컴파일을 빌드 타임에 단일 프로세스/스레드로 돌리면 되므로 영향이 작다 [I]. +- [V] 동작 증명 1, **SPIR-V 바이너리**: `gpu.module @k [#spirv.target_env<...>] { spirv.module Logical GLSL450 requires #spirv.vce {...} }` 에 `gpu-module-to-binary` 를 돌리면 `gpu.binary` 에 `\03\02#\07` (SPIR-V magic 0x07230203) 로 시작하는 바이너리 blob 이 나온다 (`t_bin.py`). 즉 **파이썬 안에서 SPIR-V 직렬화까지 도달**. 단 `vce_triple` 없으면 "module must have 'vce_triple' attribute to be serializeable" 로 실패. +- [V] 동작 증명 2, **gpu -> spirv 하향**: `gpu.func` (global_id, memref.load, arith.addf, memref.store) 에 `convert-gpu-to-spirv` 를 돌리면 `spirv.func` 가 나온다 (`t_spirv.py`; StorageBuffer 구조체 래핑, GlobalInvocationId builtin 포함). 이어 `gpu-module-to-binary` 를 같은 모듈에 바로 돌리면 "the module has no target attributes" 로 실패: 하향 결과의 `spirv.module` 이 `gpu.module` 의 형제로 남고 target 속성이 없기 때문. **하향 산출물을 `gpu.module [target]` 안에 넣는 접합 코드는 우리가 써야 한다** (실행해 보지 않음, [I] 난이도 낮음). 이 end-to-end(f32 add -> .spv 바이트) 는 **검증하지 못했다**. +- [V] 동작 증명 3, **PTX**: `gpu.module @k [#nvvm.target]` 에 `gpu.module(convert-gpu-to-nvvm,canonicalize),gpu-module-to-binary{format=isa}` 를 돌리면 `.version 7.0 .target sm_80 ... .visible .entry f(` 로 시작하는 PTX 텍스트가 나온다 (`t_ptx.py`). wheel 에 NVPTX 백엔드가 들어 있다는 뜻. +- [V] `spirv.KHR.CooperativeMatrixLoad/Store/MulAdd/Length` op 이 파이썬 dialect 모듈(`mlir.dialects._spirv_ops_gen`)에 있음. `gpu-lower-to-nvvm-pipeline`, `convert-vector-to-spirv`, `convert-arith-to-spirv`, `convert-nvgpu-to-nvvm` 패스가 파싱됨. +- [V] 확인 못 한 패스 이름: `spirv-lower-abi-attrs`, `spirv-update-vce`, `convert-gpu-to-nvvm` (Context 없이 파싱 시도했거나 이름이 다를 수 있음; `convert-gpu-to-nvvm` 은 위 PTX 실험에서 `gpu.module(...)` 안에서는 동작). 결론 가능한 것은 "일부 패스가 파이썬에서 호출 가능" 까지. + +### 1.2 xdsl + +- [V] `uv pip install xdsl` -> xdsl 0.72.0 (PyPI 2026-10-02 릴리스). 순수 파이썬, 의존성 `immutabledict`, `ordered-set`, `typing-extensions` 셋뿐. Requires-Python >=3.10, classifier 3.14 까지. **py3.14t venv 에서도 설치/임포트되고 `sys._is_gil_enabled()` = False** [V]. 3.15t 는 시험하지 못함 (순수 파이썬이라 [I] 가능성 높음). +- [V] 라이선스: PyPI METADATA 는 "Apache License v2.0 with LLVM Exceptions". GitHub API 의 SPDX 는 NOASSERTION (표기 불일치, 사소함). GitHub star 595, 최근 push 2026-10-03, PyPI 릴리스 126 개, 최근 8 개가 2026-05-29 ~ 2026-10-02 사이 (약 2~4 주 간격) [V]. +- [V] dialect 목록 (`pkgutil.iter_modules(xdsl.dialects.__path__)`): acc, affine, arith, arm, arm_neon, bufferization, builtin, cf, comb, complex, csl, dlti, emitc, func, **gpu**, linalg, llvm, math, memref, memref_stream, ml_program, omp, pdl, ptr, riscv*, scf, snitch*, stencil, tensor, tosa, **transform**, ub, **vector**, wasm, x86*, 등. **`spirv` dialect 없음**: `grep -ril "spirv\|spir-v" site-packages/xdsl` 결과 0 건 [V]. +- [V] 백엔드 (`xdsl.backend`): `llvm` (xdsl -> LLVM IR 텍스트 변환, 약 950줄), `wgsl/wgsl_printer.py` (252줄), `riscv`, `x86`, `csl`, **`mps` 는 스텁** (`print_to_mps` 가 `NotImplementedError("MPS backend not yet implemented")`, AIR 비트코드를 목표로 한다고 docstring). `xdsl-opt` 의 대상 목록: csl, llvm, mps, riscv-asm, wgsl, x86-asm. SPIR-V, MSL, PTX 방출기는 없음. +- [V] gpu dialect(796줄)는 launch/func/module/global_id/barrier/subgroup_id 등 기본 op 위주이며, `subgroup_mma` 계열 op 이름은 상위 60 개 심볼 출력에서 확인되지 않음 [V 한계: 전체 grep 하지 않음]. 협력 행렬 관련 dialect 는 `grep` 상 없음. + +## 2. 선택지 비교 + +### 옵션 1. 완전 자체 IR + 자체 방출기 + +- 대상별 경로 [I]: + - SPIR-V: 우리 IR -> SPIR-V 워드를 직접 방출(binary writer 수백 줄 + 구조화 제어 흐름(OpLoopMerge/OpSelectionMerge), 메모리 모델, Capability/Extension 선언, 협력 행렬 `OpCooperativeMatrix*KHR`). `spirv-val` 로 검증 가능. + - MSL: 같은 IR 에서 텍스트 방출. 어느 옵션이든 MSL 은 우리가 쓴다(아래 참조), 자체 IR 에서는 **가장 자연스럽다**. + - WGSL: 텍스트 방출. SPIR-V 로 가는 것보다 단순. + - PTX: 텍스트 방출(작은 부분집합) 또는 LLVM 에 의존. 자체 PTX 방출은 레지스터 할당/주소 공간이 일. + - CPU SIMD: C(벡터 확장)로 방출, 이는 §4.3 에서 이미 정한 `typed IR -> C` 와 합칠 수 있다 [R typedpython.md §4.3]. +- 타일 변환: **전부 직접.** 레이아웃 추론, 소프트웨어 파이프라인, swizzle, autotune 대체물. 이것이 가장 큰 일이다 (이전 조사 §5.5). +- 도구 무게: 최소. 순수 파이썬 코드, 3.15t 문제 없음 [I]. +- 위험: 수년 걸리는 변환 인프라를 혼자 쌓는다. 반면 MSL/SPIR-V 직접 제어 가능, 디버깅이 쉽다. +- 첫 슬라이스 (f16 elementwise 또는 softmax -> SPIR-V -> `vulkan.rs` 디스패치 -> CPU 참조와 비교): 파이썬 AST -> 타입 IR(어차피 필요), SPIR-V 방출 약 800 ~ 1500 줄, 호스트 하니스. 추정 2 ~ 4 주(1인) [I, 측정 아님]. f16 은 `Float16` + `StorageBuffer16BitAccess` capability 필요(기기 의존). +- 타일 matmul + cooperative matrix: IR 에 tile 타입과 `mma` op, `OpCooperativeMatrix*` 방출, 공유메모리 타일링, 정밀한 레이아웃. 추정 2 ~ 4 개월 [I]. Adreno/Mali 에서는 기본값 아님(선택 경로) [V 이전 조사, llama.cpp PR 요약]. + +### 옵션 2. MLIR (upstream, 파이썬 바인딩, C++ 코어) + +- 획득 [V §1.1]: PyPI 에 없음. `llvm/eudsl` release 의 nightly wheel(매일, cp314t 포함, macOS arm64 / Linux x86_64+aarch64 / Windows x86_64). 커밋 해시가 버전에 붙는 **nightly 스냅샷**이라 안정 릴리스 번호가 없고 고정하려면 특정 날짜 wheel URL 을 lockfile 에 박아야 한다 (uv 의 `--find-links` 또는 URL 의존성). 직접 빌드는 LLVM 전체 빌드(`MLIR_ENABLE_BINDINGS_PYTHON=ON`, nanobind) 이며 시간/디스크가 크다 [I, 빌드해 보지 않음]. +- 버전 churn: 매일 갱신되는 upstream main 추종. Python 바인딩 API, dialect op 이름, 패스 옵션이 자주 바뀐다 (예: 검색에서 확인된 SPIR-V KHR CooperativeMatrix Load/Store Stride 옵션화 PR #214194, 협력 행렬 operand 수 수정 PR #118014) [V 검색요약]. 완화: 특정 날짜 wheel 로 고정하고 분기별로 올리는 정책 [I]. +- 3.15t: **wheel 없음** (cp314t 까지) [V]. 기반이 nanobind 라 소스 빌드로 3.15t 가능성은 있으나 미확인 [I]. 사용자 목표가 3.15t 이므로, 컴파일러(빌드 타임)를 3.14t 또는 GIL 파이썬 별도 인터프리터에서 돌리고 앱 런타임과 분리하는 것이 현실적 [I]. 공식 문서가 free-threading 에서 GPU dialect/PassManager 동시 사용을 안전하지 않다고 명시한 점도 단일 스레드 컴파일로 우회 [V 문서]. +- 대상별 경로: + - SPIR-V: gpu/vector/arith/memref/scf -> `convert-gpu-to-spirv` -> `spirv` dialect -> `gpu-module-to-binary` 로 직렬화 [V 부분: 각 단계 개별 실행. 연결은 접합 코드 필요]. 협력 행렬: `gpu.subgroup_mma_*` 가 `spirv.KHR.CooperativeMatrixMulAdd` 로 내려가는 변환이 있다는 포럼 글 [V 검색요약], `spirv.KHR.CooperativeMatrix*` op 은 설치본에 존재 [V]. 실제로 f16 협력 행렬 matmul 이 변환을 통과해 Adreno/Mali 에서 도는지는 **검증 못함**. + - MSL: **upstream 에 MSL 타깃 없음**. 우리가 확인한 설치본에서 dialect 목록(spirv, gpu, nvvm, amdgpu, emitc 등)에 metal/msl 이 없다 [V 목록 확인]. 외부 사례만 있음: Triton 외부 백엔드(triton-ext PR #126, bledden/triton-msl)가 TTGIR 에서 MSL 을 방출하는 패스를 가진다 [V 검색요약]. 가능한 길 둘: (a) `emitc` dialect 로 C++ 형 텍스트를 뽑는 구조를 흉내 내 **MSL 프린터를 우리가 작성**, (b) SPIR-V -> SPIRV-Cross -> MSL (협력 행렬은 simdgroup_matrix 8x8 float 로만 제한되고 최근 initial 지원 단계 [V 검색요약, SPIRV-Cross PR #2645]; Metal 4 cooperative tensor / matmul2d 는 도달 불가 [V 이전 조사]). 따라서 Apple 은 어느 쪽이든 MSL 방출기를 직접 쓴다 [I]. + - WGSL: upstream 에 WGSL 타깃 없음 [V 목록]. SPIR-V -> naga 또는 직접 방출. + - PTX: **되는 것을 확인** [V §1.1, `t_ptx.py`]. 가장 성숙한 경로. + - CPU SIMD: `vector` + `llvm` dialect -> LLVM IR -> ExecutionEngine/오브젝트 [V: `mlir.execution_engine` 임포트됨. 실제 NEON 오브젝트 방출은 시험 안 함]. 또는 `emitc` 로 C. iOS 앱엔 LLVM 을 싣지 않고 빌드 타임에 오브젝트/라이브러리만 [I]. +- 타일/레이아웃/파이프라인 변환: **여기가 MLIR 의 최대 이점.** `linalg`(structured ops) + `transform` dialect(`structured.tile_using_for`, vectorize 등), `vector.contract` 패턴, `gpu` dialect 의 mma/예약 [V 검색요약 + 설치본에 `_structured_transform_ops_gen`, `_gpu_transform_ops_gen`, `_nvgpu_transform_ops_gen` 존재 [V]]. 소프트웨어 파이프라이닝(cp.async 급) 은 nvgpu 쪽에 일부, Vulkan 쪽은 IREE 가 가진 별도 코드젠(IREE 의 SPIR-V 파이프라인은 upstream 이 아님 [V iree.dev 요약])이라 upstream 만으로 Vulkan 에서 Triton/IREE 급 matmul 성능이 나온다는 보장은 없다 [I]. +- 프런트엔드 접합: 파이썬 AST -> MLIR 를 `ir.Module`/`OpBuilder` 로 구성. "valid Python 그대로 실행" 요건은 프런트엔드가 AST 를 읽는 방식이므로 충돌하지 않는다 [I]. +- 도구 무게: wheel 66 ~ 76 MB / 설치 253 MB, 빌드 타임 전용이므로 앱 크기와 무관. 단 uv/pypackpack 설치 흐름에 PyPI 가 아닌 외부 index(find-links) 가 들어가고 날짜 고정이 필요하다 [V + I]. +- 위험: (1) nightly 스냅샷 의존과 API churn, (2) 3.15t wheel 부재, (3) eudsl 이 자체를 "alpha" 로 표기 [V], (4) MSL/WGSL 은 어차피 직접 방출, (5) 디버깅이 C++ 패스 안에서 일어나 에러 메시지가 불친절 [I], (6) wheel 이 단일 메인테이너 프로젝트(makslevental -> llvm/eudsl) 의 유지에 달림 [I]. +- 첫 슬라이스: 접합 코드(gpu.module 에 spirv.module + target_env/vce 부착) + 프런트엔드 -> MLIR 빌더 + 형상/버퍼 ABI 정리(`spirv.interface_var_abi`, push constant 는 MLIR 에서 `spirv.StorageClass PushConstant` 구성 필요 [I]). 추정 2 ~ 4 주 [I]; 환경/버전 고정과 f16/ABI 문제 때문에 옵션 1 과 비슷하거나 약간 더 걸릴 수 있다. 이득은 슬라이스 이후. +- 타일 matmul + cooperative matrix: 변환 파이프라인을 upstream 패스로 조립. Vulkan KHR 협력 행렬 경로는 존재하나 성능 튜닝 정도는 모름. 추정 1 ~ 3 개월, 하지만 위험이 옵션 1 보다 "변환이 안 맞을 때 고칠 수 없음" 쪽에 쌓인다 [I]. + +### 옵션 3. xdsl (순수 파이썬 MLIR 호환) + +- [V §1.2] 설치 쉬움(PyPI, 의존 3개, 3.10 ~ 3.14t), 활발(2~4 주 릴리스), 라이선스 Apache-2.0 with LLVM Exception. 3.15t 는 미시험이나 순수 파이썬이라 [I] 문제 가능성 낮음. +- MLIR 텍스트 호환: `xdsl-opt` 로 MLIR 텍스트를 읽고 쓴다(프린터/파서). 따라서 **xdsl <-> MLIR(mlir-opt, 위 wheel) 상호운용을 텍스트로 연결**할 수 있다 [I, 두 도구를 실제로 연결해 보지 않음]. 즉 옵션 3 은 옵션 2 와 배타적이지 않다: 프런트엔드/고수준 변환은 xdsl, SPIR-V 직렬화나 NVPTX 는 MLIR wheel 에 텍스트로 넘기는 혼합 구성이 가능 [I]. +- 대상별 경로: + - SPIR-V: **방출기 없음, spirv dialect 도 없음** [V]. 텍스트를 MLIR 에 넘겨야 하거나 SPIR-V 방출을 직접 써야 한다. + - MSL: 없음. `mps` 백엔드는 AIR 비트코드를 목표로 한 NotImplemented 스텁 [V]. + - WGSL: 있음(252줄) 단 범위는 작음(소스에서 `gpu.FuncOp`, `memref.StoreOp` 처리 등 기본 연산) [V 부분: 코드 헤드만 확인, 지원 범위 전체는 미확인]. + - PTX: 없음. LLVM IR 텍스트 변환(`xdsl.backend.llvm`)은 있어 `llvm` 의 NVPTX 로 이어 붙일 수 있으나 llvmlite 등 별도 필요 [I]. + - CPU SIMD: `vector`, `x86`, `arm_neon`, `riscv`, `llvm` 백엔드 등 있음. 실사용 성숙도는 미확인. +- 변환: `transform` dialect(1042줄) 와 `vector`(1470줄), `linalg` 가 있으나 upstream MLIR 의 `structured.*` 변환 패스 수준인지 확인 못 함. 협력 행렬/타일링 패스 존재는 확인 못 함 [V 한계]. 패스를 파이썬으로 쓰므로 새 패스 작성/수정은 쉬움(`pattern_rewriter`) [I]. +- 성숙도: 연구/컴파일러 교육 중심 프로젝트(스텐실, CSL, RISC-V/Snitch 등) 이며 GPU 코드젠 프로덕션 사례는 확인 못함 [I, 근거: dialect 구성]. +- 위험: 필요한 후단(SPIR-V/MSL/PTX 방출, 협력 행렬 dialect)이 거의 비어 있어, 옵션 1 에서 "IR 클래스 + 패스 인프라 + 프린터" 를 xdsl 로 빌려 오는 정도의 이득. IR 인프라(SSA, 파서, verifier, 패턴 리라이터, 정확성 검사)를 공짜로 얻는 점이 장점 [I]. +- 첫 슬라이스: 옵션 1 과 같은 SPIR-V 방출기를 어차피 써야 하므로 추가 이득은 IR 인프라 재사용뿐. xdsl 학습 비용 + 방출기 직접 작성 -> 2 ~ 4 주 [I]. Cooperative matrix: spirv dialect 를 xdsl 에 추가(IRDL 로 정의 가능 [I]) + 직렬화기 작성 필요. 옵션 1 대비 이득 작음. + +### 기타 (간단히) + +- Triton 의 MLIR dialect(TTIR/TTGIR) 재사용: 서브모듈/LLVM 포크에 묶여 있고 JIT 전제, Apple MSL 은 외부 포크(triton-ext)가 TTGIR -> MSL 을 했다 [V 검색요약]. 우리 파이썬-서브셋 프런트엔드와 TTGIR 의 레이아웃 인코딩을 쓰려면 Triton 빌드 전체가 따라온다 [I]. 비추천. +- IREE: SPIR-V 코드젠(Mali Valhall, Adreno 640+ 대상 triple) 과 Vulkan 런타임을 가진 AOT 컴파일러 [V iree.dev 요약]. 입력은 모델 단위(linalg/StableHLO)이고 산출물은 `.vmfb`(flatbuffer) 이며 SPIR-V 만 라이브러리로 뽑아 쓰는 공식 경로는 확인하지 못했다 [V 부재는 검색 한계]. 타일링/매트릭스 파이프라인이 가장 성숙한 오픈 구현이므로 **알고리즘 참조 대상**, 의존 대상은 아님 [I]. +- Mojo: 닫힘, MAX 가 Apple GPU 초기 [V 이전 조사]. +- SPIRV-Cross(MSL): 협력 행렬 초기 지원, simdgroup_matrix 는 8x8 float 계열만 [V 검색요약]. 즉 SPIR-V 경유 Apple 은 f16 matmul 에서 불리할 수 있고 Metal 4 tensor ops 는 불가 [V 이전 조사 + I]. + +## 3. 요약 표 + +| 항목 | 1. 자체 IR | 2. MLIR wheel | 3. xdsl | +|---|---|---|---| +| SPIR-V 방출 | 직접 작성 | **있음** (gpu/spirv dialect + `gpu-module-to-binary`, 접합 필요) [V] | 없음 [V] | +| MSL | 직접 (자연스러움) | upstream 없음, 직접 [V] | 없음, mps 는 AIR 스텁 [V] | +| WGSL | 직접 | upstream 없음 | 있음(소형) [V] | +| PTX | 직접(작게) | **있음, 실행 확인** [V] | 없음 (llvm 텍스트 경유 [I]) | +| CPU SIMD | C 벡터 확장 | vector+llvm / emitc [I] | arm_neon/x86/llvm 있음 [V 존재] | +| 타일/레이아웃/파이프라인 | 전부 직접 | linalg+transform+vector+gpu 가 큼 [V 존재, 성능은 미확인] | transform/vector 있음, 수준 미확인 | +| 설치 | 없음 | wheel 66~76MB, PyPI 아님, nightly find-links [V] | `uv add xdsl` [V] | +| 3.14t | 해당 | **wheel 있음, GIL 유지 off** [V] | **됨** [V] | +| 3.15t | 문제 없음 [I] | **wheel 없음** [V] | 가능성 높음 [I] | +| 버전 churn | 없음 | 높음 (매일) [V] | 중 (2~4 주 릴리스, 0.x) [V] | +| 첫 슬라이스 | 2~4 주 [I] | 2~4 주 [I] | 2~4 주 [I] | +| 타일 matmul + coop mat | 2~4 개월 [I] | 1~3 개월 (고장 시 수정 불가 위험) [I] | 2~4 개월 [I] | + +## 4. 권고 (불확실성 명시) + +1. **IR 의 정본은 우리가 소유한다(옵션 1 의 IR), 단 MLIR 호환 형태로 설계한다.** 이유: (a) §2.1 요건과 프런트엔드는 우리 몫이며 어떤 외부 IR 도 대신하지 못한다 [I], (b) MSL, WGSL 은 MLIR 을 써도 직접 써야 하므로 MLIR 이 줄여 주는 방출기는 SPIR-V 와 PTX 둘뿐이다 [V], (c) 3.15t wheel 이 없고 nightly 고정이라 도구 체인의 기초로 삼기에는 변동이 크다 [V]. 이 판단은 "방출기가 일의 대부분" 이라는 가정에 의존한다 [I]. +2. **MLIR wheel 은 "SPIR-V 와 PTX 직렬화기/검증기, 그리고 비교 오라클" 로 선택적 사용을 우선 시험한다.** 우리 IR 을 MLIR 텍스트로 내리는 얇은 방출기(수백 줄)만 쓰면, `convert-gpu-to-spirv` + `gpu-module-to-binary` 로 SPIR-V 를, `gpu-module-to-binary{format=isa}` 로 PTX 를 얻는다 [V 각 단계 개별]. 막히면 직접 SPIR-V 방출로 후퇴할 수 있어 위험이 낮다. 3.15t 와는 별도 프로세스(3.14t 또는 GIL 인터프리터) 로 분리 실행 [I]. +3. **타일 변환(tiling, 협력 행렬, pipelining) 에서 MLIR 의 transform/linalg/vector 를 쓸지는 슬라이스 2 이후 별도 시험으로 결정.** 지금은 약속하지 않는다. 이 영역이 MLIR 의 진짜 이점이지만 Vulkan 에서의 성능은 IREE 쪽 코드젠에 있다 [I, V 검색요약]. +4. **xdsl 은 채택하지 않되 IR/패스 인프라 설계의 참고, 그리고 MLIR 텍스트 상호운용 후보로 남긴다.** 후단이 비어 있어 SPIR-V/MSL/PTX 에 이득이 없다 [V]. +5. **Apple 은 MSL 직접 방출.** 이전 조사와 동일 결론이고 이번에도 어느 옵션도 대신하지 않는다 [V + I]. + +미결/불확실: (a) `gpu.func` 하향 결과를 `gpu-module-to-binary` 로 연결하는 접합이 실제로 쉬운지 (실행 안 함), (b) f16/협력 행렬 SPIR-V 가 실기(Adreno, Mali)에서 도는지, (c) wheel 이 NEON/x86 CPU 오브젝트를 만들 수 있는지(LLVM 타깃 구성 미확인; NVPTX 는 확인), (d) 3.15t 에서의 MLIR 소스 빌드 가능성, (e) eudsl wheel 의 장기 유지, (f) 시간 추정은 모두 측정 아닌 추정. + +## 5. 다음 시험 (구체적 첫 슬라이스) + +1. `t_spirv.py` 를 확장: 하향된 `spirv.module` 을 `gpu.module [target_env]` 안에 옮기고 `requires` 를 붙여 직렬화, `spirv-val` 로 검증 (`spirv-val` 은 이 환경에 설치하지 않았음). +2. 그 `.spv` 를 `torch_c/shaders/` 의 체크인 모양대로 넣고 `dispatch_kernel_words` 경로로 실행, CPU 참조(torchnative flash.rs 식 허용 오차 비교)와 대조. 이 단계는 torchnative 쪽 빌드가 필요하므로 이번 조사에서는 하지 않았다. +3. 같은 커널을 직접 방출한 SPIR-V 와 크기/명령어/기기 호환성 비교(옵션 1 대비 MLIR 산출물의 위험 평가). + +## 6. 이 조사 방법이 찾지 못하는 것 + +wheel 설치와 개별 패스 실행으로 "존재와 단일 단계의 동작" 만 확인했으므로, 연결된 전체 파이프라인의 정확성, 실기 GPU 드라이버 동작, SPIR-V 산출물의 성능, 장기 API 안정성, 그리고 xdsl/MLIR 변환 패스의 실제 품질은 찾지 못한다 (시간 추정도 측정이 아니다).