모니터링 기간: 2026-10-10 07:00 ~ 2026-10-12 07:00 (베이징 시간, 48시간; 월요일 조간, 기간은 직전 데일리 리포트(10-10 대체 휴일 보충 보고) 마감 시점부터 산정하여 토요일, 일요일 이틀 종일과 월요일 새벽을 포함).

출처: GitHub(tile-ai 조직 29개 저장소 푸시 전수 점검, 기간 내 8개 저장소에 푸시 발생: tilelang, TileOPs, TileFoundry, TileOPs-nightly, tilelang-hygon, tilelang-mlir-ascend, tilelang.github.io, TileOPs.github.io; 조직 기간 내 병합 34건, 신규 오픈 42건; 메인 저장소 9건 병합 및 9건 신규 오픈, TileOPs 21건 병합, TileFoundry 2건, Ascend / Hygon / MLIR 어댑테이션 저장소 저장소별 점검, 야간 파이프라인 신규 스냅샷 3건, 문서 사이트 빌드 2건), 커뮤니티 저장소 점검(tilelang-debugger, TileLang-Awesome, PTO-ISA 튜토리얼 및 벤치마크 저장소), Google News RSS 중영문 다중 쿼리(프록시 경유), Hacker News, arXiv, 문서 사이트 및 릴리스 페이지 점검.

이번 호 인덱스

  • 오늘의 중점: 주말 전면 추진——TileOPs 21건 병합(Op 베이스 클래스 리팩터링 확정, MLA 디코딩 가속 배치), 메인 저장소 SM121 착지, TileFoundry 희소 디코딩 TIR 관통(10-10 ~ 10-12)
    1. 핵심 프로젝트 진행
      • 1.1 메인 저장소: SM121(DGX Spark) 점등 블록 스케일링 MMA; CuTeDSL 펜스 선출 및 파라미터 아이덴티티 수정(10-10 ~ 10-11)
      • 1.2 메인 저장소: GEMM dtype 화이트리스트 마무리——지원되지 않는 조합을 CUDA codegen 전에 거부(10-12)
      • 1.3 메인 저장소: For 스텝, JIT 출력 인덱스, popcount, 비트 리덕션 및 테스트 인프라 5건 확정(10-10 ~ 10-11)
      • 1.4 메인 저장소 신규 오픈: Host IR JIT 2건과 병렬 컴파일, 부분 웨이브 리덕션 등 9건(10-10 ~ 10-11)
      • 1.5 TileFoundry: 희소 디코딩 TIR 관통(#223) 및 mesh 스코프 수정(#220)(10-10 ~ 10-12)
    1. 멀티 백엔드 어댑테이션(Ascend / MetaX / Hygon / Moore Threads)
      • 2.1 Ascend: Sparse MLA 마이그레이션 SeeSaw 및 Pipelined 2건 오픈; 장기 실행 마이그레이션 배치 집중 업데이트(10-10 ~ 10-11)
      • 2.2 Ascend: 캐시 클래스 수정 4연타——JIT 특화 키, autotune 키, NPU 및 시뮬레이터 캐시 분리(10-10)
      • 2.3 Ascend: MLIR 경로 전면 성능 리포트 병합——7개 연산자 112개 케이스 전부 통과; NPU 프로파일링 지원 심사 중(10-10)
      • 2.4 Ascend: 메인 저장소 PTO / CANN 심사 중 커밋 클러스터——리덕션 헬퍼 vendor, 문서 요건 및 예제 타깃 선택(10-10 ~ 10-11)
      • 2.5 Hygon: 레지스터 파이프라인 업데이트; MetaX / Moore Threads 무음(10-10 점검)
    1. 생태계 및 채택처
      • 3.1 TileOPs: 21건 병합——Op 베이스 클래스 일괄 리팩터링 확정(-5200줄 규모)(10-10 ~ 10-11)
      • 3.2 TileOPs: MLA / KDA / NSA 성능 배치——SM89 디코딩 최고 약 4.6배 가속, top-k 비트 정렬 및 WGMMA 디스크립터 통일(10-10 ~ 10-11)
      • 3.3 TileOPs: attention sinks 5연타 및 masked FP8 GEMM 신규 오픈(10-10)
      • 3.4 TileOPs: DSA / 페이지드 GQA 수정 배치——공유 메모리 레이스 및 낡은 행 제외(10-10)
      • 3.5 야간 파이프라인: 신규 스냅샷 3건——정확성 1032건, 벤치마크 1947건(52개 스위트) 무실패(10-10 ~ 10-12)
      • 3.6 문서 및 채택처: Op 베이스 클래스 가이드 리팩터링과 함께 사이트 착지; 메인 문서 사이트 빌드 2건; TileKernels / FlashQLA 무음(10-10 ~ 10-11)
    1. 커뮤니티, 튜토리얼 및 활동
      • 4.1 tilelang-debugger: 오프라인 HTML 디버그 리포트 출시, 0.1.12 / 0.1.15 두 버전 검증(10-10)
      • 4.2 PTO-ISA: SuperNPUBench 8건 추진——mega-moe 및 tail_ocp_fp8 전용(10-10 ~ 10-11)
      • 4.3 커뮤니티: TileLang-Awesome 신규 저장소 개설(10-11)
      • 4.4 미디어 및 학술: 중국어 심층 해설 기사 출시; 해외 종합 기사 지속; HN 및 arXiv 신규 없음(10-10 ~ 10-11)
    1. 트렌드 관찰
      • 5.1 완전한 주말: 34건 병합이 토요일, 일요일과 월요일 새벽에 걸쳐 실행
      • 5.2 리팩터링 마무리와 버전 무음: 릴리스 기간 임박
      • 5.3 하드웨어 커버리지 동주 확장: SM121 점등, SM89 성능 정책 성문화
  • 부록: 자료 및 점검 설명
  • 전체 출처 목록

오늘의 중점: 주말 전면 추진——TileOPs 21건 병합 마무리, 메인 저장소 SM121 착지, TileFoundry 희소 링크 관통(10-10 ~ 10-12)

날짜: 2026-10-10 ~ 2026-10-12　출처: TileOPs PR 목록／ Op 기반 클래스 리팩터링(#2520)／ SM121 블록 스케일링 MMA(#3430)／ 희소 디코딩 TIR 관통(#223)／ 야간 스냅샷(fc2a2a91)

모니터링 기간 48시간 내 조직은 34건을 병합했고(메인 저장소 9, TileOPs 21, TileFoundry 2, MLIR 경로 1, 문서 1), 42건을 새로 열었으며, 8개 저장소에 푸시가 있었고, 리듬은 토요일 종일, 일요일과 월요일 새벽까지 이어졌다——병합 분포는 토요일 23건, 일요일 9건, 월요일 새벽 2건으로, 주말 병합량이 이전 모니터링 기간(금요일)의 21건보다 뚜렷하게 많았다.

  • TileOPs 주말 하루 만에 21건 병합: Op 기반 클래스 「직무별 단일 체계」 대규모 리팩터링(#2520, 363개 파일, 순감 약 1865줄)이 릴리스 전에 마무리되었다. MLA, KDA, NSA 세 가지 성능 라인이 동시에 마감되었고, attention sinks 5연속과 masked FP8 GEMM 등 새로 열린 건이 추가되었다(자세한 내용은 3장 참조).
  • 메인 저장소 9건은 플랫폼 확장과 수정이 중심: SM121(GB10 / DGX Spark) 블록 스케일링 MMA 점등(#3430), GEMM dtype 화이트리스트 마무리(#3245, 10-12 새벽 병합), CuTeDSL 펜스 선출 및 NVRTC / CuTeDSL 파라미터 아이덴티티 수정(#3400 / #3414), For 스텝과 JIT 출력 인덱스 등 수정 배치(자세한 내용은 1장 참조).
  • TileFoundry는 희소 디코딩을 TIR까지 밀어붙였다: #223(84개 파일, +4424/-1543)은 DSA 디코딩의 손실 루프 작업량, 조회 청구 과다와 단계별 업데이트 오류를 수정하고, 아래첨자 방식 비동기 수집, Bitcast와 Where를 구현했다(자세한 내용은 1.5 참조).
  • 야간 파이프라인이 세 번의 스냅샷을 연속 생성(최신 병합 지점 #2525): 정확성 1032건, 벤치마크 1947건(52개 스위트), 실패 0건 오류 0건. 이전 스냅샷 대비 각각 2건 증가했고, 이틀에 세 번의 스냅샷이 중단 없이 이어졌다.
  • 미디어 측: 중국어권에서 TileLang을 서사의 핵심으로 삼은 심층 해설 한 편이 등장했고(Zhiwei, 10-10 펑원 / 10-11 Zhihu), 해외 종합 기사는 보도 흐름을 이어갔다(자세한 내용은 4.4 참조).

1. 핵심 프로젝트 진행 상황

1.1 메인 저장소: SM121(DGX Spark) 블록 스케일링 MMA 점등; CuTeDSL 펜스 선출과 파라미터 아이덴티티 수정(10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: SM121 블록 스케일링 MMA(#3430)／ CuTeDSL 펜스 선출(#3400)／ NVRTC / CuTeDSL 파라미터 아이덴티티(#3414)

  • SM121 지원 추가(#3430, 10-11 00:32 머지): SM121(GB10, DGX Spark)은 실제로 SM120급 warp 블록 스케일링 MMA를 지원하지만, tl::sm120_mma_sync_blockscaled는 CUTLASS_ARCH_MMA_SM120A_ENABLED만 검사하며(이는 __CUDA_ARCH__ == 1200일 때만 정의됨), 이로 인해 sm_121a에서 모든 블록 스케일링 MMA가 OMMA.SF 대신 BPT.TRAP으로 컴파일되었다. 이번 커밋은 arch와 패밀리 기준으로 SM120 / SM121 타깃을 동시에 허용하도록 변경하고, NVF4 테스트를 SM12x 패밀리 기준으로 게이팅하며, CuTeDSL TMA 아키텍처 테이블에서 sm_121 / sm_121a를 허용한다. 작성자는 DGX Spark(드라이버 580.178.04, CUDA 13.0)에서 실측 통과했다.
  • 펜스 선출(#3400, 10-10 10:47 머지): sync_grid()는 이전에 모든 tid.x == 0 스레드를 선출하여, 2D / 3D 스레드 블록에서 여러 스레드가 동일 블록 도달을 중복 발행하고 펜스가 조기에 해제될 수 있었다. 이제 물리 스레드 (0, 0, 0)만 선출하며, 카운터 / 세대 프로토콜과 협력적 시작은 변경되지 않는다. 회귀 테스트는 1D / 2D / 3D 그리드와 반복 재사용을 커버한다(RTX 5090, CuTeDSL 4.7.0).
  • 파라미터 아이덴티티(#3414, 10-10 13:58 머지): NVRTC와 CuTeDSL 호스트 래퍼는 이전에 원시 Var 이름으로 시그니처를 생성하고 출력 이름으로 파라미터를 매칭하여, 중복 이름(stream, device_id 등)과 디바이스 측 이름 변경이 서로 간섭할 수 있었다. 이제 Var 아이덴티티 기준으로 래퍼 파라미터 이름을 할당하고 위치 ABI를 유지하며, CuTeDSL은 동일 커널을 서로 다른 스칼라 바인딩으로 반복 호출할 수 있게 되어 동일 이름의 shape / stride 심볼을 구분할 수 있다.

1.2 메인 저장소: GEMM dtype 화이트리스트 마무리 — 지원되지 않는 조합은 CUDA codegen 전에 거부(10-12)

날짜: 2026-10-12　출처: GEMM dtype 화이트리스트(#3245)

  • 일부 T.gemm 피연산자 dtype 조합(예: bf16 x bf16 → fp16)은 이전에 프런트엔드에서 허용되어 nvcc 내부에서 불투명한 static_assert로 실패했다(결함 티켓 #3004). 이번 커밋은 Ampere / Ada mma.sync 경로에 명시적 dtype 화이트리스트를 추가하여, 프런트엔드(tileop/gemm)와 C++ 명령 선택 양쪽에서 강제하며, 지원되지 않는 조합은 CUDA codegen 전에 명확한 메시지와 함께 거부된다.
  • 부정 예시 1건과 긍정 예시 3건의 회귀 테스트를 첨부했으며, 지원되는 조합과 비 Ampere 타깃의 동작은 변경되지 않는다. 이 커밋은 10-12 01:21에 머지되어 모니터링 기간 마감 전 최신 커밋이다.

1.3 메인 저장소: For 스텝, JIT 출력 인덱스, popcount, 비트 리덕션 및 테스트 인프라 다섯 건 확정(10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: For 루프 스텝(#3427)／ JIT 출력 인덱스(#3318)／ popcount(#3300)／ 비트 리덕션(#3332)／ 테스트 인프라(#3313)

  • For 스텝 크기（#3427, 10-10 12:54 병합）: 루프 전달 충돌 분석이 이전에는 스텝 크기를 항상 1로 간주했으나（주석 「기본값 1, 미지정 시」와 상충），비단위 / 부호 스텝 크기에서 유효 반복 쌍 판정이 왜곡되었습니다. 이제 명시적 스텝 크기에 따라 반복 쌍 범위를 축소하고, CUDA와 HIP codegen 모두 스텝 크기를 유지하며 CodeGenC에 위임하여 발행하고, 루프 전개 시의 어휘적 할당 스코프도 함께 고려합니다.
  • JIT 출력 인덱스（#3318, 10-10 12:02）: out_idx=[-1]이 파라미터 개수가 다른 커널 간에 재사용될 때 정규화로 인해 오염되었습니다（2-파라미터 커널이 [-1]을 [1]로 정규화하고, 4-파라미터 커널이 이후에 입력 슬롯을 잘못 점유）. 이제 정규화 전에 호출자 목록을 복사하고, 파라미터 개수와 동일한 스칼라 인덱스를 거부하며, 유효 범위를 제공합니다.
  • popcount（#3300, 10-10 12:02）: 부호 있는 정수와 8 / 16비트 좁은 정수의 CUDA 하강을 보완했습니다. 먼저 원래 폭에 따라 부호 없는 정수로 변환한 후 확장합니다（int8(-1)은 32비트가 아닌 8비트로 계산）. 64비트는 __popcll으로 처리하며, 8가지 dtype 회귀를 커버합니다. #2984를 수정했고, H100（CUDA 13.0）에서 16개 테스트 전부 통과했습니다.
  • 비트 리덕션（#3332, 10-11 14:13）: reduce_bitand / bitor / bitxor가 이전에는 대상 버퍼만 검사하고 부동소수점 소스를 허용했습니다（[1.0, 2.0, 4.5, 8.9]가 경고 없이 1 2 4 8 = 15로 조용히 잘림）. 이제 소스와 대상 모두 정수 / 불리언 규칙에 따라 거부합니다（refs #3301）.
  • 테스트 인프라（#3313, 10-11 18:25）: 공유 IR 검사 도우미를 추가했습니다（노드 수집, 정확한 연산자 호출 수 계산, 심볼별 외부 호출 수집）. issue와 transform 테스트를 로컬 순회 코드에서 도우미로 마이그레이션했습니다. lower_ldgstg 테스트에 구조적 단언을 추가하고, 선택된 transform 테스트는 sm_90 대상을 고정하고 가용성 스킵을 추가했습니다.

1.4 메인 저장소 신규 오픈: Host IR JIT 2건과 병렬 컴파일, 부분 웨이브 리덕션 등 9건 (10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: 호스트 IR 통일（#3473）／ 직접 연결 Host IR JIT（#3481）／ 병렬 JIT 컴파일（#3476）／ 부분 웨이브 리덕션（#3478）

  • Host IR 라인（#3473, #3481）: 먼저 CUDA와 CuTeDSL 두 실행 경로를 Host IR로 통일하고（#3473）, 이어서 직접 연결 Host IR JIT와 네이티브 캐시를 시도합니다（#3481, 실험적, MIR 방향）. JIT 호스트 측 재구성이 메인 저장소의 새로운 메인 라인이 됩니다.
  • 컴파일 처리량（#3476）: autotuner가 상호 독립적인 JIT 구성을 병렬로 컴파일하여 큰 구성 공간의 튜닝 월클록 시간을 단축합니다.
  • 수정류 신규 오픈: 부분 웨이브 스칼라 리덕션（#3478）, 호스트 표현식 강제 변환 그룹화（#3477）, 런타임에야 트립 카운트를 알 수 있는 파이프라인 에필로그 전개（#3472）, Ascend 예제 대상 선택 가능（#3474）. 2건의 WIP（#3482 희소 역방향 그래디언트 원자적 누적, #3483 MoE 스테일 라우팅 전문가）가 10-11 심야에 오픈되었습니다.
  • 큐 현황: 메인 저장소 오픈 PR은 총 170건입니다（지난 회차 177건, 7건 감소）.

1.5 TileFoundry: 희소 디코딩 TIR 관통; mesh 스코프 수정 (10-10 ~ 10-12)

날짜: 2026-10-10, 2026-10-12　출처: 희소 디코딩 TIR 관통（#223）／ mesh 스코프（#220）／ 파서 속성（#221）

  • #223（10-12 00:52 머지, 84개 파일, +4424/-1543）: #218의 DSA 디코딩 잔여 문제를 이어받아 — 손실된 루프 작업량, 과도하게 청구된 lookup 비용, 단계별 업데이트 오류를 수정하고, 인덱스 기반 비동기 수집, Bitcast 및 Where 하강을 적용했습니다. 할당 제약은 단일 접근자(Storage 반환)로 변경되고, 캡처 업데이트는 tile(16B 공유 메모리)에 직접 기록되며, 파라미터 초기화의 캐리 값은 독립 저장소와 사전 복사를 통해 입력을 보장합니다. 16개 FixtureCase 그룹이 각각 CLI 호출과 출력 검사를 포함하여 가드하며, 생성 튜토리얼 보고서를 동기적으로 갱신합니다.
  • #220（10-10 11:59 머지）: #214(창 시작점이 이미 로드된 스칼라일 때 ISL 공간 비교 충돌)와 #215(mesh 외부 초기화 후 루프 내 재바인딩 시 이전 값이 조용히 반환되는 문제)를 수정했습니다. HIR 호출이 mesh 스코프를 벗어나면 즉시 오류를 보고하고, 타입화된 호출 지점별 mesh 전파가 추론 캐시 키에 포함됩니다.
  • #221 신규 오픈(검토 중): 파서가 연산자 속성과 어휘 레이아웃 슈거를 지원합니다.

2. 멀티 백엔드 어댑테이션 (Ascend / MetaX / Hygon / Moore Threads)

2.1 Ascend: Sparse MLA 이관으로 SeeSaw와 Pipelined 두 건 오픈; 장기 이관 배치 집중 업데이트 (10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: Sparse MLA SeeSaw（#1869） / Sparse MLA Pipelined（#1870） / NSA Decode（#1701） / WY Fast（#1827）

  • Sparse MLA 포워드 두 건 신규 오픈 (#1869 / #1870): SeeSaw(시소 이중 소비자)와 Pipelined(파이프라인) 두 버전의 Sparse MLA 포워드 연산자가 동시에 오픈되었습니다 — DeepSeek 희소 어텐션 연산자의 Ascend 이식 라인이 계속 심화되며, 두 버전이 병존하여 비교됩니다.
  • 장기 이관 배치 집중 업데이트 (10-10): NSA Decode(#1701, 08-27 오픈), NSA Bwd(#1730), WY Fast 포워드/백워드(#1827 / #1828), Chunk O Forward(#1866) — “먼저 이관, 이후 연마” 장기 PR의 과반이 진전되었습니다.
  • 문서 및 테스트: #1657(T.tile.relu 등의 문서 문자열과 사용 예), #1867(예제 정밀도 검사 정렬)이 동기적으로 업데이트되었습니다. Ascend 어댑테이션 저장소의 오픈 PR은 총 230건입니다(지난 기간 225건).

2.2 Ascend: 캐시류 수정 4연타 — JIT 특수화 키, autotune 키, NPU와 시뮬레이터 캐시 분리 (10-10)

날짜: 2026-10-10　출처: 튜닝 파라미터 JIT 키 반영（#1871） / autotune 키 정렬（#1872） / 캐시 분리（#1873） / 저장소 재작성 디커플링（#1874）

  • 4건이 같은 날 오픈되었습니다: #1871은 튜닝 파라미터를 JIT 특수화 캐시 키에 포함시키고, #1872는 autotune 캐시 키를 유효 프로파일링 파라미터와 일치시키며, #1873은 Ascend NPU와 시뮬레이터 캐시를 분리하고, #1874는 저장소 재작성이 더 이상 호스트 NPU 가용성에 의존하지 않도록 합니다.
  • 공통적으로 “동일 파이프라인이 실기와 시뮬레이터, 서로 다른 튜닝 구성 간에 캐시가 섞이는” 유형의 문제를 겨냥합니다 — Ascend 측 캐시 시맨틱이 강화되었습니다(커밋에 자동 표시).

2.3 Ascend: MLIR 경로 전체 성능 보고서 머지 — 7개 연산자 112개 사례 전부 통과; NPU 프로파일링 지원 검토 중 (10-10)

날짜: 2026-10-10　출처: 전체 보고서（#197） / NPU 프로파일링 지원（#198）

  • #197 (10-10 14:34 머지): TileOPs 전체 성능 보고서 업데이트 — 7개 연산자, 112개 정확성 테스트 케이스 모두 통과 (CANN 8.5.0, Ascend 910B2C, msprof 프로파일링); 이전 버전 보고서 정리, T.arange + T.vbrc 특정 시나리오 쓰기 경쟁 조건 설명 및 1차원 arange + reshape 회피 방안 보충.
  • #198 심사 중: 선택적 Ascend NPU 프로파일링 지원 — torch_npu 프로파일러 연동 (Level1, AiCMetrics.PipeUtilization), 타임라인에 tilelang::<kernel> 및 tilelang.launch::<kernel> 표기로 호출 계층 노출, CANN 시작 프로브 첨부; 비활성화 시 기존 시작 동작 유지.

2.4 Ascend: 메인 저장소 PTO / CANN 심사 중 커밋 클러스터 — 리덕션 헬퍼 vendor, 문서 요구사항 및 예제 대상 선택 (10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: PTO 리덕션 헬퍼 vendor (#3469)／ CANN 9.2.0 문서 (#3460)／ 예제 대상 선택 (#3474)

  • PTO 라인: #3469는 업스트림 PTOAS에서 전체 리덕션 헬퍼 모듈이 삭제됨에 따라 SIMT 전체 리덕션 헬퍼를 메인 저장소 PTO 디렉터리로 회수하고 테스트 보충 — 업스트림 인터페이스 변경이 메인 저장소 PTO 경로에 영향을 미친 수정 커밋; #3316 (PTO 실행을 PTO codegen으로 전환)은 계속 심사 중.
  • CANN / 툴체인: #3460 (Ascend 950 백엔드의 CANN 9.2.0 요구 문서 커밋), #3461 (bisheng 래퍼가 임시 파일 정리 스위치 준수), #3468 (L0C fp32에서 f16/bf16 이중 복사 경로 변환), #3466 (스칼라 sqrt 하강).
  • 기타: #3474 (예제 선택적 대상, 10-10 신규 오픈), #3324 (torch_npu 작업 큐), #3326 (aibin 단일 단계 컴파일에서 ld.lld 제거), #3437 (SimdVF dim=1 리덕션), #3455 (SimdVF 다중 버퍼 채우기), #3321 (호스트 연산자와 디바이스 커널 연관 프로파일링).

2.5 Hygon: 레지스터 파이프라인 업데이트; MetaX / Moore Threads 조용 (10-10 확인)

날짜: 2026-10-10　출처: 레지스터 파이프라인 (#18)／ MetaX 저장소／ Moore Threads 저장소

  • Hygon #18 (10-10 업데이트): tilelang-hygon에 레지스터 파이프라인 추가 (기존 공유 메모리 파이프라인과 협력, 파이프라인 단계별 커밋 및 배리어 명령 삽입), warp 수가 4를 초과하면 warp-diverge 활성화; 자체 설명에 따르면 example_gemm이 rocblas 수준에 도달; 48개 파일, +4055/-318, 심사 중.
  • MetaX (최근 푸시 09-24), Moore Threads (09-30) 모두 모니터링 기간 내 새로운 활동 없음.

3. 생태계 및 채택자

3.1 TileOPs: 21건 머지 — Op 기본 클래스 일괄 리팩터링 확정 (-5200줄급) (10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: Op 기본 클래스 리팩터링 (#2520)／ Op 기본 클래스 가이드 (TileOPs.github.io #69)

  • #2520 (10-11 06:36 머지, 자칭 “릴리스 전” 마무리, 363개 파일, +3362/-5227): 10개 커밋이 “먼저 동작, 이후 기계적 이름 변경” 순으로 배열 — 컴파일 경계가 있는 연산자와 없는 연산자가 모두 _run_call을 통해 실행되도록 통일, 실패 시 이 호출에서 대상이 선택되었는지 여부에 따라 롤백 결정; 하위 연산자 호출이 어떤 delegate_for 단계에도 속하지 않으면 복합 호출 전체를 실패시키고 더 이상 조용히 무시하지 않음; _validate_dtypes 제거, select_kernel_key / select_implementation을 key_for로 병합; 컴파일 경계를 manifest에서 도출하도록 변경 (compile_op_names 및 콜드 컴파일 테스트 세트 검증), default_kernel_map 제거; Gelu / Div 등 생성자 매개변수가 호출 스펙에 포함됨.
  • 문서 사이트는 다음 날 《Op 기본 클래스 가이드》(#69, 10-11)로 후속 조치하여 코드 리팩터링과 동기화 유지.

3.2 TileOPs: MLA / KDA / NSA 성능 배치 — SM89 디코딩 최대 약 4.6배까지 가속 (10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: MLA 디코딩 SM89 (#2511)／ MLA 튜닝 분할 확장 (#2524)／ MLA 후보 제거 (#2525)／ WGMMA 디스크립터 (#2523)

  • MLA 디코딩 (#2511, 10-10 22:31 머지): SM89에서 배치 및 헤드 블록 수에 따라 2/4/8/16 KV 분할을 선택(목표 128 CTA, 상한 16), 기존 스캔 및 FP32 분할 병합 커널 재사용 — 단일 요청 bf16 4k 기본 지연 0.158 → 0.034 ms(약 4.6배), 32k 1.170 → 0.162 ms(약 7.2배).
  • MLA 튜닝 면 (#2524 / #2525): 튜닝 후보 분할을 1/2/4/8/16으로 확장(후보 24 → 40, 실측 FlashInfer 대비 2.00x 행 존재); 후보 쿼리 타일과 임시 KV가 초과 공유 메모리 예산에 해당하는 경우 제거(D=512일 때 40 중 20 유지), 콜드 튜닝 및 첫 튜닝이 348.6 s에서 151.6 s로 단축, 실패 후보 0으로 정리.
  • WGMMA 디스크립터 (#2523 / #2504): MLA 디코딩과 DSA 시소 커널의 디스크립터를 lane 0 브로드캐스트 구성으로 변경(HGMMA의 R2UR 대기 제거, 오버플로 없음) — 26개 MLA 벤치마크 0.7% ~ 4.1% 빨라짐; 부수적으로 SM90에서 bs1 부분 RoPE 디코딩이 분할 마지막 블록 부족 시 멈추는 문제 수정.
  • KDA 프리필 (#2510): SM89 단일 시퀀스 4096 ~ 64k token, 최대 8개 값 헤드 시나리오에서 16열 값 타일로 분할 — 4k / 32k / 64k 지연 각각 22.9% / 19.8% / 18.5% 감소.
  • NSA top-k (#2514): 각 압축 블록을 하나의 CTA에 할당(쿼리 헤드가 64행 WGMMA 타일을 채움), 레지스터 내 비트 정렬 네트워크로 전체 행 정렬(홀짝 전치 정렬기 대체), 중요도는 더 낮은 블록 번호로 타이브레이크.
  • 페이지드 GQA (#2512): SM90 warp 특화 TMA 커널(TMA 프로듀서 warp 1개 + 컨슈머 warpgroup 2개가 QK와 PV 중첩), 영구 CTA는 SM당 1개; 캐시 끝에서 잘린 타일의 오래된 V 행을 먼저 0으로 초기화 후 읽기; 패킹 레이아웃이 평균적으로 요청당 최소 하나의 128행 타일을 가질 때 우선 디스패치, H200에서 점유된 4행 소요 시간이 원래의 0.44 ~ 0.88배로 감소.

3.3 TileOPs: attention sinks 5연작과 masked FP8 GEMM 신규 개설 (10-10)

날짜: 2026-10-10　출처: 덴스 GQA sinks (#2515)／ 페이지드 GQA RoPE / 윈도우 / sinks (#2516)／ 패킹 varlen sinks (#2517)／ 페이지드 GQA sinks (#2518)／ 페이지드 KV 수집 (#2519)

  • superAngGao의 attention sinks 시리즈 5건이 같은 날 개설: 덴스 GQA(#2515), 페이지드 GQA의 RoPE / 윈도우 / sinks(#2516), 패킹 varlen(#2517), 페이지드(#2518), 페이지드 KV 수집(#2519) — 어텐션 “싱크 token” 지원이 모든 GQA 형태로 확대.
  • 별도로 masked FP8 그룹 GEMM(#2521, MoE 노선)이 개설됨; Hopper 페이지드 GQA 단일 커널(#2499, 자체 설명 H200에서 17/17 비교 가능 사례가 FlashInfer보다 빠름)은 여전히 리뷰 중.

3.4 TileOPs: DSA / 페이지드 GQA 수정 배치 — 공유 메모리 레이스와 오래된 행 제외 (10-10)

날짜: 2026-10-10　출처: DSA SM90 경쟁 상태(#2503) / 페이지드 GQA 스테일 행(#2480) / 캐시 재구축 배치(#2501)

  • #2503: 구형 SM90 DSA 디코드 커널이 is_kv_valid 마스크를 공유하는 문제(인과 경계 밖 키가 활성으로 판정될 수 있음), cp.async.mbarrier.arrive가 일반 저장을 정렬하지 않는 문제, 소비자 간 행 합 경쟁 상태 등 세 곳의 경쟁 상태 수정; 부수적으로 DSADecodeBasicKernel 테일 버퍼 구축 결함 수정.
  • #2480: 페이지드 GQA 출력에서 캐시 끝단의 스테일 행(K / V에 NaN 포함 시 제로 가중치가 차단하지 못함)을 경로별로 각각 클램프 또는 제로화 — 두 개의 NaN 요청 회귀(페이지 64 / 256).
  • 소항목 배치: #2501 / #2513(csrc 헤더 및 튜닝 설정 변경 시 커널 캐시 재구축), #2522(매 자동 튜닝 검색이 미튜닝 설정 소요 시간에 포함됨), #2502(csrc 네임스페이스 스타일), #2498(KDA 청크 프리필이 SM90에서 벗어남), #2506 ~ #2509(SM89 테스트 그룹 및 샘플링 테스트 마무리).

3.5 야간 파이프라인: 세 개의 새 스냅샷 — 정확성 1032건, 벤치마크 1947건(52개 스위트) 무실패(10-10 ~ 10-12)

날짜: 2026-10-10 ~ 2026-10-12　출처: 야간 스냅샷(fc2a2a91) / 스냅샷 실행 메타데이터

  • 모니터링 기간 내 세 개의 스냅샷: 병합 지점 c3146365(10-11 04:16)에서 fe5995e0(10-11 11:46)으로, 다시 70cfdf84 즉 #2525(10-12 03:40, 실행 38161982847)로.
  • 최신 스냅샷: 정확성 1032건 무실패 무오류(1건 건너뜀은 H200에서 SM89에 적용되지 않는 자동 튜닝 케이스), 벤치마크 1947건(52개 스위트) 무실패 무오류(1건 건너뜀은 지원되지 않는 KDA 설정); 이전 스냅샷 대비 각각 2건 증가.
  • 환경: H200, CUDA 13.2, 드라이버 595.71.05, torch 2.13.0+cu132, TileLang 0.1.12(tileops 0.0.1.dev1286+g70cfdf848).

3.6 문서 및 도입처: Op 기본 클래스 가이드가 리팩터링과 함께 사이트에 반영됨; 메인 문서 사이트 두 건 빌드(10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: 메인 문서 사이트 빌드(9afde05a) / TileRT 벤치마크 문서(#58) / TileKernels / FlashQLA

  • 메인 문서 사이트 두 건 자동 빌드(44ac4d7f, 9afde05a, 10-10 및 10-11); TileOPs 문서 사이트는 리팩터링과 함께 후속 반영(3.1 참조).
  • 도입처 및 릴리스: TileKernels, FlashQLA는 모니터링 기간 내 푸시 없음(모두 09-30에 머무름); 메인 저장소 최신 태그는 여전히 v0.1.15(09-30), 모니터링 기간 내 새 릴리스 없음; TileRT 벤치마크 문서 #58(GLM-5.1과 vllm의 PD 배포 지연 / 처리량)는 검토 중(10-10 업데이트 있음).

4. 커뮤니티, 튜토리얼 및 활동

4.1 tilelang-debugger: 오프라인 HTML 디버그 리포트 출시, 0.1.12 / 0.1.15 두 버전 검증(10-10)

날짜: 2026-10-10　출처: tilelang-debugger / HTML 리포트 커밋(07911c42)

  • 모니터링 기간 내 2건: 통합 소스 접근 캡처(소스 접근 관측으로 고정 lowered trace를 대체하고, 표현식 평가 시맨틱을 유지하며, 동적 스텝 파라미터를 지원하고, 프런트엔드 호환성 프로브와 논리 원자 / 주소 어댑터를 추가) — TileLang 0.1.12와 0.1.15에서 각각 117개 CPU / TIR 테스트를 실행하고, H200에서 42회 수집과 66쌍 베이스라인 / 계측 방출을 수행했으며, 독립 코드와 설치 패키지 검수 재확인을 통과했다. 이후 링크 증거와 진단이 포함된 오프라인 HTML 디버그 리포트를 추가했다.
  • 이 프로젝트는 10-08에 설립된 이후 연속 네 개 모니터링 기간 동안 활성을 유지했다(이번 기간 2건, 이전 기간 9건).

4.2 PTO-ISA: SuperNPUBench 8건 추진 — mega-moe와 tail_ocp_fp8 특화(10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: SuperNPUBench / tilelang-puzzles-ascend

  • 모니터링 기간 내 8건: mega-moe의 combine / dispatch 버전 2(#217), FP4 피크와 MFU 지표 기준 수정, rms_norm 32k 단일 블록 경로 재최적화(gamma 향상, 두 개 R 블록 내 리덕션), tail_ocp_fp8 동적 / 정적 커널의 M32 삼중 가드(inf / zero / special) 보완과 V3 2차원 분할(행 타일 입도 곱하기 N축 추가 분할), matmul 재사용 워크북 갱신 및 perf_data로 이동.
  • 교육 저장소 tilelang-puzzles-ascend는 이번 모니터링 기간에 커밋이 없다.

4.3 커뮤니티: TileLang-Awesome 새 저장소 설립(10-11)

날짜: 2026-10-11　출처: TileLang-Awesome

  • 커뮤니티에 SwHwCodesign/TileLang-Awesome이 새로 추가되었다(10-11 설립, 현재 초기 커밋만 존재) — 「Awesome 목록」 형식으로 TileLang 생태계를 정리하는 움직임이 나타났다. tilelang-debugger의 지속적 활성과 결합해, 제3자가 메인 프로젝트를 둘러싼 도구와 자료 구축을 유지하는 온기가 이어지고 있다.

4.4 미디어와 학술: 중국어 심층 해설 기사 게재, 해외 종합 기사 지속, HN과 arXiv 신규 없음(10-10 ~ 10-11)

날짜: 2026-10-10, 2026-10-11　출처: Zhiwei 「DeepSeek의 이번 오픈소스가 R1만큼 중요할 수 있는 이유는 무엇인가」 / tech-insider.org 종합 기사 / Hacker News 검색 / arXiv 검색

  • 중국어 심층 해설: Zhiwei 기사(10-10 펑원 / 10-11 즈후)는 TileLang을 서사의 핵심으로 삼아 Tile 추상화, 세 단계 프로그래밍 인터페이스와 자동 튜닝 캐시 메커니즘을 설명하고, Ascend 적응이 Ascend C와 「성능이 거의 차이가 없다(A2 / A3)」는 비교를 제시하며, 이 사건의 의미가 「R1에 필적한다」고 평가했다. 기사에서는 TileLang-Ascend 코드를 인용해 Ascend 950의 Cube / Vector 구조와 메모리 아키텍처 유사성을 논의했다.
  • 해외 종합: tech-insider.org 10-11 기사는 Tom’s Hardware, The Next Web 등의 인용을 짜맞춰 「6모듈 Ascend 툴킷」 목록과 CUDA / ROCm / Ascend 삼중 스택 대조를 구성했다 — 이는 보도 흐름의 연장이며 독립 출처가 아니므로 인용 성격으로 수록한다.
  • Hacker News는 모니터링 기간 내 항목이 0건이다(최근 관련은 10-02의 「DeepSeek 오픈소스 Ascend 프로그래밍 스택」 논의). arXiv에는 새 프리프린트가 없다(최신 관련은 07-24의 TileSight 분석 모델 논문).

5. 트렌드 관찰

5.1 온전한 주말: 34건 머지가 토요일, 일요일과 월요일 새벽에 걸쳐 실행

  • 34건 머지는 토요일 23건, 일요일 9건, 월요일 새벽 2건으로 분포했고, 42건 신규 오픈도 하루 종일 분포했다(일요일 밤부터 월요일 새벽까지 신규 오픈 지속). 이전 모니터링 기간(금요일)의 21건과 비교하면 — 주말 머지량이 이전 근무일보다 많았으며, 프로젝트의 「7×24 롤링」 리듬은 이미 안정되었다.
  • 평가 기반도 동시에 롤링한다: 이틀 동안 세 차례 야간 스냅샷이 중단 없이 이어졌고, 정확성 1032건, 벤치마크 1947건 규모는 계속 상승하고 있다.

5.2 리팩터링 마무리와 버전 침묵: 릴리스 기간이 눈앞

  • 세 가지 신호가 겹쳤습니다. TileOPs의 Op 기본 클래스 리팩터링은 스스로 “릴리스 전” 마무리라고 칭하며(한 번에 약 1865줄 순감소), 문서 사이트는 다음 날 《Op 기본 클래스 가이드》로 후속 대응했으며, 메인 저장소 태그는 09-30 v0.1.15 이후 조용히 유지되고 야간 파이프라인은 계속 제로 실패로 평가 기반을 갱신하고 있습니다.
  • 엔지니어링 품질 작업(네이밍 / 캐시 시맨틱 / 리뷰 가드레일)은 이번 모니터링 기간에도 계속 집중적으로 착지했으며, “먼저 기준선을 고정하고 이후 정식 버전을 낸다”는 경로가 뚜렷합니다. 다음 버전 릴리스를 위한 정적 준비는 이미 완료되었습니다.

5.3 하드웨어 커버리지 동시 주간 확장: SM121 점등, SM89 성능 정책 대량 적용

  • 같은 기간에: SM121(GB10 / DGX Spark) 점등 블록 스케일링 MMA(#3430), SM89(RTX 40 시리즈 / L40급)에 대량 디코딩 성능 정책 적용(#2510 / #2511 / #2506), SM90(H100 / H200)은 계속 세분화(#2512 / #2497 / #2503).
  • 하드웨어 중심축이 “단일 Hopper 최적화”에서 “여러 세대 동시 수정, 소비자급 진입”으로 전환되었습니다. 국산 백엔드도 같은 박자로: Ascend 희소 MLA 마이그레이션 2건에 캐시 수정 4연타, Hygon 레지스터 파이프라인 업데이트.

부록: 자료 및 검증 설명

출처 검증 결과
tile-ai 조직(29개 저장소) 모니터링 기간 내 8개 저장소에 푸시 발생: tilelang(머지 9, 신규 9), TileOPs(머지 21, 신규 24), TileFoundry(머지 2, 신규 1 리뷰 중), tilelang-mlir-ascend(머지 1), tilelang-hygon(#18 업데이트), TileOPs-nightly(스냅샷 3건), tilelang.github.io(빌드 2건), TileOPs.github.io(머지 1). 조직 모니터링 기간 내 총 머지 34건, 신규 42건, 오픈 PR 총 498건
메인 저장소 tilelang 머지 9건(#3400, #3300, #3318, #3427, #3414, #3430, #3332, #3313, #3245), 신규 9건(#3472, #3473, #3474, #3476, #3477, #3478, #3481, #3482, #3483), 오픈 PR 170건(이전 기간 177건). 릴리스 페이지와 태그는 모니터링 기간 내 신규 없음(최신 v0.1.15, 09-30)
Ascend 저장소(ascend / mlir-ascend) ascend: 이번 모니터링 기간 내 머지 없음, 신규 6건(#1869~#1874), 장기 마이그레이션 배치 다수 업데이트. 오픈 PR 230건(이전 기간 225건). mlir-ascend: 머지 1건(#197), #198 리뷰 중
Hygon / MetaX / Moore Threads Hygon: #18(레지스터 파이프라인) 10-10 업데이트, 리뷰 중. MetaX(최근 푸시 09-24), Moore Threads(09-30) 활동 없음
TileOPs 머지 21건(#2480, #2497, #2498, #2501~#2514 구간, #2520, #2522, #2523, #2524, #2525), 신규 24건, 오픈 PR 9건(이전 기간 6건)
TileOPs-nightly 신규 스냅샷 3건: c3146365(10-11 04:16), fe5995e0(10-11 11:46), 70cfdf84(10-12 03:40, 머지 포인트 #2525). 최신 정확성 1032건, 벤치마크 1947건(52개 스위트), 제로 실패 제로 오류
TileFoundry 머지 2건(#220, #223), 신규 1건(#221, 리뷰 중)
채택자 및 커뮤니티 TileKernels, FlashQLA 모니터링 기간 내 푸시 없음(모두 09-30). TileRT #58 리뷰 중(10-10 업데이트). tilelang-debugger 2건. TileLang-Awesome 신규 저장소 생성(10-11)
PTO-ISA 조직 SuperNPUBench 모니터링 기간 내 8건, tilelang-puzzles-ascend 업데이트 없음
Google News / Hacker News / arXiv Google News 4건 적중(중국어 심층 해설 1, 해외 집계 1, 배경성 보도 2, 모두 1차 자료 아님). HN 0건. arXiv 신규 프리프린트 없음

전체 출처 목록