TileLang 데일리 리포트 (2026-10-13)
모니터링 기간: 2026-10-12 07:00 ~ 2026-10-13 07:00 (베이징 시간, 24시간).
출처: GitHub(tile-ai 조직 29개 저장소 푸시 전수 점검, 기간 내 9개 저장소에 푸시 발생: tilelang, TileOPs, TileFoundry, tilelang-ascend, tilelang-hygon, tilelang-sunrise, TileOPs-nightly, tilelang.github.io, TileOPs.github.io; 조직 기간 내 병합 20건, 신규 오픈 41건, 심사 중 업데이트 59건; 메인 저장소 7건 병합, TileOPs 10건 병합, TileFoundry #221 랜딩, Ascend NSA 이중 마이그레이션 병합, Hygon 3개 브랜치 추진, 야간 파이프라인 신규 스냅샷 1건, 메인 문서 사이트 3회 빌드), 커뮤니티 저장소 점검(PTO-ISA SuperNPUBench 커밋 및 릴리스, tilelang-debugger, SwHwCodesign 계정), Google News RSS 중영문 다중 쿼리(프록시 경유), Hacker News, arXiv, 문서 사이트 및 릴리스 페이지 점검.
이번 호 인덱스
- 오늘의 중점: ROCm 두 건과 Ascend NSA 대규모 마이그레이션 확정——TileOPs 하루 10건으로 SM89를 두 번째 플랫폼으로 구축 (10-12 ~ 10-13)
-
- 핵심 프로젝트 진행
- 1.1 메인 저장소: ROCm 두 건——DeepSeek NSA 전방/디코딩과 flash-decoding 가능 타일 (10-13)
- 1.2 메인 저장소: thread_storage_sync 정확한 스레드 경계 (#3291) 및 #2954 수정 (10-13)
- 1.3 메인 저장소: Autotune 구조화 실험 보고서, CodeRabbit 자동 할당 및 테스트 규범 문서 (10-12 ~ 10-13)
- 1.4 메인 저장소: Ascend PTODSL SIMT 리덕션 헬퍼 vendoring (#3469)으로 CANN Weekly 컴파일 수정 (10-12)
- 1.5 메인 저장소 심사 중: SM120 MXFP8/MXFP4 블록 스케일링 MMA, L2 프리페치 및 Host IR 두 건 등 (10-12)
- 핵심 프로젝트 진행
-
- 멀티 백엔드 적응 (Ascend / MetaX / Hygon / Moore Threads)
- 2.1 Ascend: NSA Decode / NSA Bwd 두 건 대규모 마이그레이션 병합——48건 계층 테스트 (10-12)
- 2.2 Ascend: cross-warp allreduce 그룹화 수정 오픈 (#3494)과 복사/정밀도/캐시 수정 배치 (10-12)
- 2.3 Hygon: FP8 브랜치 추진 (432 TFLOPS 기준), 업스트림 v0.1.15 업그레이드 브랜치와 등록 파이프라인 (10-12)
- 2.4 MLIR Ascend 두 건 심사 중; Sunrise는 CI 의존성 업그레이드만; MetaX / Moore Threads 무소식 (10-12)
- 멀티 백엔드 적응 (Ascend / MetaX / Hygon / Moore Threads)
-
- 생태계와 채택자
- 3.1 TileOPs: SM89 플랫폼 확장 배치——GQA 프리필/디코딩, GDN / DeltaNet 과 Norm (10-12)
- 3.2 TileOPs: foundry 성능 배치——NSA 영구 CTA, W4A16 GEMM 과 16비트 나눗셈 (10-12)
- 3.3 TileOPs: 선형 어텐션 수정——TMA 서스펜드 회피와 GLA 프리필 상태 스킵 (10-12)
- 3.4 TileOPs 심사 중: SM89 네 건 연속, 페이지드 KV / 어텐션 클러스터와 MoE 두 건 (10-12)
- 3.5 TileFoundry: 파서 연산 속성과 어휘 레이아웃 슈가 병합 (#221); scaffold 독립 프로젝트 심사 중 (#227) (10-12)
- 3.6 야간 파이프라인: 신규 스냅샷——정확성 1042건 (+10), 벤치마크 1947건 (52개 스위트) 무실패 (10-13)
- 3.7 채택자와 문서: TileKernels / FlashQLA 무소식; 메인 문서 사이트 3회 빌드 (10-12 ~ 10-13)
- 생태계와 채택자
-
- 커뮤니티, 튜토리얼 및 활동
- 4.1 PTO-ISA: SuperNPUBench 네 건 커밋과 두 개 릴리스 (ops-20261012 / ops-20261013) (10-12)
- 4.2 커뮤니티 계정: SwHwCodesign 여섯 개 AI-* 플레이스홀더 저장소; TileLang-Awesome 접근 불가 (10-12)
- 4.3 미디어와 학술: Google News / HN / arXiv 기간 내 신규 없음 (10-12)
- 커뮤니티, 튜토리얼 및 활동
-
- 트렌드 관찰
- 5.1 ROCm 보충: 디바이스 예산 인식 두 건 수정
- 5.2 SM89이 두 번째 플랫폼으로: 하루 6건 병합에 4건 심사 중
- 5.3 Ascend 마이그레이션 라인은 정확성 마무리 단계로 전환
- 5.4 릴리스 침묵기의 기반 성장
- 트렌드 관찰
- 부록: 자료와 점검 설명
- 전체 출처 목록
오늘의 중점: ROCm 두 건과 Ascend NSA 대규모 마이그레이션 확정——TileOPs 하루 10건으로 SM89를 두 번째 플랫폼으로 구축 (10-12 ~ 10-13)
날짜: 2026-10-12, 2026-10-13 출처: ROCm NSA(#3152)/ ROCm flash-decoding 타일(#3149)/ Ascend NSA Decode(#1701)/ Ascend NSA Bwd(#1730)/ TileOPs PR 목록/ 야간 스냅샷(6ec0d086)
모니터링 기간 24시간 이내에 조직은 20건을 병합했고(메인 저장소 7, TileOPs 10, TileFoundry 1, Ascend 적응 저장소 2), 41건을 새로 열었으며, 9개 저장소에 푸시가 있었고, 병합은 하루 종일 이어져 10-13 새벽에 집중적으로 마무리되었다(메인 저장소 4건의 병합이 베이징 시간 새벽에 이루어짐).
- 메인 저장소 ROCm 두 건 확정: DeepSeek NSA 순전파/디코드가 HIP 측에서 64 lane 단일 웨이브로 전환(#3152); flash-decoding 커널이 디바이스 LDS/VGPR 예산에 따라 128x64와 64x64 타일 사이에서 자동 선택(#3149) — 각각 gfx942에서의 「서브웨이브 block 거부」와 「타일 초과 예산으로 인한 시작 실패」를 수정(상세는 1.1 참조).
- Ascend NSA 두 건 이관 병합: Decode(#1701)와 역방향(#1730) 두 건은 합계 +2348행의 예시로, 모두 Developer mode hybrid 구현이며 28 / 20건의 계층적 테스트 스위트를 첨부(상세는 2.1 참조).
- TileOPs 하루 10건 병합: SM89는 하루 만에 GQA 프리필/디코드, GDN / DeltaNet 프리필, InstanceNorm과 Norm 재구성을 받아들였고; foundry 성능 배치(NSA 영속 CTA, W4A16, 16비트 나눗셈)는 같은 날 진행(상세는 3장 참조).
- Hygon 세 브랜치 진행: 신규 기능 브랜치 FP8 MQA(커밋에서 channelwise GEMM이 432 TFLOPS에 도달했다고 주장), 업스트림 v0.1.15 업그레이드 브랜치 적응 병합, 등록 파이프라인 #18은 계속 심사 중(상세는 2.3 참조).
- 야간 스냅샷과 문서 사이트: 정확성 1042건(이전 대비 +10), 벤치마크 1947건(52 스위트) 무실패; 메인 문서 사이트 3회 빌드.
1. 핵심 프로젝트 진행
1.1 메인 저장소: ROCm 두 건 — DeepSeek NSA 순전파/디코드와 flash-decoding 실행 가능 타일(10-13)
날짜: 2026-10-13(베이징 시간 새벽에 병합) 출처: ROCm NSA(#3152)/ ROCm flash-decoding 타일(#3149)
- DeepSeek NSA on ROCm(#3152, +111/-2): gfx942의 웨이브프론트 크기는 64이며, TileLang은 웨이브프론트보다 작은 HIP block을 거부하는데, 두 NSA 예시는 32 스레드를 하드코딩하고 있었다; 이번 건은 순전파와 단일 token 디코드 커널에 대해 단일 64 lane 웨이브를 선택하고, CUDA의 기존 32 스레드 시작 구성을 유지하며, gfx942에서 인과 순전파와 단일 token 디코드의 정확성 커버리지를 보강한다(PyTorch NSA 레퍼런스 대조, FP16, 16 Q 헤드 1 KV 헤드, 시퀀스 64).
- flash-decoding 타일 자동 적응(#3149, +82/-5): gfx942에서 128x64 타일의 생성 커널은 66048바이트 동적 LDS가 필요하며(65536 디바이스 상한 초과) 256 VGPR 한계에 도달해, 시작 시 「invalid argument」를 보고했다; 이번 건은 시작 전에 생성된 HIP LDS와 VGPR 점유를 검사하고, 예산을 초과할 때만 기존 64x64 타일로 폴백하며(디바이스 이름으로 하드코딩하지 않음), CUDA 경로는 변경하지 않고, gfx942 런타임 회귀를 추가한다.
1.2 메인 저장소: thread_storage_sync 정확한 스레드 경계(#3291) 및 #2954 수정(10-13)
날짜: 2026-10-13 출처: 정확한 스레드 경계(#3291)
thread_storage_sync에 대한 정확한 스레드 경계를 계산합니다(4개 파일, +292/-96). 작성자는 #2954가 이 커밋으로 수정되었음을 확인했으며, “동등한 부분 스레드 가드가 동일한 barrier로 환원됨”을 검증하는 테스트 케이스를 추가하고 SM89(4060 Ti)에서의 전후 컴파일 산출물 비교를 첨부했습니다.
1.3 메인 저장소: Autotune 구조화 실험 보고서, CodeRabbit 자동 할당 및 테스트 규범 문서 (10-12 ~ 10-13)
날짜: 2026-10-12, 2026-10-13 출처: 실험 보고서(#3393)/ CodeRabbit 자동 할당(#3487)/ 테스트 규범(#3489)
- Autotune 실험 보고서(#3393, +98/-7): 실험적
AutoTuner.run(on_trial=...)관찰자 인터페이스를 추가했습니다(RFC #3375에 대응). 시험마다 설정 인덱스, JSON 설정, 결과, 지연 시간, 검증 레이블, 제한된 오류 텍스트를 출력하며, 콜백은 집계 스레드와 공유 캐시 잠금 밖에서 실행됩니다. 보고/직렬화 오류는 해당 보고만 종료하고 튜닝에는 영향을 주지 않으며, 후보 생성·타임아웃·캐시 아이덴티티를 변경하지 않습니다. 작성자는 “콜백 vs 고정 JSONL 아티팩트”가 여전히 열린 API 결정 사항임을 명시했습니다. - 리뷰 가드레일(#3487, +1/-1): CodeRabbit의
reviews.auto_assign_reviewers를 true로 설정하여 실제로 review 요청을 발송하도록 했습니다(이전에는 순회 검토에서 권장 사항만 나열). 11개의 reviewer 규칙과 순서는 변경되지 않았습니다. - 테스트 규범(#3489, 8개 파일): Testing Style Guide를 신설했습니다 — 기존 커버리지에 의존할 시점, 최소 확장할 시점, 집중 커버리지를 신설할 시점.
tilelang-cpp-style은tilelang-dev-guide로 이름을 변경하고 C++ 규칙은 필요 시 참고용으로 전환했으며, CONTRIBUTING은 간소화하고 링크를 연결했습니다.
1.4 메인 저장소: Ascend PTODSL SIMT 리덕션 헬퍼 vendoring(#3469)으로 CANN Weekly 컴파일 수정 (10-12)
날짜: 2026-10-12 출처: PTODSL 리덕션 헬퍼 vendor(#3469)
- PTOAS 상류에서
ptodsl/_allreduce.py를 제거하고 TileLang이 자체적으로 vendor하기를 기대했지만, 메인 저장소의contrib/ptodsl/simt.py가 여전히 해당 모듈을 참조하고 있었습니다 — 최신 CANN Weekly(09-30)로 모든 PTO 대상 연산자를 컴파일하면 즉시ModuleNotFoundError가 발생했습니다. 이 커밋은 헬퍼를tilelang/contrib/ptodsl/allreduce.py로 vendored하고(+555/-115), 정수 부호 호환성을 monkey patch에서 리덕션 테이블 및 dtype 디스패치로 통합하여 PTO 대상의 사용 가능성을 복원했습니다.
1.5 메인 저장소 심사 중: SM120 MXFP8/MXFP4 블록 스케일링 MMA, L2 프리페치 및 Host IR 두 건 등 (10-12)
날짜: 2026-10-12 출처: SM120 MX 블록 스케일링 MMA(#3495)/ L2 프리페치(#3490)/ Host IR 직접 연결 JIT(#3481)/ Host IR 실행 통합(#3473)/ Windows ARM64 CI(#2962)
- #3495 (+884/-51):
T.gemm_blockscaled가 SM120/SM121에서 블록 스케일링 MMA 제품군을 NVFP4에서 MXFP8(kind::mxf8f6f4, 1X, UE8M0)과 MXFP4(mxf4nvf4, 2X, UE8M0)로 확장하며, FP32 누적, K 요소 32개마다 스케일 하나; 10-10에 병합된 SM121 점등(#3430)을 이어 하드웨어 매트릭스를 계속 확장한다. - #3490 (+87/-1):
tl.prefetch_l2(addr, eviction_policy=...)를 신규 추가, evict_normal / evict_last를 지원하며 sm_80 이상을 요구하고, CUDA codegen이 정책을 컴파일 타임 템플릿 파라미터로 방출한다. - Host IR 두 건(#3481 / #3473): 직접 연결 Host IR JIT와 캐시, CUDA와 CuTeDSL 실행이 Host IR을 통해 통합되며 계속 심사 중.
- 기타 심사 중: CI에 Windows ARM64 추가(#2962), JIT 파라미터 해시 가능 검증(#3417), 부분 웨이브 스칼라 리덕션(#3478), 양수 상수 제수 안전 guard 평탄화(#3488), cast 조합 검증(#3350), TMA 디스크립터 이름 검증(#3484) 등.
2. 멀티 백엔드 적응 (Ascend / MetaX / Hygon / Moore Threads)
2.1 Ascend: NSA Decode / NSA Bwd 두 건의 대규모 이관 병합——48개 계층별 테스트 (10-12)
날짜: 2026-10-12 출처: NSA Decode(#1701) / NSA Bwd(#1730)
- #1701 (+755):
example_tilelang_nsa_decode.py가 Developer mode hybrid(combineCV + auto_cv_sync + auto_sync + memory_planning 네 가지 pass_configs)로 NSA Decode 단일 코어 디코딩(seq_q=1)을 구현, 명시적 L1/L0C/UB 할당, 5개 세그먼트의 GM workspace를workspace_idx=[6..10]로 자동 할당; 컴파일러가 Cube/Vector 스코프를 자동 분할하고 크로스 코어 flag 동기화를 삽입(수동 barrier/Scope/set_flag 불필요). 계층별 테스트 28개: L0 임계값, L1 기능(DeepSeek-V3 전형적 구성 D=128, S=16, BS=64 포함), L2 예외 및 경계 특수값; 910B3, fp16 실측(B=2, SEQ_LEN=64). - #1730 (+1593):
example_tilelang_nsa_bwd.py가nsa_fwd(온라인 softmax + 희소 K 블록 선택, 엔드투엔드용)와nsa_bwd_single(단일 코어 역방향: 5 GEMM + softmax + Delta + dS 온칩 융합) 두 개의 JIT 커널을 제공; 역방향은 flashattn_bwd 패러다임에 따라 온칩 직접 연결 T.copy(L0C→UB, UB→L1)로 9개의 GM workspace를 생략하고, 2단계 보상 GEMM(ds_delta / ds_delta2)을 유지; 계층별 테스트 20개. 두 건의 이관 배치는 국경절 전후에 걸쳐 진행되었고, 이번 모니터링 기간에 적응 저장소의 기본 브랜치에 확정되었다.
2.2 Ascend: cross-warp allreduce 그룹화 수정 개시(#3494)와 복사/정밀도/캐시 수정 배치 (10-12)
날짜: 2026-10-12 출처: cross-warp allreduce 그룹화(#3494) / 예제 정밀도 정렬(#1867) / situ_glu_grad 예제(#1877) / copy_ub_to_gm 수정(#1795)
- #3494(메인 저장소, +145/-26, 리뷰 중): 하나의 cross-warp allreduce가 여러 독립 리덕션 그룹을 커버할 때 ascend와 pto 두 타깃이 첫 번째 그룹의 결과를 모든 그룹에 전송하는 문제——#3469 머지 이후 커뮤니티가 보고한 정확성 피드백(최소 재현 포함)이며, 수정은 그룹별 전송 방식이다. Ascend 라인의 “한 배치 마이그레이션, 한 배치 감사” 리듬이 뚜렷하다.
- 적응 저장소 승인 대기 중: 예시 정밀도 검사 정렬 표준(#1867), situ_glu_grad 연산자 예시(#1877),
copy_ub_to_gm이 컴팩트 sub-32B 행을 표현하지 못해 keepdim 리덕션 출력이 손상되는 문제 수정(#1795); 캐시 3분할(#1871~#1873, JIT 특화 키에 튜닝 파라미터 포함, autotune 키를 유효 프로파일 파라미터에 정렬, NPU와 시뮬레이터 캐시 분리). - 메인 저장소 Ascend 리뷰 중 클러스터: SimdVF T.fill 다중 버퍼(#3455), SimtVF remat_idx / 벡터 gather pack(#3492), SimdVF dim=1 리덕션 행별 단일 원소 저장(#3437), L0C fp32→f16/bf16 dual_copy 인플라이트 cast(#3468과 #3328 두 건 병행), 독립 AIV 실행(#3485), T.PersistentKernel과 AutoPersistent(#3325), 수동 FlashMLA 예시(#3464), VF 스코프 외 스칼라 sqrt(#3466), DeepGEMM의 PTO codegen(#3330), 예시 타깃 선택(#3474) 등.
2.3 Hygon: FP8 브랜치 추진(432 TFLOPS 기준), 업스트림 v0.1.15 업그레이드 브랜치와 등록 파이프라인(10-12)
날짜: 2026-10-12 출처: 레지스터 파이프라인(#18)/ feat/hcu-mqa-fp8 브랜치/ upgrade/upstream-v0.1.15 브랜치
- 레지스터 파이프라인(#18): 등록 브랜치
gcy_merge_main(49개 파일, +4144/-318)이 모니터링 기간 내 두 번 푸시되어 업데이트되었으며 계속 리뷰 중——Hygon 백엔드에 레지스터 파이프라인(공유 메모리 파이프라인과 협조) 및 4웨이브 초과 시 warp-diverge를 추가하며, 작성자는 예시 GEMM이 대응 라이브러리 수준의 성능에 도달했다고 밝혔다. - FP8 기능 브랜치
feat/hcu-mqa-fp8: 커밋 체인이 MLS channelwise FP8 GEMM을 커버——10-10부터 “432 TFLOPS 달성” 커밋, 10-12에 “M≤32일 때 fp8 디코딩 GEMM 대역폭이 hipblas를 초과” 추가 및 파일과 네이밍 정리(모두 브랜치 내 작성자 자체 측정 기준이며, 아직 PR 연관은 확인되지 않음). - 업그레이드 브랜치
upgrade/upstream-v0.1.15: 10-12에 “Merge upstream v0.1.15 and adapt HCU backend” 머지, Hygon 백엔드가 메인 저장소 v0.1.15와 정렬됨. - MetaX(최근 푸시 09-24), Moore Threads(09-30)는 이번 모니터링 기간 내 활동 없음.
2.4 MLIR Ascend 두 건 리뷰 중; Sunrise는 CI 의존성 업그레이드만; 나머지 백엔드는 조용(10-12)
날짜: 2026-10-12 출처: T.Pipelined 파라미터화(#199)/ FloorDiv / FloorMod codegen(#200)/ tilelang-sunrise
- MLIR Ascend:
T.Pipelined()가 num_stages / stage / order / group 파라미터를 지원하는 건 리뷰 중(#199); TileLangIR의 FloorDiv / FloorMod codegen 리뷰 중(#200). - Sunrise(Sunrise S2/S3의 TANG 백엔드 저장소): 모니터링 기간 내 Dependabot의 CI 의존성 업그레이드 두 건만 리뷰 중(upload-artifact, download-artifact)이며 코드 측면 동향 없음.
- 조직 수준의 나머지 저장소(TileRT, tilescale, DeepStack, tilelang-puzzles 등)는 모니터링 기간 내 모두 커밋 없음.
3. 생태계와 채택자
3.1 TileOPs: SM89 플랫폼 확장 배치——GQA 프리필/디코딩, GDN / DeltaNet과 Norm(10-12)
날짜: 2026-10-12 출처: D256 GQA 디코드(#2529)/ dense GQA 프리필(#2528)/ GDN / DeltaNet 프리필(#2527)/ InstanceNorm 수정(#2526)/ Norm 구성 리팩터링(#2544)
- D256 dense GQA 디코드(#2529): SM89에서 256차원 dense GQA 디코드를 개방——64×64, 단일 stage, 128 스레드 분할 프로그램을 그대로 사용하며, 16개 후보 분할이 98304바이트를 차지한다(SM89 상한 101376 미만). 단일 split 후보는 제외하고, 짧은 KV의 두 split 후보는 유지한다. 테스트는 61개 노드(+6)로 확대되어 KV 길이 1 / 63 / 2048을 커버한다. RTX 6000 Ada 실측.
- dense GQA 프리필(#2528, +260/-43): SM89에 FP16/BF16과 네이티브 FP8 프리필을 추가(패킹 어텐션 프로그램 재사용, 헤드 차원 128). 14개 사례가 main에서 실패하고 본 커밋에서 통과(10 FP16/BF16 + 4 FP8). 사각형 마스크, 그룹 헤드, 사용자 정의 scale, softcap, RoPE와 FP8 역양자화를 유지하며, 기본 구성과 메모리 회계를 공유하여 중복 구현을 피한다.
- GDN / DeltaNet 프리필(#2527, 11개 파일): KDA의 chunk prepare와 scan 프로그램을 재사용하여 SM89 프리필을 구현. 21개 범용 프리필 사례가 ‘실패하여 배포 불가’에서 통과로 전환. raw-beta 반올림, 패킹 시퀀스 경계, 그룹 헤드, FP32 초기/최종 상태와 GDN의 value-major 상태 레이아웃을 유지하며, 안정성 삼각 솔버를 재사용하고 상관 키 사례의 허용 오차는 완화하지 않는다.
- Norm 두 건: #2526은 192×192 공간 입력의 InstanceNorm 학습 컴파일을 수정(여섯 batch / 정밀도 조합이 실패에서 통과로, GroupNorm의 너비 경계와 두 번의 리덕션을 재사용). #2544는 공유 시프트 리덕션 너비 상수(32768)를 norm 구성 모듈로 되돌린다(87개 사례 전부 통과, 성능 주장 없음).
3.2 TileOPs: foundry 성능 배치——NSA 영구 CTA, W4A16 GEMM과 16비트 나눗셈(10-12)
날짜: 2026-10-12 출처: NSA 영구 CTA(#2531)/ W4A16 GEMM(#2534)/ 16비트 나눗셈(#2543)
- NSA 압축 영구 CTA(#2531, +391/-2): SM90에서
NSACompressedVarlenFwdOp가 새 커널NSACompressedFwdVarlenTMAKernel로 교체——기존 커널은 (token, KV 헤드)당 32 스레드 CTA 하나로 토큰별로 압축 K/V를 다시 로드했으나, 새 커널은 512 영구 CTA로 각 CTA가 64행 타일(4 token × 16 헤드)을 순회하며, 쿼리는 이중 stage TMA 링을 통해, 요청 K/V 블록은 공유 메모리에 상주하고, QK는 n=16 SS WGMMA로, PV는 n=64 RS WGMMA로 처리하며, O는 QK 읽기 완료 후 TMA를 통해 log-sum-exp와 함께 출력된다. SM당 4개의 48KB CTA가 상주한다. - W4A16 GEMM(#2534, +24/-0):
_legal_configs가 K 타일이 K를 초과하지 않도록, 링 깊이가 K의 타일 수를 초과하지 않도록 제한하며, K<512일 때 더 이상 0으로 채우지 않는다. single-tile(K=128)과 multi-tile(K=256) 행은 8x64 타일에 128/256 너비 K 타일(두 stage)을 사용하도록 변경되어 torch 역양자화 matmul을 앞선다. 나머지 13개 행의 구성은 그대로이며, 3개 행은 여전히 뒤처진다(작성자가 스스로 한계를 명시). - 16비트 나눗셈(#2543, 9개 파일): Reciprocal과 bf16 sigmoid를 IEEE float32 나눗셈에서
__fdividef고속 나눗셈으로 교체(bf16은 범위 스케일링 버전 사용). 16M 행 fp16/bf16 지연 시간 17.86 / 17.79 마이크로초로, 구버전 대비 1.02x, torch-compile 대비 0.991x / 0.995x.
3.3 TileOPs: 선형 어텐션 수정——TMA 행 회피와 GLA 프리필 상태 건너뛰기(10-12)
날짜: 2026-10-12 출처: 단일 타일 디코드 상태 로드(#2538)/ GLA 프리필 상태 건너뛰기(#2540)
- #2538(+66/-6): SM90에서
k_tile == dim_k인 GLA / DeltaNet 재귀 디코드 프로그램이 첫 실행 시 바로 중단됨——TileLang 0.1.12가 단일 단계 상태 복사를 1D 배치 복사로 낮추면서, 그 producer가 warp specialization 하에서 배리어에 영원히 도달하지 못함. 네 개의 builder를 비-TMA 로드로 변경한 후, dim_k를 나누어떨어뜨리는 임의의 k_tile이 모두 실행 가능해졌고,k_tile < dim_k동작은 변하지 않음. 코드에 FIXME를 남겨 두고, TileOPs가 TileLang 0.1.13+를 고정한 뒤 우회를 회수할 예정. - #2540(+77/-67): 네 세트의 GLAFwdOp 프리필 구현을
has_initial_state에 따라 특화——초기 상태가 없을 때 로컬 재귀 타일을 비우고 읽지 않은 스칼라로 자리만 차지하게 하여, 전체 FP32 초기 상태 텐서의 할당, 제로 초기화, 읽기를 생략함. 상태가 있는 빌드의 디바이스 소스 코드는 main과 바이트 단위로 동일하고, 공용 인터페이스와 디스패치 도메인은 변하지 않음(#2380 종료).
3.4 TileOPs 심사 중: SM89 네 연속, 페이지드 KV / 어텐션 클러스터와 MoE 두 건(10-12)
날짜: 2026-10-12 출처: SM89 네 연속(#2546 ~ #2549)/ varlen sinks와 슬라이딩 윈도우(#2537)/ PagedKVCacheWrite(#2539)/ masked FP8 그룹 GEMM(#2521)/ DeepSeek 그룹 top-k(#2535)
- SM89 네 연속(#2546 ~ #2549): GLA 역방향 상태 수명, DeltaNet 역방향 공유 메모리 적응, KDA 융합 프리필, MoE 컴팩트 행별 그룹 GEMM——SM89 커버리지가 머지된 건에서 심사 중인 건으로 확장됨.
- 페이지드 KV / 어텐션 클러스터(여섯 건): varlen GQA 헤드 sinks와 파이프라인 슬라이딩 윈도우(#2537——윈도우 행이 기존 최적 베이스라인보다 1.03~1.13배 빠르고, sink 혼합이 FlashInfer보다 1.51~1.53배 빠르며, 부분 RoPE와 2행 FP8은 여전히 참조보다 낮음); PagedKVCacheWriteFwdOp(#2539, 호출자 슬롯별 제자리 스캐터, FP8 독립 K/V 스케일과 포화 반올림, vLLM 벤치마크 어댑터 포함); paged GQA에 RoPE/윈도우/sinks 추가(#2516), 페이지드 KV gather(#2519), LSE 가중 어텐션 상태 병합(#2530), 긴 인과 KV의 dense FP8 GQA 파이프라인(#2541).
- MoE 두 건: masked FP8 그룹 GEMM(#2521,
deep_gemmmasked 시맨틱에 정렬, MoE roofline 공식이 전체 텐서 이동으로 인한 트래픽 과대평가를 수정, 14행 부하); DeepSeek 그룹 top-k 라우팅(#2535,num_expert_group/topk_group/routed_scaling_factor를 vLLMgrouped_topk에 정렬, T=512일 때 32 / 132 SM만 차지하는 스케줄링 문제도 수정). - GEMM 성능 배치(네 건): FP8 타일 재배열과 1x128 B 스케일 웨이브 커널(#2533), 짧은 컴팩트 그룹으로 SM 채우기(#2536), 모듈 상수 스코프 수렴(#2542), W4A16 스코어링과 K 슬라이싱 결합(#2545)이 계속 심사 중.
3.5 TileFoundry: 파서 연산 속성과 어휘 레이아웃 슈가 머지(#221); scaffold 독립 프로젝트 심사 중(#227)(10-12)
날짜: 2026-10-12 출처: 파서 속성과 레이아웃 슈가(#221)/ scaffold 독립 프로젝트(#227)
-
#221 병합(222개 파일, +4473/-4951): op 속성과 레이아웃에 대한 통합 작성자 인터페이스 구축——schema화된 op 속성, 레이아웃 조합 및 명시적 값 상태를 파싱하고, fixture와 튜토리얼을 마이그레이션하고 printer 골든 샘플을 재생성하며, mesh 선택을 조합 생성으로 변경하고
specialize(mesh=...)를 지원; parser 문법/규칙을 실행 가능한 패턴에서 생성하고 pre-commit 검증을 첨부; buffer-alias 변경과 기반을 합친 후 Reshard 별칭과 영역 메모리 기록을 정렬. -
#227 심사 중(47개 파일):
tilefoundry scaffold SOURCE OUT_DIR신규 추가——Target codegen에 따라 편집 가능한 독립 CMake 프로젝트(파라미터 뷰, 특수화 디스패치, 런처, 호스트 등록 및 자리 표시자 CUDA 바디)를 생성하여 작성자가 호출 규약을 깨지 않고 커널 바디를 인계받을 수 있도록 함; codegen 산출물과 링크를 6개 모듈 파일로 분리하고,LinkedModule.functions가 단일entry를 대체.
3.6 야간 파이프라인: 신규 스냅샷——정확성 1042건(+10), 벤치마크 1947건(52개 스위트) 제로 실패(10-13)
날짜: 2026-10-13 출처: 야간 스냅샷(6ec0d086)/ 스냅샷 실행 메타데이터
- 신규 스냅샷 6ec0d086(run 38288018755, 10-13 03:53): 정확성 1042건(61개 스위트, 이전 스냅샷 대비 +10, 0 실패 0 오류 1 스킵), 벤치마크 1947건(52개 스위트, 0 실패); 환경 H200, CUDA 13.2, 드라이버 595.71.05, 스냅샷 내 tileops 0.0.1.dev1296+g6ec0d086d, tilelang 0.1.12.
- 정확성 증분은 모니터링 기간 내 SM89 시리즈 병합과 같은 박자로, 평가 기반은 계속 “한 건 병합, 기반 한 단계 상승”의 롤링 리듬 유지.
3.7 채택자와 문서: TileKernels / FlashQLA 조용; 메인 문서 사이트 3회 빌드(10-12 ~ 10-13)
날짜: 2026-10-12, 2026-10-13 출처: 메인 문서 사이트/ TileKernels/ FlashQLA
- 채택자: TileKernels(DeepSeek)와 FlashQLA(Qwen)는 모니터링 기간 내 푸시 없음(모두 09-30 이후 조용); TileRT는 모니터링 기간 내 동향 없음.
- 문서 사이트: 메인 사이트 3회 빌드(10-12 2회, 10-13 새벽 1회); TileOPs 사이트 1회 gh-pages 배포(10-12 오전).
4. 커뮤니티, 튜토리얼 및 활동
4.1 PTO-ISA: SuperNPUBench 4건 커밋과 2건 릴리스(ops-20261012 / ops-20261013)(10-12)
날짜: 2026-10-12 출처: SuperNPUBench/ ops-20261013 릴리스
- 모니터링 기간 내 4건 커밋: rms_norm이 16 KiB 타일로 R=8192 리덕션 분해(#223), gfsim RES_CHECK 파일 IO 락스텝 안전 수정, matmul_test를 4-PE로 단순화(Local-A / Shared-B, #222), #225 병합.
- 2건 ops 릴리스: ops-20261012와 ops-20261013——후자는 main
4c66c2ab기반, 스위트별 2개 소형 패키지로 분할, 전체 컴파일 산출물 ELF 포함(gfrun 게이트 미설정); 릴리스 노트에서 “툴체인 네 항목 변동 없음, 이전 버전과의 산출물 차이는 전부 코드에서 기인”임을 명확히 밝힘.
4.2 커뮤니티 계정: SwHwCodesign 6개 AI-* 자리 표시자 저장소; TileLang-Awesome 접근 불가(10-12)
날짜: 2026-10-12 출처: SwHwCodesign 계정/ 원래 TileLang-Awesome 링크
- 10-11에 등장한 서드파티 정리 저장소 TileLang-Awesome(이전 보고 기간에 수록)은 현재 접근할 수 없음(404, 동일 이름의 신규 저장소도 검색되지 않음). 같은 계정이 10-12 저녁에 6개의 플레이스홀더 저장소(AI-Chip / AI-Compiler / AI-Infra / AI-Instruction / AI-Kernel / AI-Model, 모두 초기 커밋)를 새로 생성 — 서드파티 자료 구축은 아직 형태를 갖추지 못함.
- tilelang-debugger(superAngGao)는 모니터링 기간 내 업데이트 없음(최근 커밋 10-10). 다른 신규 서드파티 도구 저장소는 발견되지 않음.
4.3 미디어와 학술: Google News / HN / arXiv 모니터링 기간 내 신규 0건(10-12)
날짜: 2026-10-12 출처: Google News/ Hacker News 검색/ arXiv 검색
- 중영문 여러 쿼리 모두 모니터링 기간 내 신규 히트 0건: 10-11의 tech-insider 종합 기사와 10-10 / 10-11의 Zhiwei 심층 해설은 모두 이전 호에 수록되었고, 이번 모니터링 기간에는 후속 신규 기사가 없음.
- Hacker News 96시간 0건, arXiv 신규 프리프린트 없음(최신 관련 항목은 여전히 07-24의 TileSight 분석 모델 논문).
5. 트렌드 관찰
5.1 ROCm 보강: 디바이스 예산 인식 관련 두 건의 수정
- 이번 모니터링 기간의 두 건의 ROCm 수정은 동일한 주제, 즉 “디바이스 현실에 맞춘 적응”을 가리킴: gfx942의 64 lane 웨이브프런트(#3152)와 LDS/VGPR 예산(#3149). 둘 다 사용 사례 측(examples)에서 수정되었고 커널 내부 생성 전략은 건드리지 않았으며, 이는 ROCm 지원이 여전히 “실행 성공 및 회귀 추가” 단계에 있음을 보여줌. 메인 저장소의 최근 SM121, SM89 다세대 동시 수정과 합류한 후, TileLang의 백엔드 토폴로지는 “NVIDIA 우선”에서 “다중 디바이스 동시 유지보수”로 전환 중임.
5.2 SM89가 두 번째 플랫폼으로: 하루 6건 머지에 4건 심사 중
- TileOPs가 하루 만에 6건의 SM89 머지(어텐션 프리필/디코드, 선형 어텐션, Norm 전면 커버)에 4건의 SM89 심사 중을 더함. RTX 40 시리즈와 L40 / 6000 Ada급 워크스테이션이 연산자 라이브러리 측에서 Hopper와 나란히 상시 목표가 됨. 평가 측도 동기화(야간 기준선 +10 항목)되어, SM89의 정확성 커버리지는 “컴파일 가능”에서 “회귀 보유”로 나아가는 중.
5.3 Ascend 마이그레이션 라인은 정확성 마무리 단계로 전환
- NSA 정방향/역방향 두 건의 대규모 마이그레이션 머지(합계 48개 계층별 테스트), 캐시 3분할과 cross-warp allreduce 그룹화 수정이 이어짐 — Ascend 라인의 중심이 “예제와 연산자 이식”에서 “정확성, 캐시 의미론, 컴파일 사용성”의 마무리 감사로 이동. #3469 수정 직후 커뮤니티 피드백에 따른 정확성 수정(#3494)이 등장해 협업 루프가 정상 작동함.
5.4 릴리스 침묵기의 기준선 성장
- 메인 저장소는 09-30 v0.1.15 이후 신규 릴리스 없음. 조직 오픈 PR은 총 507건(이전 호 498건), 단일 모니터링 기간 신규 41건, 머지 20건 — “머지 게이트는 조이고 심사 대기열은 두꺼워지는” 구조가 지속됨. TileOPs의 SM89 확장과 야간 기준선 +10은 정식 버전이 침묵하는 동안에도 머지와 평가 두 라인의 “정적 준비”가 계속 축적되고 있음을 보여주며, 릴리스 창이 임박했다는 판단은 변함없음.
부록: 자료와 검증 설명
| 출처 | 검증 결과 |
|---|---|
| tile-ai 조직(29개 저장소) | 모니터링 기간 내 9개 저장소에 푸시 발생: tilelang(병합 7, 리뷰 중 업데이트 29), TileOPs(병합 10, 리뷰 중 업데이트 16), TileFoundry(병합 1, 리뷰 중 1), tilelang-ascend(병합 2, 리뷰 중 업데이트 7), tilelang-hygon(3개 브랜치 푸시, #18 업데이트), tilelang-sunrise(CI 의존성 2건 리뷰 중), TileOPs-nightly(스냅샷 1건), tilelang.github.io(빌드 3건), TileOPs.github.io(배포 1건); 조직 모니터링 기간 내 총 병합 20건, 신규 오픈 41건, 리뷰 중 업데이트 59건; 오픈 PR 총 507건(이전 기간 498건) |
| 메인 저장소 tilelang | 병합 7건(#3152, #3149, #3291, #3489, #3487, #3393, #3469); 리뷰 중 업데이트 29건; 오픈 PR 168건(이전 기간 170건); 릴리스 페이지와 태그 모니터링 기간 내 신규 없음(최신 v0.1.15, 09-30) |
| Ascend 저장소(ascend / mlir-ascend) | ascend: 병합 2건(#1701, #1730); 리뷰 중 업데이트 7건; 오픈 PR 228건(이전 기간 230건). mlir-ascend: #199, #200 리뷰 중 업데이트 |
| Hygon / MetaX / Moore Threads | Hygon: 3개 브랜치 푸시(gcy_merge_main, feat/hcu-mqa-fp8, upgrade/upstream-v0.1.15); #18 리뷰 중. MetaX(최근 푸시 09-24), Moore Threads(09-30) 활동 없음 |
| Sunrise | 모니터링 기간 내 Dependabot의 CI 의존성 업그레이드 2건(#7, #8)만 리뷰 중; 오픈 PR 3건 |
| TileOPs | 병합 10건(#2526~#2529, #2531, #2534, #2538, #2540, #2543, #2544); 리뷰 중 업데이트 16건; 오픈 PR 19건 |
| TileOPs-nightly | 신규 스냅샷 1건 6ec0d086(10-13 03:53): 정확성 1042건(61개 스위트, +10), 벤치마크 1947건(52개 스위트), 실패 0건 오류 0건 |
| TileFoundry | 병합 1건(#221); 리뷰 중 1건(#227); 오픈 PR 2건 |
| 채택자 및 커뮤니티 | TileKernels, FlashQLA 모니터링 기간 내 푸시 없음(모두 09-30); TileRT 모니터링 기간 내 동향 없음; tilelang-debugger 업데이트 없음(10-10 이후); SwHwCodesign AI-* 플레이스홀더 저장소 6개(10-12) |
- [1] 메인 저장소 PR 목록 — https://github.com/tile-ai/tilelang/pulls
- [2] ROCm DeepSeek NSA(#3152) — https://github.com/tile-ai/tilelang/pull/3152
- [3] ROCm flash-decoding 타일(#3149) — https://github.com/tile-ai/tilelang/pull/3149
- [4] thread_storage_sync 스레드 경계(#3291) — https://github.com/tile-ai/tilelang/pull/3291
- [5] 테스트 규범 및 개발 가이드(#3489) — https://github.com/tile-ai/tilelang/pull/3489
- [6] CodeRabbit 자동 할당(#3487) — https://github.com/tile-ai/tilelang/pull/3487
- [7] Autotune 실험 보고서(#3393) — https://github.com/tile-ai/tilelang/pull/3393
- [8] PTODSL 리덕션 헬퍼 vendor(#3469) — https://github.com/tile-ai/tilelang/pull/3469
- [9] SM120 MX 블록 스케일링 MMA(#3495) — https://github.com/tile-ai/tilelang/pull/3495
- [10] L2 프리페치 명령(#3490) — https://github.com/tile-ai/tilelang/pull/3490
- [11] cross-warp allreduce 그룹화(#3494) — https://github.com/tile-ai/tilelang/pull/3494
- [12] Host IR 직접 연결 JIT(#3481) — https://github.com/tile-ai/tilelang/pull/3481
- [13] Host IR 실행 통합(#3473) — https://github.com/tile-ai/tilelang/pull/3473
- [14] Windows ARM64 CI(#2962) — https://github.com/tile-ai/tilelang/pull/2962
- [15] 디버그 저하 트레이스 호스트/디바이스 분리(#3496) — https://github.com/tile-ai/tilelang/pull/3496
- [16] Ascend 이중 마이그레이션: NSA Decode(#1701) — https://github.com/tile-ai/tilelang-ascend/pull/1701
- [17] Ascend 이중 마이그레이션: NSA Bwd(#1730) — https://github.com/tile-ai/tilelang-ascend/pull/1730
- [18] Ascend 예제 정밀도 정렬(#1867) — https://github.com/tile-ai/tilelang-ascend/pull/1867
- [19] Ascend situ_glu_grad 예제(#1877) — https://github.com/tile-ai/tilelang-ascend/pull/1877
- [20] Ascend copy_ub_to_gm 수정(#1795) — https://github.com/tile-ai/tilelang-ascend/pull/1795
- [21] Ascend 캐시 수정 배치(#1871 ~ #1873) — https://github.com/tile-ai/tilelang-ascend/pull/1871
- [22] Ascend MLIR: T.Pipelined 파라미터화(#199) — https://github.com/tile-ai/tilelang-mlir-ascend/pull/199
- [23] Ascend MLIR: FloorDiv / FloorMod(#200) — https://github.com/tile-ai/tilelang-mlir-ascend/pull/200
- [24] Hygon 레지스터 파이프라인(#18) — https://github.com/tile-ai/tilelang-hygon/pull/18
- [25] Hygon FP8 MQA 브랜치 — https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-mqa-fp8
- [26] Hygon v0.1.15 업그레이드 브랜치 — <ht
tps://github.com/tile-ai/tilelang-hygon/tree/upgrade/upstream-v0.1.15>
- [27] Sunrise 저장소 — https://github.com/tile-ai/tilelang-sunrise
- [28] MetaX 저장소 — https://github.com/tile-ai/tilelang-metax
- [29] Moore Threads 저장소 — https://github.com/tile-ai/tilelang-musa
- [30] D256 GQA 디코딩(#2529) — https://github.com/tile-ai/TileOPs/pull/2529
- [31] dense GQA 프리필(#2528) — https://github.com/tile-ai/TileOPs/pull/2528
- [32] GDN / DeltaNet 프리필(#2527) — https://github.com/tile-ai/TileOPs/pull/2527
- [33] InstanceNorm 수정(#2526) — https://github.com/tile-ai/TileOPs/pull/2526
- [34] Norm 설정 리팩터링(#2544) — https://github.com/tile-ai/TileOPs/pull/2544
- [35] NSA 압축 영속 CTA(#2531) — https://github.com/tile-ai/TileOPs/pull/2531
- [36] W4A16 GEMM(#2534) — https://github.com/tile-ai/TileOPs/pull/2534
- [37] 16비트 나눗셈(#2543) — https://github.com/tile-ai/TileOPs/pull/2543
- [38] 단일 타일 디코딩 상태 로드(#2538) — https://github.com/tile-ai/TileOPs/pull/2538
- [39] GLA 프리필 상태 스킵(#2540) — https://github.com/tile-ai/TileOPs/pull/2540
- [40] SM89 네 연속(#2546 ~ #2549) — https://github.com/tile-ai/TileOPs/pull/2546
- [41] varlen sinks와 슬라이딩 윈도우(#2537) — https://github.com/tile-ai/TileOPs/pull/2537
- [42] PagedKVCacheWrite(#2539) — https://github.com/tile-ai/TileOPs/pull/2539
- [43] masked FP8 그룹 GEMM(#2521) — https://github.com/tile-ai/TileOPs/pull/2521
- [44] DeepSeek 그룹 top-k(#2535) — https://github.com/tile-ai/TileOPs/pull/2535
- [45] TileFoundry 파서 속성과 레이아웃 슈가(#221) — https://github.com/tile-ai/TileFoundry/pull/221
- [46] TileFoundry scaffold 독립 프로젝트(#227) — https://github.com/tile-ai/TileFoundry/pull/227
- [47] 야간 스냅샷(6ec0d086) — https://github.com/tile-ai/TileOPs-nightly/commit/6ec0d086d3312e34ad92c71f728ff93e6fac04d8
- [48] 스냅샷 실행 메타데이터 — https://github.com/tile-ai/TileOPs-nightly/blob/snapshots/meta.json
- [49] 메인 문서 사이트 — https://github.com/tile-ai/tilelang.github.io
- [50] TileOPs 문서 사이트 — https://github.com/tile-ai/TileOPs.github.io
- [51] SuperNPUBench — https://github.com/PTO-ISA/SuperNPUBench
- [52] SuperNPUBench ops-20261013 — https://github.com/PTO-ISA/SuperNPUBench/releases/tag/ops-20261013
- **
[53]** SwHwCodesign 계정 — https://github.com/SwHwCodesign
- [54] tilelang-debugger — https://github.com/superAngGao/tilelang-debugger
- [55] TileKernels — https://github.com/deepseek-ai/TileKernels
- [56] FlashQLA — https://github.com/QwenLM/FlashQLA
- [57] Hacker News 검색 — https://hn.algolia.com/
- [58] arXiv 검색 — https://arxiv.org/
- [59] Google News 검색 — https://news.google.com/
- [60] 메인 저장소 릴리스 페이지 — https://github.com/tile-ai/tilelang/releases