모니터링 기간: 2026-10-20 07:00 ~ 2026-10-21 07:00 (베이징 시간, 24시간).

출처: GitHub tile-ai 조직 저장소별 점검(Atom 커밋 전체, tilelang / TileOPs / TileFoundry / tilelang-ascend〔기본 브랜치 ascendc_pto〕/ tilelang-hygon / tilelang-metax〔dev〕/ tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack 포함) + 머지 및 심사 중 실값 검색(is:merged / is:open, 저장소별 그룹 통계) + 릴리스 페이지 및 VERSION 파일 점검 + 통계 브랜치 점검 + 문서 사이트 커밋 점검 + 채택 측 저장소 점검(TileKernels / FlashQLA) + 커뮤니티 신규 저장소 검색 + Google News 중영문 다중 쿼리(프록시 경유) + Hacker News + arXiv.

이번 호 인덱스

  • 오늘의 핵심: 메인 저장소의 다섯 건 컴파일러 정확성 수정이 같은 날 반영되었고, Ascend 측에서는 다섯 건의 역방향 및 재귀 연산자가 보충되었으며, TileOPs는 SM90에서 세 개의 foundry 어텐션 연산자를 warp 특화 프로그램으로 교체했다 (10-20)
    1. 핵심 프로젝트 진행 상황
      • 1.1 메인 저장소: 벡터화 분석에 비영(non-zero) 루프 하한을 포함시키고, unit-step 루프와 혼합 정밀도 guard를 수정 (#3579) (10-21)
      • 1.2 메인 저장소: 공유 메모리 스레드 간 WAW 검출을 이중 스레드 인스턴스로 전환하고, 누락된 배리어를 보완 (#2821) (10-20)
      • 1.3 메인 저장소: 지원되지 않는 MMA tile에 대해 명시적 K 패딩을 요구하여, 조용한 K 축소와 레이아웃 추론 실패를 방지 (#2818) (10-20)
      • 1.4 메인 저장소: fragment reduce에서 차원 유지 목표를 지원하여 (64,1) 목표 크래시를 수정 (#2815) (10-20)
      • 1.5 메인 저장소: TIR 프레디킷이 정수 비트 폭 상승에 따라 동기적으로 재작성되도록 수정 (#3550) (10-20)
      • 1.6 메인 저장소: CuTeDSL scan에 참여 스레드 오프셋 파라미터를 보충하여 cumsum 구성 실패를 수정 (#3576) (10-20)
      • 1.7 메인 저장소 CI: 성능 회귀 작업을 H100으로 고정하고, H20 기종은 fail-fast로 차단하도록 설정 (#3569) (10-20)
    1. 멀티 백엔드 적응 (Ascend / MetaX / Hygon / Moore Threads)
      • 2.1 Ascend 메인 저장소: L0 로컬 복사에서 명시적 pitch를 지원하여, 하나의 GEMM tile을 여러 번의 복사로 구성 가능 (#3511) (10-20)
      • 2.2 Ascend 메인 저장소 CI: 단일 check 내에서 빌드와 테스트 단계를 분리하고, 실패 후에도 나머지 스위트를 계속 실행 (#3564) (10-20)
      • 2.3 Ascend 저장소: DSA 희소 미세조정 역방향 연산자 indexer_bwd 반영, AscendC와 PTO 이중 백엔드 (#1887) (10-20)
      • 2.4 Ascend 저장소: SiTU-GLU 역방향 연산자 저장소 반영, Kimi K3 게이팅 활성화와 대조하고 910B3에서 교차 검증 (#1877) (10-20)
      • 2.5 Ascend 저장소: cann-bench에 LSTM과 GRU 추가, 12 / 8개의 jit kernel로 다층 양방향과 프로젝션 지원 (#1892 / #1902) (10-20)
      • 2.6 Ascend 저장소: CrossEntropy 벤치마크 중국어 버전 보완, en/zh 쌍 관례 유지 (#1801) (10-20)
      • 2.7 Ascend 심사 중: PTO 벡터화 프레임워크 심사 진입 (#3512), tvm_ffi를 pto 기본 실행 백엔드로 지정 예정 (#3516), DeepGEMM 휴리스틱 정렬 (#3577) (10-20)
      • 2.8 Hygon: MLS 윈도 상수 K 원점 수정 병합 (#25), fp4 파티션과 분산 두 건 심사 중 (#26 / #19) (10-20)
      • 2.9 MetaX / Moore Threads / MLIR-Ascend: 모니터링 기간 내 신규 푸시 없음, 세 라인 모두 기존 버전에서 정지 (10-20)
    1. 생태계 및 채택 측 (TileOPs / TileFoundry / 연산자 라이브러리)
      • 3.1 TileOPs: 페이지드 GQA 순방향을 paged warp 특화 프로그램으로 교체, FP8 풀을 16비트 페이지 미러를 통해 읽기 (#2668) (10-20)
      • 3.2 TileOPs: FP8 lightning indexer를 이중 consumer warpgroup이 교대로 64-key tile을 처리하도록 변경 (#2666) (10-20)
      • 3.3 TileOPs: MLA 디코딩을 오프셋 시소 프로그램으로 전환, keys를 SM 수에 따라 분할한 후 log-sum-exp로 병합 (#2665) (10-20)
      • 3.4 TileOPs: 텐서 코어 성능 상한을 피연산자 집합 포락선으로 변경, 야간 배치 실행의 105%/108% 오탐 수정 (#2661) (10-20)
      • 3.5 TileOPs: 백엔드 테스트와 torch-compile 베이스라인이 비 CUDA 기기에서 실행 가능 (#2658) (10-20)
      • 3.6 TileOPs: SM89 단일 행 단일 열을 벡터 GEMM 경로로 전환, 후보를 공유 메모리 예산에 따라 필터링 (#2660) (10-20)
      • 3.7 TileOPs: 세 건의 builder 캐시 및 튜닝 계약 수정 (#2662 / #2663 / #2664) (10-20)
      • 3.8 TileOPs 심사 중 11건: SM89의 NSA top-k, 페이지드 DSA, W4A16 좁은 tile 세 라인 병행 진행 (10-20)
      • 3.9 TileFoundry: 0.0.3 공지를 README에 작성, 테스트 리팩터링으로 컴파일 캐시 재사용 (#247 / #246) (10-20)
      • 3.10 채택 측: TileKernels와 FlashQLA는 모니터링 기간 내 신규 푸시 없음 (10-20)
    1. 커뮤니티, 튜토리얼 및 활동
      • 4.1 커뮤니티 저장소: 중국어 시스템 튜토리얼 tilelang-tutorial 저장소 생성, 연산자, 컴파일 타임 및 두 대형 백엔드 포함 (10-20)
      • 4.2 커뮤니티 저장소: CANN 커뮤니티 ops-tilelang 저장소 생성, 연산자 자산 방식으로 Ascend 구현과 베이스라인 관리 (10-20)
      • 4.3 툴체인: devkit과 VS Code 확장 활발, 메인 저장소에서 lower trace 전체를 이전하는 건 심사 중 (#3583) (10-20)
      • 4.4 커뮤니티 저장소: tilelangGym에서 PTO lowering 민감도 테스트, Hexagon 백엔드 실험 저장소 등장 (10-20)
      • 4.5 문서

버전: 문서 사이트가 4차례 자동 업데이트되었고, 메인 저장소의 VERSION은 여전히 0.1.15(10-21)

    1. 트렌드 관찰

오늘의 주요 내용

이번 모니터링 기간에 26건이 머지되었습니다(tilelang 9, TileOPs 9, tilelang-ascend 5, TileFoundry 2, tilelang-hygon 1). 신규 릴리스는 없습니다. 세 갈래가 동시에 진행되고 있습니다. 메인 저장소는 하루 동안의 대부분의 머지를 컴파일러 정확성에 집중했습니다(벡터화 루프 하한, 공유 메모리 스레드 간 WAW, MMA tile의 K 원자, fragment reduce의 유지 차원, TIR 조건자가 비트 폭에 따라 상향). 이는 기능 확장이 아니라 “조용한 오류” 정리에 해당합니다. Ascend 측의 다섯 건 머지는 전부 학습 측과 재귀 연산자(DSA 희소 미세조정 역방향, SiTU-GLU 역방향, LSTM, GRU)에 해당하며, 2만 1천 줄 규모의 PTO 벡터화 프레임워크 한 건이 검토 중입니다. PTO는 AscendC 이외의 두 번째 실행 경로가 되어가고 있습니다. TileOPs는 SM90에서 페이지드 GQA 순방향, FP8 lightning indexer, MLA 디코딩 세 개의 foundry 연산자를 같은 날 warp 특화 프로그램으로 교체했고, 동시에 텐서 코어 성능 상한의 보정 기준을 피연산자 집합 포락선으로 변경하여 야간 배치 실행에서 나타난 105%의 허위 초과를 수정했습니다.

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

1.1 메인 저장소: 벡터화 분석에 0 아닌 루프 하한 포함(#3579)

10-21 00:27에 병합되었고, 10개 파일 +818/-198이며, #2819의 guard 수정을 이어받는 동시에 #2669, #2616 두 이슈를 닫습니다. 기존 구현은 f(i)와 f(floor(i / W) * W)를 비교하여 표현식이 각 벡터 그룹 내에서 동일한 값인지 판단했는데, 이는 루프가 0에서 시작한다는 암묵적 가정을 내포합니다. 루프 하한이 0이 아닐 때 unit-step 벡터화는 서로 다른 lane의 주소를 같은 그룹으로 계산하게 되어, 브로드캐스트 로드와 전체 그룹 공유 조건의 증명이 더 이상 유효하지 않게 됩니다. 새로운 판정은 루프 최솟값을 분석에 포함하고, 혼합 정밀도 변환을 옮길 때 중첩 guard를 유지합니다. 조용히 수치 결과를 바꿀 수 있는 컴파일러 결함에 해당합니다.

출처: https://github.com/tile-ai/tilelang/pull/3579

1.2 메인 저장소: 공유 메모리 스레드 간 WAW 검출(#2821)

10-20 23:41에 병합되었고, 3개 파일 +153/-19이며, #2698을 수정합니다. 공유 메모리에서 S[tx]와 S[tx + 10]은 단일 스레드에게는 서로 다른 두 주소이지만, 서로 다른 스레드의 쓰기는 겹칠 수 있습니다. 기존 WAW 검사는 두 접근에 동일한 스레드 변수를 재사용했기 때문에 판정을 놓치고 배리어 삽입을 누락했습니다. 수정 방법은 동일 반복의 WAW에 서로 다른 두 스레드 인스턴스를 사용하고 각자의 Bind 정의를 이름 변경하며, 기존 RAW/WAR 모델을 재사용하고 루프 전달 WAW는 기존 처리를 유지하는 것입니다. 테스트는 “삽입해야 하는” 경우와 “삽입하지 말아야 하는” 경우를 모두 포함하며, guarded 쓰기, Bind 파생 인덱스, 서로소 파티션, threadIdx.y/z 치환을 포함하고, 스레드 동기화와 issue 두 세트의 테스트 케이스 95건이 통과했습니다.

출처: https://github.com/tile-ai/tilelang/pull/2821

1.3 메인 저장소: 지원되지 않는 MMA tile에 명시적 K 패딩 요구(#2818)

10-20 16:41에 병합되었고, 5개 파일 +57/-14이며, #2715를 수정합니다. 범용 MMA emitter는 피연산자 레이아웃을 구성하기 전에 K 원자를 조용히 축소할 수 있어, FP16 K=8이 잘못된 결과를 계산하게 됩니다. K=20, K=28 같은 불규칙 크기는 공유 메모리 레이아웃 추론 단계에서 명확한 안내 없이 바로 실패합니다. 이제 emitter 초기화 시 아키텍처별로 고정된 K 원자에 따라 tile의 K를 검증하고, 나누어떨어지지 않으면 ValueError를 발생시키며, 입력 dtype, 필요한 나눗셈 관계, 권장 패딩 값(예: SM80에서 FP16 K=20은 원자 16으로 보고, K=32까지 채우고 남는 원소를 0으로 채우도록 권장)을 포함합니다.

출처: https://github.com/tile-ai/tilelang/pull/2818

1.4 메인 저장소: fragment reduce가 유지 차원 목표 지원(#2815)

10-20 16:35에 병합되었고, 8개 파일 +343/-63이며, #2530을 수정합니다. T.reduce_sum 등이 (64, 64)를 dim=1을 따라 (64, 1)로 축약할 때 컴파일 타임에 크래시합니다(Check failed: vars.size() == InputDim() (2 vs. 1)). 반면 reduce_op.py의 문서 계약은 [X, Y]와 [X, 1, Y] 두 출력 형태를 모두 명시적으로 허용하고 이미 대응하는 기대 형태를 구성하고 있으므로, 이는 오용이 아니라 백엔드의 결함입니다. 수정은 두 lowering 경로를 동시에 처리하고, 유지 차원 목표가 없는 기존 동작을 보존합니다.

출처: https://github.com/tile-ai/tilelang/pull/2815

1.5 메인 저장소: TIR 조건자가 정수 비트 폭 상향에 따라 동기적으로 재작성(#3550)

10-20 09:59에 병합되었으며, 4개 파일에서 +950/-27입니다. 정수 인덱스가 더 넓은 비트 폭으로 승격된 후에도 제어 버퍼 접근을 검사하는 경계 검사는 여전히 이전 비트 폭으로 계산되어, 이론적으로 범위를 벗어난 접근을 통과시킬 수 있습니다. 새로운 IndexPredicateRewriter를 추가하여 ConfigIndexBitwidth와 FlattenBuffer에서 각 접근의 승격 정보를 기록하고, 접근 문장 처리가 끝난 후 관련 guard를 재작성합니다. 매칭은 동등한 덧셈 및 곱셈 형태로 제한되며, 무관한 조건, 데이터 값 표현식 및 고정 폭 경계는 보존됩니다. 인덱스 승격, guard 재작성, 접근 범위, 벡터 엔드포인트 및 충돌 승격의 다섯 가지 테스트를 새로 추가했습니다.

출처: https://github.com/tile-ai/tilelang/pull/3550

1.6 메인 저장소: CuTeDSL scan 참여 스레드 오프셋 보완 (#3576)

10-20 12:49에 병합되었으며, 2개 파일에서 +30/-6입니다. 공유 메모리 scan lowering은 #3536에서 참여 스레드 오프셋을 도입했지만, CuTeDSL 측 helper는 여전히 이전 생성자 인수로 호출을 생성하여 tl.CumSum1D(32, False, 0).run(...)가 직접 생성자 인수 오류를 보고합니다. 같은 배치에서 무효화된 codegen 어서션도 #2815의 두 번의 CUDA CI 실패 원인입니다. 수정 방법은 1차원, 2차원 cumsum/cummax helper에 선택적 thread_offset(기본값 0으로 호환성 유지)을 추가하고, 참여 lane과 warp를 선택하기 전에 오프셋을 빼며, CUDA 전용 테스트에 requires_cuda를 지정하는 것입니다.

출처: https://github.com/tile-ai/tilelang/pull/3576

1.7 메인 저장소 CI: 성능 회귀 작업 H100 고정 (#3569)

10-20 16:00에 병합되었으며, 1개 파일에서 +15/-1입니다. PR 회귀 봇은 커널별로 지연 시간과 베이스라인을 비교하는데, 숫자는 베이스라인을 수집한 동일 모델 카드에서만 의미가 있습니다. 조직은 현재 상용 정확성 CI를 위해 자체 호스팅 H20 용량 모델을 도입하고 있으며, 해당 레이블 집합이 기존 [self-hosted, nvidia]와 겹칩니다. 구분하지 않으면 기준 작업을 조용히 가져가게 됩니다. H20의 FP16 밀집 연산 성능은 약 148 TFLOPS로 H100의 약 990 TFLOPS와 약 7배 차이가 나며, FP64 경로도 크게 축소되어 모든 기준에서 존재하지 않는 회귀가 보고됩니다. 이제 pr-regression 작업은 반드시 h100 레이블을 가져야 하며, 작업 시작 전에 할당된 카드가 실제로 H100인지 검증하고, 레이블이 잘못되면 오해의 소지가 있는 숫자를 게시하는 대신 빠르게 실패하도록 요구합니다.

출처: https://github.com/tile-ai/tilelang/pull/3569

2. 멀티 백엔드 적응 (Ascend / MetaX / Hygon / Moore Threads)

2.1 Ascend 메인 저장소: L0 로컬 복사 explicit pitch 지원 (#3511)

10-20 14:16에 병합되었으며, 24개 파일에서 +1458/-187로 이 모니터링 기간 Ascend 방향에서 가장 큰 변경입니다. 이전에는 하나의 GEMM tile을 단일 L1→L0A/L0B 복사로만 채울 수 있었고, 0이 아닌 대상 원점은 거부되었으며, 각 복사는 자체 영역에서 컴팩트 pitch를 추론했기 때문에 분리하여 옮긴 두 절반이 대상 전체 블록으로 합쳐지지 않았습니다. MajorK/MajorMN 전치 로드도 제한되었습니다. 새로운 Ascend 전용 파라미터 T.copy(..., l0_pitch=...)를 추가하여 fractal 블록별로 L0 pitch를 지정하고, 호출자가 소비자 GEMM의 컴팩트 pitch(동적 테일 블록 포함)를 제공합니다. 이 파라미터를 생략하면 컴팩트 영역의 원래 동작을 유지합니다. 예제에서 64행 A tile의 두 절반 모두 l0_pitch=4를 씁니다. 정렬된 대상 오프셋 lowering, 물리적 경계와 전치 그룹 검증, 할당 및 버퍼 버전 간격도 함께 처리합니다.

출처: https://github.com/tile-ai/tilelang/pull/3511

2.2 Ascend 메인 저장소 CI: 단일 check 내 빌드와 테스트 분리 (#3564)

10-20 12:06에 병합되었으며, 1개 파일에서 +61/-9입니다. Ascend의 빌드와 테스트 결과를 동일한 Test for Ascend PR 검사로 통합하고, 단계별로 실패 단계를 표시하며, workflow summary에 각 단계 결론을 나열합니다. CANN 환경은 한 번만 구성하고 한 번만 빌드하며, AscendC 예제, PTO 예제 및 Ascend 단위 테스트는 동일 job의 세 단계로 각각 실행됩니다. 빌드가 성공하고 워크플로가 취소되지 않은 경우, 한 스위트가 실패해도 나머지 스위트를 계속 실행하며 최종적으로 전체는 여전히 실패로 판정됩니다. CANN 변수는 GITHUB_ENV를 통해 전달됩니다.

출처: https://github.com/tile-ai/tilelang/pull/3564

2.3 Ascend 저장소: DSA 희소 미세 조정 역방향 연산자 indexer_bwd (#1887)

10-20 15:26에 병합되었으며, 파일 2개 +1294줄입니다. Ascend NPU를 위해 indexer_bwd를 추가하여 DSA 희소 미세조정을 지원하며, AscendC와 PTO 두 백엔드를 동시에 지원합니다. 연산자는 lightning-indexer KL 미세조정 손실의 dIndexQ, dWeights, dIndexK에 대한 그래디언트를 계산하며, 가변 길이 패킹 시퀀스를 지원합니다. 구현은 세 단계로 나뉩니다. 지속화 GEMM의 빠른 kernel은 연속 및 공유 행 인덱스 레이아웃을 대상으로 하며 파이프라이닝과 더 적은 이동을 갖추고, 범용 kernel은 임의의 인덱스 레이아웃을 지원하며 -1 불법 인덱스의 그래디언트 누적을 포함하고, 별도로 cast kernel을 두어 정밀도 변환을 처리합니다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1887

2.4 Ascend 저장소: SiTU-GLU 역방향 연산자 (#1877)

10-20 15:22에 병합되었으며, 파일 5개 +854줄입니다. 의미론적으로 CANN ops-nn의 aclnnSituGluGrad와 정렬되며, Kimi K3 Stable LatentMoE의 게이트 활성화 역방향에 사용됩니다: t = tanh(gate / beta), s = sigmoid(gate), situ_a = beta * t * s, 업스트림 분기는 linear_beta가 양수인지에 따라 tanh 또는 직통 경로를 타며, grad_gate와 grad_up을 출력합니다. 정밀도 기준은 독립 구현한 fp64 PyTorch 참조이며, 910B3에서 공식 aclnn과 교차 검증하여 최대 상대 오차가 약 3e-3으로 bf16 반올림 예상치에 부합합니다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1877

2.5 Ascend 저장소: cann-bench에 LSTM과 GRU 추가 (#1892 / #1902)

두 건이 같은 날 병합되었으며 각각 +2304줄과 +2267줄입니다. LSTM은 i/f/g/o 네 게이트와 cell state 시계열 재귀를 커버하며, 다층 스택(numLayers ≤ 8), 양방향, batchFirst 레이아웃, 프로젝션 차원 축소(projSize > 0, LSTMP)와 선택적 bias를 지원하고, 가중치는 “레이어별 방향별 하나의 tensor” 방식의 TensorList로 전달되며, 출력은 (y, hn, cn)이고 12개의 @tilelang.jit kernel로 구현됩니다. 주력은 융합된 seq 경로를 타며, (레이어, 방향)당 한 번의 launch로, 커널 내에서 T.serial(S)에 따라 시간 스텝별로 이중 GEMM, bias 브로드캐스트, 게이팅과 상태 업데이트를 완료하고, 재귀와 무관한 x 측 게이트 기여는 독립적인 xgemm으로 전체 시간 스텝에 걸쳐 병렬 사전 계산됩니다. GRU는 8개의 kernel로, z/r/n 세 게이트를 하나로 합친 GEMM은 Cube에서, 게이팅과 상태 업데이트는 Vector에서 수행하며, 크로스 코어 데이터는 GM workspace를 통해 중계되고 AUTO_SYNC를 수동으로 비활성화했습니다 — 커밋 설명에 따르면 자동 모드는 시간 스텝당 약 58개의 PipeBarrier를 생성합니다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1892, https://github.com/tile-ai/tilelang-ascend/pull/1902

2.6 Ascend 저장소: CrossEntropy 벤치마크 중국어판 (#1801)

10-20 15:26에 병합되었으며, 파일 2개 +91줄입니다. #1567로 병합된 CrossEntropy benchmark 문서의 중국어판을 추가하며, FA(bench_mark.md / bench_mark_zh.md)와 SFA(README.md / README_zh.md)의 en/zh 쌍 관례를 참조하여 examples/cross_entropy_loss/ 아래에 benchmark_summary_zh.md를 신설하고 최적화 경로 8단계, 내부 성능 비교, CANN 비교, 알려진 제한 사항과 테스트 커버리지를 완전히 번역했으며, 영어판 상단에 언어 전환 링크를 추가했습니다. 순수 문서 변경입니다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1801

2.7 Ascend 심사 중: PTO 벡터화 프레임워크, tvm_ffi 백엔드와 DeepGEMM 정렬

세 가지 메인라인이 동시에 리뷰에 들어갔다. #3512(50개 파일 +21022/-72)는 사용자가 T.SimdVF(lanes=L) 안에서 스칼라 병렬 T.Parallel 프로그램(원소 단위, 제어 흐름, 리덕션)을 작성할 수 있게 하며, 컴파일러가 이를 PTO VMI 명령어로 벡터화하고 PTOAS 컴파일을 거쳐 Ascend 950에서 실행한다. 전체 기능은 여러 PR로 나누어 반영되며, 이번 PR은 파이프라인 프레임워크와 원소 단위 지원의 이정표다. #3516(18개 파일 +1473/-22)은 target="pto"에 execution_backend="auto"를 설정하면 tvm_ffi로 해석되어 기본값으로 설정되게 하며, 호스트 측은 MakePackedAPI와 C host 빌드를 거치고, Ascend 런타임은 aclrtLaunchKernelWithHostArgs를 통해 PTODSL이 컴파일한 aibin을 기동한다. #3577(29개 파일 +1392/-1036)은 DeepGEMM-Ascend의 구성 선택과 epilogue를 보완하여 밀집, 배치, 그룹 GEMM의 융합 bias, Alpha, BF16 무작위 반올림과 FP8 양자화를 포괄한다. 그 외에 #3530(NZ 버퍼의 부분 UB→L1 복사), #3573 / #3580 / #3581 / #3582 그룹의 PTO 프레디케이트, MAD 방향 분석 및 redux 부호성 수정이 있다.

출처: https://github.com/tile-ai/tilelang/pull/3512, https://github.com/tile-ai/tilelang/pull/3516, https://github.com/tile-ai/tilelang/pull/3577

2.8 Hygon: MLS 윈도 상수 K 원점 수정 (#25)

10-20 09:53에 병합, 2개 파일 +346. 루프 언롤 후 MLS 리소스 승격이 상수 윈도 오프셋을 잃어버려 이후 접근이 잘못된 구간에 떨어지는 문제가 있었다. 수정 방법은 단일 축 업데이트 또는 전방 증분 주소 지정을 선택하기 전에 업데이트되지 않은 차원을 먼저 검증하고, 필요 시 보수적으로 윈도를 재초기화하는 것이다. 또한 기존 MLS 테스트에 회귀 커버리지를 추가했다. Hygon 저장소 모니터링 기간 내 리뷰 중인 것은 세 건이다: #26은 fp4 MLS의 warp 파티션 처리(14개 파일 +261/-144), #18은 레지스터 파이프라인과 warp-diverge 지원 보완, #19는 분산 관련 변경이다.

출처: https://github.com/tile-ai/tilelang-hygon/pull/25

2.9 MetaX / Moore Threads / MLIR-Ascend: 모니터링 기간 내 신규 푸시 없음

세 라인 모두 이번 모니터링 기간에 커밋이 없다. MetaX 저장소의 최근 항목은 09-29의 0.1.15 버전 준비, Moore Threads 저장소는 09-30의 0.1.15 공개 릴리스 준비에서 멈춰 있고, MLIR-Ascend의 최근 항목은 10-19의 TVM 격리 수정(이번 모니터링 기간 외)이다. 세 저장소 모두 리뷰 중인 PR 수가 0이며, 새로운 브랜치 활동도 보이지 않는다.

출처: https://github.com/tile-ai/tilelang-metax, https://github.com/tile-ai/tilelang-musa, https://github.com/tile-ai/tilelang-mlir-ascend

3. 생태계와 도입처 (TileOPs / TileFoundry / 연산자 라이브러리)

3.1 TileOPs: 페이지드 GQA 프리필을 paged warp 특화 프로그램으로 교체 (#2668)

10-20 22:41에 병합, 8개 파일 +901/-103. 새로 추가된 GQAPrefillPagedWSKernel(SM90)이 GQAPrefillPagedWithKVCacheFwdOp를 대체한다: 한 번의 append launch로 chunk를 페이지에 기록하고(key는 융합 RoPE 하에서 회전되며, 둘 다 FP8 풀에서 양자화됨) query를 회전시킨 뒤, 이어서 GQAPagedFwdOp의 warp 특화 프로그램이 확장된 캐시에서 어텐션을 수행한다. 이 프로그램은 head dim 256(단일 query slot)과 단일 KV head에서 16개 헤드를 초과하는 query 그룹((T * H, D)의 query를 행으로 평탄화)을 서비스한다. FP8 프로그램이 커버하지 않는 차원 64와 256의 풀은 페이지의 16비트 미러를 통해 읽으며, 단일 회차 상한은 4 GiB다. 요청이 분할될 수 있는 경우 캐시 길이로 정렬하여 가장 긴 요청이 먼저 처리되며, 최대 128개 요청이다.

출처: https://github.com/tile-ai/TileOPs/pull/2668

3.2 TileOPs: FP8 lightning indexer를 듀얼 warpgroup 소비로 변경 (#2666)

10-20 22:41에 병합되었으며, 6개 파일 +448/-2입니다. 원 구현은 하나의 warpgroup이 64-key tile을 곱한 뒤 곱과 같은 길이의 epilogue를 실행했습니다(각 head에 key scale을 곱하고 FMUL과 FADD를 분리). 또한 2.5 TB/s로 창 밖의 -inf를 쓰는 클린 스캔이 한 번 있었는데, 커밋 설명에 따르면 이는 chunk-prefill-kscale 472마이크로초 중 425마이크로초를 차지합니다. 디코드 호출에서는 단일 query가 하나의 CTA를 차지해 16384개의 key를 끝까지 처리하며 289마이크로초가 걸립니다. 새로운 FP8LightningIndexerWSKernel(SM90)은 TMA warp로 64-key tile을 스트리밍하여 두 consumer warpgroup이 번갈아 처리하도록 하며, tile마다 한 번의 m64n256 FP8 WGMMA를 수행하고 consumer epilogue는 csrc helper로 수렴됩니다.

출처: https://github.com/tile-ai/TileOPs/pull/2666

3.3 TileOPs: MLA 디코드가 어긋난 시소 프로그램으로 전환(#2665)

10-20 22:39에 병합되었으며, 7개 파일 +881/-590입니다. MLADecodeSeesawKernel(SM90, 512 + 64 latent, 단일 KV head)이 MLADecodeWithKVCacheFwdOp를 이어받아 MLAPagedFwdOp의 어긋난 시소 방식을 따르며, 요청별로 TMA를 통해 k와 k_pe를 읽습니다. 프로그램 본체는 kernels/attention/mla/seesaw.py로 이전되어 페이지드와 덴스 두 진입점을 제공하며, MLAPagedFwdOp의 분할 구성은 기존과 동일한 저장소를 산출합니다. 한 요청의 key는 min(sm_count // blocks, tiles // 2)개 세그먼트로 분할되어 log-sum-exp로 병합되며, 덴스 세그먼트는 출력 dtype에 따라 부분 결과를 남기고, 16비트 출력은 warpgroup 자체 key 버퍼에 임시 저장된 후 TMA로 기록됩니다.

출처: https://github.com/tile-ai/TileOPs/pull/2665

3.4 TileOPs: 텐서 코어 상한을 피연산자 집합 포락선으로 변경(#2661)

10-20 22:27에 병합되었으며, 4개 파일 +125/-82입니다. 야간 배치 39299002056이 GemmW4A8FwdOp를 tensor_core.fp8 상한의 105%와 108%로 표시했는데, 확인 결과 상한 자체가 낮았습니다. 기존 보정은 무작위 피연산자에서 cuBLAS의 속도를 취했으나, 700 W 전력 상한에서는 피연산자 값이 클럭을 직접 결정하므로 무작위 cuBLAS보다 전력이 낮은 kernel은 해당 속도 이상으로 동작할 수 있습니다. 이제 gemm_throughput.py는 전부 0인 피연산자도 함께 계측하여 각 피연산자 집합의 최고 속도를 calibration_operands에 기록하고, 그중 최댓값을 보정값으로 사용합니다. 이는 HBM 보정의 기존 포락선 규칙과 일치하며, h200.yaml도 이에 따라 재조정됩니다.

출처: https://github.com/tile-ai/TileOPs/pull/2661

3.5 TileOPs: 백엔드 테스트와 torch-compile 기준선이 비 CUDA 디바이스에서 실행 가능(#2658)

10-20 13:17에 병합되었으며, 10개 파일 +241/-48이며, #2595에서 MTT S5000 배치 실행이 드러낸 TileOPs 측 세 가지 실패를 수정합니다. 첫째, torch_musa가 PrivateUse1의 이름을 변경하므로 test_backend.py와 test_kernel_dispatch.py에서 "privateuseone"으로 등록된 reader가 이를 읽지 못했는데, torch.device("privateuseone").type을 사용하도록 변경했습니다. 둘째, torch_musa가 임포트 시 내용을 출력하여 전체 stdout을 비교하는 임포트 검사가 실패했으므로, tests/fresh_interpreter.py를 새로 추가하여 표시가 있는 결과 한 줄만 읽도록 하고 일곱 곳의 해당 검사를 일괄 전환했습니다. 셋째, CUDA가 없는 환경에서 Triton 없이 torch-compile 경로를 다시 사용할 수 있게 되었습니다.

출처: https://github.com/tile-ai/TileOPs/pull/2658

3.6 TileOPs: SM89 단일 행 단일 열이 벡터 GEMM으로 전환(#2660)

10-20 10:35에 병합되었으며, 3개 파일 +87/-11입니다. 기존 CUDA-core GEMV 구현을 SM89로 확장했습니다. 좌측 피연산자 단일 행과 전치된 가중치, 우측 피연산자 단일 열과 비전치 두 경우가 동일한 kernel body를 재사용하며, 두 행의 교차점은 여전히 SM90 보정을 따릅니다. 동시에 선택된 디바이스의 공유 메모리 예산에 따라 튜닝 후보를 필터링하며, 허용되는 최대 staged ring은 96 KiB이고, 실제 생성된 프로그램은 최대 192바이트의 리덕션 임시 저장 공간을 추가로 선언합니다. 단일 행/단일 열 입력이 정렬되지 않은 요소 오프셋에서 시작하여 연속 저장되는 경우 기존 정렬 helper를 재사용해 처리합니다.

출처: https://github.com/tile-ai/TileOPs/pull/2660

3.7 TileOPs: 세 건의 builder 캐시 및 튜닝 계약 수정(#2662 / #2663 / #2664)

세 건 모두 10-20 13:16에 병합되었다. #2662는 SM89 FP8 grouped GEMM builder에 32개 항목 LRU 캐시를 추가했으며, 기존 SM90 builder 및 저장소 규약과 정렬하여 네 개의 형태 파라미터가 동일한 반복 호출이 JIT 팩토리를 재사용하도록 했다. #2663은 dense GLA subchunk prefill의 두 builder에도 마찬가지로 32개 항목 LRU를 추가했으며, 캐시 키는 형태, scale, 활성화 dtype과 gate dtype을 유지하고 launch 스레드 구성은 내부 팩토리 파라미터로 남겨두었다. #2664는 페이지드 GQA 일반 튜닝 supplier의 12개 입력 계약을 복원했다——이전에는 8개 입력을 기대했으나 9개를 반환하여 세 개의 query/key/value scale을 누락했으며, 지원되지 않는 num_split 후보 키도 제거했다. 세 건 모두 디바이스 측 성능 주장이 없고 연산 본체는 변경되지 않았음을 선언한다.

출처: https://github.com/tile-ai/TileOPs/pull/2662, https://github.com/tile-ai/TileOPs/pull/2663, https://github.com/tile-ai/TileOPs/pull/2664

3.8 TileOPs 심사 중 11건: SM89 3개 라인 병행

모니터링 기간 내 업데이트된 심사 중 PR은 SM89에 집중되었다: NSA 희소 어텐션의 top-k 선택(#2675는 선택 범위를 2048 chunk로 확장, #2674는 group48 기준으로 tile 분할, #2673은 요청별 top-k 출력 슬롯 보완), 페이지드 DSA(#2669 최초 지원, #2670 scale 재사용 및 value 공유), W4A16 좁은 tile(#2671은 작은 M에 더 좁은 tile 사용, #2672는 M64까지 포함). 그 외에 #2667은 공유 메모리를 초과하는 GQA 튜닝 tile을 제거하고, #2654는 작은 토큰 수의 W4A8 작업량을 줄이며, #2515는 dense GQA에 attention sinks를 추가하고, #2490은 KV cache append가 포함된 구 GQA 전방을 퇴역시키려 한다.

출처: https://github.com/tile-ai/TileOPs/pull/2675, https://github.com/tile-ai/TileOPs/pull/2669, https://github.com/tile-ai/TileOPs/pull/2672

3.9 TileFoundry: 0.0.3 공지 및 테스트 리팩터링(#247 / #246)

#247은 README의 Latest News에 세 개의 10/2026 항목을 추가했다: 0.0.3 릴리스, Qwen3.8-27B 프리필 예시(#229 / #236), Apple AMX 프리필 예시(#244)이며, 두 개의 08/2026 항목을 접힌 Earlier 목록으로 옮겨 최신 다섯 개를 유지했다. 커밋 설명에 따르면 README가 곧 PyPI long description(pyproject.toml의 readme = "README.md")이므로 공지는 tag를 찍기 전에 반영되어야 한다. #246은 51개 파일 +655/-366의 테스트 리팩터링이다: 로컬과 CI가 ccache와 torch extension 캐시를 재사용하고(PR 간 캐시 격리, job별 통계), 두 곳의 테스트를 호스트 측에서 컴파일한 CuTe 참조로 전환하여 MMA 레이아웃을 비교함으로써 수작업 참조와 수천 개의 파라미터화 케이스를 대체하며, torch는 지연 프록시 import로 변경하고 필요에 따라 dtype을 해석한다.

출처: https://github.com/tile-ai/TileFoundry/pull/247, https://github.com/tile-ai/TileFoundry/pull/246

3.10 채택자: TileKernels와 FlashQLA 모니터링 기간 내 푸시 없음

DeepSeek의 TileKernels와 Qwen의 FlashQLA는 이번 모니터링 기간에 커밋이 없었으며, 최근 시점은 각각 10-14의 v2.1.0 릴리스 병합과 10-14의 SM120 수정 병합에 멈춰 있다. 두 저장소의 심사 중 PR 수는 0이다.

출처: https://github.com/deepseek-ai/TileKernels, https://github.com/QwenLM/FlashQLA

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

4.1 커뮤니티 저장소: 중국어 시스템 튜토리얼 tilelang-tutorial 저장소 생성

huandu/tilelang-tutorial은 10-20에 저장소가 생성되어 푸시되었으며, “제로에서 원까지의 TileLang 중국어 시스템 튜토리얼”을 표방한다. 범위는 연산자 작성, 컴파일 타임 메커니즘, NVIDIA와 Ascend 두 백엔드를 포함한다. 이전의 중국어 자료는 단편 블로그와 공식 문서 번역이 주를 이루었으며, 체계적인 튜토리얼 저장소는 이번이 처음이다.

출처: https://github.com/huandu/tilelang-tutorial

4.2 커뮤니티 저장소: CANN 커뮤니티 ops-tilelang 저장소 생성

Atlas-Ascend-Fork/ops-tilelang은 10-20에 저장소를 생성했으며, CANN 커뮤니티가 Ascend NPU를 대상으로 하는 TileLang 고성능 연산자 저장소로 자리매김하고, “Kernel 구현, 자동화 테스트 케이스, 표준 워크로드와 성능 기준선의 통합 관리” 방식으로 연산자 자산을 축적하며, 목표는 Ascend 측 TileLang 연산자가 재사용 가능하고 검증 가능하며 지속적으로 진화할 수 있도록 하는 것이다. 이는 이전에 Ascend 연산자가 examples/와 cann-bench에 흩어져 있던 상태와 대조를 이룬다.

출처: https://github.com/Atlas-Ascend-Fork/ops-tilelang

4.3 툴체인: devkit과 VS Code 확장 활발, lower trace 전반 메인 저장소에서 이전

easy-tilelang/tilelang-devkit(7월 저장소 생성)과 tilelang-devkit-vscode(10-19 저장소 생성, 10-20 푸시)는 이번 모니터링 기간 모두 활동이 있었으며, 후자는 VS Code에서 lowering pass, IR 차이와 CPU C JIT 컴파일을 탐색하는 것으로 자리매김한다. 이와 호응하여 메인 저장소에서 심사 중인 #3583(13개 파일 +106/-4390)은 lower trace 구현, 전용 환경 설정과 관련 테스트를 메인 저장소에서 전반적으로 분리해 내고 공유 컴파일러 계측 인터페이스만 남겨 추적 도구가 독립적으로 진화하도록 한다. 이는 메인 저장소가 7월 이래로 이어온 “디버깅 도구 외부화”의 연장이다.

출처: https://github.com/easy-tilelang/tilelang-devkit-vscode, https://github.com/tile-ai/tilelang/pull/3583

4.4 커뮤니티 저장소: tilelangGym과 Hexagon 백엔드 실험

PTO-ISA/tilelangGym(10-14 저장소 생성, 10-20 푸시)은 서로 다른 benchmark와 마이크로 벤치마크가 PTO lowering 하에서의 성능 민감도를 테스트하는 데 사용되며, Ascend PTO 라인의 추진과 동기화된다. SoilSoi1/tilelang-hexagon은 10-20에 저장소를 생성했으며 Hexagon 백엔드의 커뮤니티 실험이다. 메인 저장소의 심사 대기열에서 이에 대응하는 것은 바로 #2155 [feat] Add Hexagon HMX backend support(작성자 Calaweh, 5월부터 “각 백엔드 분리”에 대한 후속 경로를 논의한 바 있음)이다. 또한 sablin39/tilelang-cuda-skills(146 스타)는 기간 내에도 계속 업데이트되고 있으며, TileLang의 CUDA 기술 세트를 작성하고 디버깅하는 데 사용된다.

출처: https://github.com/PTO-ISA/tilelangGym, https://github.com/SoilSoi1/tilelang-hexagon, https://github.com/tile-ai/tilelang/pull/2155

4.5 문서와 버전: 문서 사이트 네 차례 자동 업데이트, VERSION은 여전히 0.1.15

tilelang.github.io는 기간 내 네 차례 자동 업데이트되었으며(10-20 12:16, 13:01, 16:08과 10-21 00:49), 모두 메인 저장소의 머지에 따라 트리거되었다. 메인 저장소의 VERSION은 여전히 0.1.15이며, README의 최신 마일스톤은 09-30의 Ascend 950 백엔드에 멈춰 있다. TileOPs의 통계 브랜치는 마지막 머지에 따라 10-20 22:42에 갱신되었고 새로운 tag는 보이지 않으며, TileFoundry의 0.0.3은 아직 tag 생성 대기 중이다.

출처: https://github.com/tile-ai/tilelang.github.io, https://github.com/tile-ai/tilelang

5. 트렌드 관찰

컴파일러 정확성이 이번 모니터링 기간의 주된 흐름이 되었으며, 기능 확장이 아니다. 메인 저장소의 9건 머지 중 5건은 조용한 오류 정리였다: 벡터화가 루프가 0에서 시작한다고 가정(#3579), 공유 메모리 WAW 누락 판정(#2821), MMA emitter의 조용한 K 원자 축소(#2818), fragment reduce 보존 차원 충돌(#2815), 인덱스 비트 폭 향상 후 프레디케이트 미추종(#3550). 이러한 결함의 공통점은 오류를 보고하지 않고 잘못된 결과만 내거나 배리어를 누락한다는 것으로, 사용자가 생산 환경에서 가장 찾기 어려운 유형에 속한다. 같은 날 TileOPs의 #2661도 “성능 상한 기준”을 수정하고 있어, 프로젝트가 “숫자의 신뢰성”을 “빠르게 실행됨”과 동급의 엔지니어링 목표로 삼고 있음을 보여준다.

Ascend 라인이 예제 보충에서 훈련 측으로 전환되고 있다. 5건의 머지 중 3건은 역방향 또는 재귀 연산자(indexer_bwd, situ_glu_grad, LSTM/GRU)이며, situ_glu_grad는 Kimi K3 Stable LatentMoE의 게이트 활성화에 직접 정렬하고 910B3에서 aclnn과의 교차 검증 데이터를 제시한다. #3577의 DeepGEMM-Ascend 구성 선택과 epilogue 정렬과 함께, Ascend 측의 커버리지는 추론 전방에서 훈련 역방향, 상태 재귀와 대형 모델 MoE 게이팅으로 확장되고 있다.

PTO는 AscendC 이외의 두 번째 실행 경로가 되고 있다. 심사 중인 #3512는 2.1만 줄 규모의 벡터화 프레임워크(T.Parallel → PTO VMI, 목표 Ascend 950)이며, #3516은 tvm_ffi를 pto 타깃의 기본 실행 백엔드로 설정하려는 안건이다. 여기에 #3530, #3573, #3581, #3582의 복사 및 프레디킷 수정 묶음을 더하면, PTO 라인의 엔지니어링 규모는 이미 단일 패치 수준을 뚜렷이 넘어섰다. 메인 저장소는 동시에 #3511로 AscendC 측의 L0 로컬 복사 기능을 보강하고 있으며, 두 경로는 대체 관계가 아니라 병행 추진 관계다.

툴체인이 독립 저장소로 외연 확장되고 있다. 하루 만에 세 개의 신규 저장소(중국어 튜토리얼, CANN 커뮤니티 연산자 저장소, Hexagon 백엔드 실험)와 VS Code 확장 한 세트가 등장했으며, 메인 저장소는 lower trace를 통째로 devkit으로 이전했다(#3583, 4390줄 삭제). 커뮤니티 측의 확장 속도는 이미 메인 저장소의 모듈 구획보다 빨라졌으며, 이는 7월의 “멀티 백엔드 언어 방언” 리팩터링 방향과 일치한다. 즉 메인 저장소는 의미론과 IR로 좁히고, 도구와 백엔드 실험은 외부로 분리하는 것이다.

CI 기준의 명시화는 주목할 만하다. #3569는 성능 회귀를 H100에 고정하고 오표기 카드에 대해 빠르게 실패하도록 했으며, #3564는 Ascend 빌드와 테스트를 동일 check의 단계별로 분리했고, #2661은 피연산자 집합 포락으로 텐서 코어 상한을 재표기했다. 세 건 모두 “숫자를 비교 가능하게 만드는” 작업으로, 이후 PR의 진입 기준에 직접 영향을 준다.

부록: 자료 및 검증 설명

출처 검증표

출처 검증 결과
tile-ai/tilelang (main) 모니터링 기간 내 8회 커밋, 9건 머지. VERSION은 여전히 0.1.15. 심사 중 167건, 그중 163건이 모니터링 기간 내 갱신됨
tile-ai/TileOPs (main) 모니터링 기간 내 9회 커밋, 9건 머지. 심사 중 16건, 11건이 모니터링 기간 내 갱신됨. 통계 브랜치 22:42 갱신
tile-ai/tilelang-ascend (ascendc_pto) 모니터링 기간 내 5회 커밋, 5건 머지. 심사 중 236건, 7건이 모니터링 기간 내 갱신됨
tile-ai/tilelang-hygon (main) 모니터링 기간 내 1회 커밋, 1건 머지. 심사 중 3건
tile-ai/TileFoundry (main) 모니터링 기간 내 2회 커밋, 2건 머지. 심사 중 1건. 0.0.3 tag 대기 중
tile-ai/tilelang-metax (dev) / musa / mlir-ascend 세 저장소 모두 모니터링 기간 내 커밋 없음
tile-ai/TileRT / tilescale / DeepStack 세 저장소 모두 모니터링 기간 내 커밋 없음(최신 각각 09-24, 08-03, 10-05)
deepseek-ai/TileKernels / QwenLM/FlashQLA 모두 커밋 없음, 최신 모두 10-14에 정지
tilelang.github.io 문서 사이트 모니터링 기간 내 4회 자동 갱신
커뮤니티 및 서드파티 저장소 신규 저장소 3개(tilelang-tutorial, ops-tilelang, tilelang-hexagon), 활성 3개(devkit, devkit-vscode, tilelangGym)
Google News / Hacker News / arXiv 중영문 다수 조회에서 TileLang 직접 관련 보도 없음. arXiv 최신 관련 논문은 여전히 2026-07-24

보충 설명: 검색 인터페이스가 심사 대기열에 대해 제시한 updated_at 중 64건이 동일하게 10-20 11:30 전후에 몰려 있었으며, 그중 세 건을 표본 확인한 결과 해당 시각에 코멘트, 리뷰, 브랜치 푸시 이벤트가 없어 플랫폼 측의 일괄 재계산으로 판정하고 개발 활동에 포함하지 않았다. 심사 중 건수는 각 저장소의 is:open 집계를 기준으로 한다.

전체 출처 목록

2658 백엔드 테스트와 torch-compile 기준선 CUDA 분리 — https://github.com/tile-ai/TileOPs/pull/2658