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

출처: GitHub tile-ai 조직 저장소별 점검 (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon 브랜치 / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 문서 사이트 3곳) + 머지 및 심사 중 실측 검색 (is:merged와 is:open 두 그룹, 저장소별 그룹 통계) + 커밋 수준 diff 및 릴리스 페이지 tag 점검 + 채택측 저장소 점검 (TileKernels / FlashQLA) + 커뮤니티 신규 저장소 검색 + Google News 중영문 다중 쿼리 (프록시 경유) + Hacker News + arXiv + 산업 매체 보도.

이번 호 인덱스

  • 오늘의 핵심: 35건의 머지 중 메인 저장소는 분석기 라이프사이클에 중점을 두었고, DeepSeek V3.2의 top-k 선택기 범위 초과 쓰기가 CUDA_ERROR_ILLEGAL_ADDRESS로 특정되었습니다. Ascend NZ 레이아웃은 일부 UB→L1 복사(18개 파일)를 연결했습니다. CANN 커뮤니티는 ops-tilelang 연산자 저장소를 새로 개설했습니다.
    1. 핵심 프로젝트 진행 상황
      • 1.1 메인 저장소: 벡터화 재작성 전에 analyzer 복제(#3612)(10-22)
      • 1.2 메인 저장소: 병렬 루프 분할 후 기존 analyzer 바인딩 무효화(#3226)(10-22)
      • 1.3 메인 저장소: deepseek_v32 top-k 선택기 후보 버퍼 범위 초과(#3280)(10-22)
      • 1.4 메인 저장소: CUDA 16비트 역삼각함수 및 이진 수학 폴백 보완(#3172)(10-22)
      • 1.5 메인 저장소: sm8x 희소 메타데이터 비전단사 형태를 거부로 변경(#3125)(10-22)
      • 1.6 메인 저장소: eager 프런트엔드에서 Python 소스 위치 보존(#3558)(10-22)
      • 1.7 심사 중: pass 능력 검사, 부호 없는 mulhi 및 assume 술어(10-22)
    1. 다중 백엔드 적응(Ascend / Hygon / MetaX / Moore Threads)
      • 2.1 Ascend: NZ 버퍼에서 부분 UB→L1 복사 지원(#3530)(10-22)
      • 2.2 Ascend PTO: MAD 방향 술어의 지수적 증가로 컴파일 정체(#3581)(10-22)
      • 2.3 Ascend PTO: SIMT 벡터 지역 변수를 lane별로 할당(#3523)(10-22)
      • 2.4 Ascend 저장소: cann-bench에 Dilation2D 추가, 같은 날 depthwise_conv_2d 롤백(#1916 / #1919)(10-22)
      • 2.5 Ascend 저장소 심사 중: sparse_mla 순전파 및 역전파(#1929 / #1928)(10-22)
      • 2.6 Hygon: 병합 브랜치에서 enable_register_pipeline 제거(10-22)
      • 2.7 MetaX / Moore Threads / MLIR-Ascend: 모니터링 기간 내 푸시 없음(10-22)
      • 2.8 산업 측: SCMP가 SemiAnalysis 보고서를 인용, TileLang이 6개 저장소 중 유일하게 사용 가능(10-22)
    1. 생태계 및 도입 측(TileOPs / TileFoundry / 연산자 라이브러리)
      • 3.1 TileOPs: FP8 페이지드 warp 특화 kernel의 공유 메모리 앨리어싱으로 NaN 발생(#2708)(10-22)
      • 3.2 TileOPs: dense prefill을 영속 순회로 변경, D64 세 소비자(#2706)(10-22)
      • 3.3 TileOPs: SM90 dense GQA decode가 warp 특화 사용(#2694)(10-22)
      • 3.4 TileOPs: 인과 슬라이딩 윈도우에서 하나의 KV head에 속한 두 key tile로 전체 query head 그룹 서비스(#2696)(10-22)
      • 3.5 TileOPs: 페이지드 decode 자동 튜닝이 이전에 block table 하나를 공유(#2703)(10-22)
      • 3.6 TileOPs: masked softmax의 FMA 오버플로 및 varlen RoPE 홀수 반폭(#2699 / #2700)(10-22)
      • 3.7 TileOPs: FP8 grouped GEMM 계층별 scale 임시 저장 및 gate/up 핑퐁 warp group(#2704 / #2693)(10-22)
      • 3.8 TileOPs: 술어 브로드캐스트 및 불리언 피연산자를 16바이트 벡터 읽기로 변경(#2702 / #2697)(10-22)
      • 3.9 TileOPs: SM89 세 경로 같은 날 진행(#2705 / #2691 / #2690)(10-22)
      • 3.10 TileOPs: 리덕션 폴딩 단일 발행 및 빈 DSA 선택 검증(#2689 / #2698)(10-22)
      • 3.11 TileOPs: 벤치마크 기준을 float32에서 계산하도록 변경(#2692)(10-22)
      • 3.12 TileFoundry: sm_89 아키텍처 및 RTX 6000 Ada 디바이스 설치(#259)(10-22)
      • 3.13 TileFoundry: gmem 과금을 풋프린트와 실행 mesh 기준으로 변경, 피호출자 차원 바인딩(#258)(10-22)
      • 3.14 TileFoundry: 디바이스가 kernel 시작을 거부할 때 반드시 오류 발생(#256)(10-22)
      • 3.15 배포판: Sunrise가 0.1.15+sunrise.1.2.0 릴리스 PR 생성(#9, 심사 중)(10-22)
      • 3.16 도입 측: TileKernels 및 FlashQLA 모니터링 기간 내 푸시 없음(10-22)
    1. 커뮤니티, 튜토리얼 및 활동
      • 4.1 CANN 커뮤니티에서 ops-tilelang 연산자 저장소 개설(10-22)
      • 4.2 tilelang/agent: 연산자 최적화를 agent의 SKILL 워크플로로 작성(10-22)
      • 4.3 커뮤니티 신규 저장소: Sail 래퍼, 원격 자동 튜닝 agent 및 튜토리얼 저장소(10-22)
      • 4.4 문서 및 버전: 문서 사이트 3회 업데이트, VERSION은 여전히 0.1.

15 (10-22)

    1. 트렌드 관찰

오늘의 핵심

이번 모니터링 기간에 tile-ai 조직은 35건을 머지했으며(TileOPs 20, tilelang 9, TileFoundry 3, tilelang-ascend 2, TileFoundry 문서 사이트 1), 별도로 71건의 심사 중 PR이 기간 내 업데이트되었다. 메인 저장소의 9건은 지난주 “형상과 기하 검증”에서 분석기 수명주기로 중심이 옮겨갔다: 벡터화 재작성 전 analyzer 복제(#3612, 단일 파일 2줄), 병렬 루프 분할 후 구 바인딩 무효화(#3226, 6개 파일), eager 프런트엔드의 Python 소스 위치 보존(#3558) — 세 건 모두 kernel 계산을 틀리게 만들지는 않지만, 재작성 후 컴파일 시점에서 만료된 상태를 읽게 만든다. 수치 측에서 가장 기록할 만한 것은 #3280이다: deepseek_v32의 top-k 선택기가 radix 후보를 고정 [2, 4096] 공유 버퍼에 임시 저장하면서도 첫 라운드에서 선택된 임계값 버킷 크기를 제약하지 않아, 낮은 가수 비트가 동일한 FP32 점수나 완전히 동일한 입력에서 범위를 넘어 기록되어 CUDA_ERROR_ILLEGAL_ADDRESS로 크래시한다. Ascend 측 #3530은 당일 규모가 가장 큰 건(18개 파일 +993/-256)으로, 압축 NZ 패딩을 저장 정규화까지 유지하고 일부 UB→L1 복사에 대해 C0 그룹 행 간격을 명시적으로 제공한다; #3581은 PTO에서 컴파일을 10분 이상 정체시킬 수 있는 프레디케이트 표현식의 지수적 증가를 수정했다. 생태계 측 두 가지 새로운 라인: CANN 커뮤니티는 10-20에 ops-tilelang 연산자 저장소를 만들고 10-22에 첫 수정을 푸시했다; Sunrise 배포판은 0.1.15+sunrise.1.2.0 릴리스 PR을 올려 TANG 백엔드 wheel을 배포에 포함했다.

1. 핵심 프로젝트 진행

1.1 메인 저장소: 벡터화 재작성 전 analyzer 복제(#3612)

10-22 23:43에 머지되었으며, 벡터화 해당 소스 한 곳만 변경, +2/-1. 루프 벡터화는 IR을 재작성할 때 전달받은 산술 분석기를 재사용하는데, 분석기는 재작성 전의 상수와 구간 제약을 캐시하고 있다; 벡터화가 루프 경계나 인덱스 표현식을 바꾸면 이후 pass는 여전히 구 결론을 얻는다. 수정 방식은 벡터화 재작성에 진입하기 전에 분석기를 복사하여 재작성이 복사본에만 적용되도록 하는 것이다. 이 건에는 회귀 테스트가 포함되지 않았으며, 가치는 #3226과 동일한 근원의 문제를 다른 진입점에서 막는다는 데 있다 — 같은 날 그 건은 병렬 루프 분할 후의 무효화된 바인딩을 처리했으며, 두 건을 함께 보면 분석 컨텍스트의 수명주기가 독립적인 결함 유형으로 정리되고 있다.

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

1.2 메인 저장소: 병렬 루프 분할 후 구 analyzer 바인딩 무효화(#3226)

6개 파일 +55/-19, loop_partition, loop_vectorize, lower_tile_op 세 pass 및 해당 헤더 파일에 관련된다. 병렬 루프가 분할된 후 루프 본문의 이터레이션 변수가 재바인딩되지만, pass는 여전히 분할 전 변수로 분석기를 조회하여 구 변수의 구간과 나눗셈 정보를 얻으므로 단순화와 범위 판정이 왜곡된다. 수정은 분석 컨텍스트를 pass별로 격리하고 루프 분석 스냅샷의 초기화 경로를 단순화했으며, 최소화된 병렬 루프 분석 회귀 케이스 세트를 함께 추가했다. 제출자에는 Codex 공동 서명이 포함되어 있다.

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

1.3 메인 저장소: deepseek_v32 top-k 선택기 후보 버퍼 범위 초과(#3280)

2개 파일 +91/-10, #1351 수정. tl_topk_impl은 radix 후보를 고정 [2, 4096] 공유 버퍼에 담지만, 첫 라운드 radix에서 선택된 임계값 버킷이 4096개 요소를 초과하지 않도록 제약하는 것이 아무것도 없다 — FP32 점수의 낮은 가수 비트가 동일하거나 입력이 완전히 동일할 때 원자적 추가가 버퍼 끝을 넘어 기록되어 CUDA_ERROR_ILLEGAL_ADDRESS로 크래시한다. 수정 방식은 tl_topk 래퍼 계층에서 (batch, 2, seq_len) 크기의 전역 workspace를 할당하여 두 radix 버퍼 모두 seq_len으로 감싸지도록 하는 것이다; 공개된 tl_topk() 시그니처는 변경되지 않는다. 79줄 재현 케이스가 포함되어 있으며, #1351의 입력 구성에 따른다.

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

1.4 메인 저장소: CUDA 16비트 역삼각 함수 및 이진 수학 하향 보완(#3172)

3개 파일 +463/-8. asin, acos, asinh, acosh, atanh, log1p, atan2, hypot, copysign, ldexp, nextafter에 CUDA 전용 하향 규칙이 없었고, 범용 규칙은 float32/float64만 처리하여 float16과 bfloat16 호출이 codegen 단계에서 아무도 처리하지 않았다. 수정은 이들을 CUDAMath에 등록하고 half_t / bfloat16_t 오버로드를 추가한다: 열 개 연산자는 float32로 평가한 뒤 다시 변환하며, nextafter는 16비트 인코딩 기준 정수 비교 스텝으로 바꾸어 fast math에서도 bfloat16 서브노멀을 보존한다. float32/float64의 함수 이름은 그대로 유지되며, common.h 측에 182줄의 템플릿 구현이 새로 추가되었다.

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

1.5 메인 저장소: sm8x 희소 메타데이터 비전단사 형상을 거부로 변경 (#3125)

단일 파일 tilelang/layout/gemm_sp.py +6/-0, #2606 수정. make_cutlass_metadata_layout_sm8x의 interleave는 M이 그룹 크기의 정수 배이고 K가 짝수일 때만 전단사이며, 조건을 만족하지 않을 때 기존에는 조용히 잘못된 메타데이터를 생성하여 희소 GEMM 결과가 아무 표시 없이 틀렸다. 이제 두 조건 모두 명시적으로 거부하도록 변경했다. 원래 PR과 함께 추가되었던 희소 메타데이터 형상 테스트는 리뷰에서 제외 요구되어 메인 브랜치에 반영되지 않았다.

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

1.6 메인 저장소: eager 프런트엔드의 Python 소스 위치 보존 (#3558)

4개 파일. eager 프런트엔드가 Python의 for / with 문을 생성형 호출로 내린 뒤, 오류 및 디버깅 매핑이 생성 코드의 줄 번호로 되돌아갔다. 수정은 생성된 루프 호출을 for 문에, 컨텍스트 래핑을 with 문에 매핑하고, 이터러블 인자 언패킹 위치도 보완한다. #3125와 마찬가지로 리뷰에서 동봉된 소스 위치 회귀 테스트의 제외를 요구했다.

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

1.7 심사 중: pass 능력 검사, 부호 없는 mulhi 및 assume 술어 (10-22)

메인 저장소 심사 대기열은 모니터링 기간 내 50건이 갱신되었다. 주목할 만한 항목: [Transform] Locate VerifyBufferInit reads outside tile ops(#3620), [Codegen] Add unsigned mulhi intrinsic(#3617), [Transform] Keep assumption checks with local dependencies in scope(#3611)와 [Transform] Preserve tl.assume predicates that simplify to true(#3610), [CPU] Support clz in C and C host code generation(#3609); Ascend 측 [BugFix][Ascend] Reject odd FP4 view geometry in GM->L1 copies(#3608)와 [Ascend] Support exact placement of versioned on-chip buffers(#3607)가 같은 날 갱신되었다. 별도로 #3613은 “리뷰하지 말 것”이 명시된 네이티브 argmin/argmax 리덕션 초안이다.

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

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

2.1 Ascend: NZ 버퍼의 부분 UB→L1 복사 지원 (#3530)

당일 가장 큰 규모의 변경: 18개 파일 +993/-256. 압축 NZ 레이아웃의 패딩이 기존에는 논리 레이아웃 단계에서 평탄화되어, 부분 UB→L1 복사 시 소스와 목적지 각각의 C0 그룹 행 간격을 표현할 수 없었다. 수정은 컴팩트 NZ 패딩을 저장 정규화 단계까지 유지한 뒤 용량에 따라 정확히 한 번 구체화하고, 복사 영역이 내려갈 때 32바이트 블록 단위의 소스/목적지 C0 그룹 행 간격을 명시하며, 행 간격을 scatter 템플릿, Ascend/PTO DMA codegen 및 접근 풋프린트까지 관통시킨다. 동시에 기하·정렬·경계 검사를 추가하고, 물리 뷰 소유자를 보존하며, 호출자를 논리 할당 형상과 명시적 행 간격으로 이전했다. src/ascend/op/nd2nz.h에 83줄 추가.

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

2.2 Ascend PTO: MAD 방향 술어의 지수적 증가로 컴파일 정체 (#3581)

2개 파일 +48/-5. 중복 가변 가드를 가진 PTO 커널은 행렬 곱이 없어도 분석에 10분 이상 걸릴 수 있다. 이는 NPU 실행 이전의 컴파일 타임 정체다. 근본 원인은 PTOMmadDirectionAnalyzer가 가드가 가변 메모리를 읽을 때 진입 상태를 보수적으로 유지하는 데 있다. 두 분기가 모두 MAD 방향을 유지하면 Merge가 동일한 술어 S를 받고, 구현은 항상 S || S를 구성·간소화한다. 산술 간소화기가 재현 사례에서 이 중복을 제거하지 못하므로 연속 분기에서 술어 표현식이 배수로 불어난다. 수정은 양측 술어가 동일할 때 재구성을 건너뛰고, 20줄의 표적 테스트를 추가한다.

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

2.3 Ascend PTO: SIMT 벡터 지역 변수의 lane별 할당 (#3523)

4개 파일 +273/-29. SIMT 벡터 로컬 변수가 스칼라 PTODSL 버퍼로 할당되어 lane 수를 계산하지 않았음에도 Python 인덱스로 접근하고 있었다. AllocatedBufferValue가 항목별 할당을 지원하지 않아, 루프 인덱스 쓰기가 PTODSL AST 재작성 단계에서 더 이른 시점에 실패했다. 수정은 각 벡터 슬롯마다 충분한 스칼라 요소를 할당하고 슬롯 인덱스를 스칼라 오프셋으로 변환하며, 기존의 연속 load/store 경로를 재사용하고, AscendC의 벡터 로컬 시맨틱을 유지하며, 루프 언롤링을 바꾸지 않고 새로운 하강 pass를 추가하지 않는다. float32x2 / float32x4 두 그룹의 codegen 테스트를 추가했으며 총 209줄이다.

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

2.4 Ascend 저장소: cann-bench에 Dilation2D 추가, 같은 날 depthwise_conv_2d 롤백 (#1916 / #1919)

#1916은 Dilation2D 포워드를 병합했다(cann-bench level3 작업): NHWC, fp16/fp32, 임의 kernelSize/stride/rate 및 SAME/VALID, 경계를 벗어난 tap의 기여는 -inf + filter로 golden과 비트 단위로 일치하며, NaN 입력 전파, +inf filter tap은 경계를 벗어난 기여를 NaN으로 만든다. 순수 Vector 커널과 Python wrapper로 구현했으며, Ascend 복사 엔진의 32B 행 스트라이드 정렬 제약에 따라 디스패치하고, grid는 코어 수를 하드코딩하지 않고 기하 구조에 따라 적응한다. 같은 날 #1919는 10-21에 막 병합된 depthwise_conv_2d(#1915)를 전면 롤백했으며, 커밋 메시지에는 이유가 없다. 예제 저장소는 “하나 들어오고 하나 나가는” 상태에 있다.

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

2.5 Ascend 저장소 심사 중: sparse_mla 포워드 및 역전파 (#1929 / #1928)

#1929는 sparse_mla_fwd(DSA 희소 MLA 포워드: varlen packed 랜덤 행 gather, 온라인 softmax, QK/PV 이중 GEMM 융합)를 추가하여 examples/dsa_sparse_finetune/에 두었고, 이중 커널(fast 특화 + portable 범용)을 host 디스패치와 함께 구성했다. fast 경로의 적용 범위는 H=128 및 G=1, dim/tail은 512/64, topk는 256의 정수 배, NI는 1에서 6 사이로 고정되며, 범위 밖은 범용 경로로 간다. #1928은 sparse_mla_bwd(dQ GEMM 체인, P/dP 벡터 체인, dKV 랜덤 행 산란)를 추가하여 examples_experiment/dsa/에 두고 연산자 테스트 목록에 등록했으며, 포워드 결과 k1은 이미 이 연산자에 접혀 들어갔다. Ascend 저장소 심사 중 합계 11건.

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

2.6 Hygon: 병합 브랜치에서 enable_register_pipeline 제거 (10-22)

gcy_merge_main은 10-22 11:11에 6개 파일 +10/-24의 정리를 푸시하여 enable_register_pipeline 스위치를 전면 삭제했으며, HCU 레지스터 소프트웨어 파이프라인의 주입과 파이프라인 계획 두 pass, pipeline_utils.h 및 언어 계층의 loop 확장이 관련되며 예제와 성능 스크립트도 호출을 동기화해 변경했다. 스위치 삭제는 보통 레지스터 파이프라인 경로가 이미 기본으로 적용되어 이중 트랙 유지가 더 이상 이득이 없다는 것을 의미하지만, main 브랜치는 여전히 10-20의 preserve constant K origins in hoisted MLS windows(#25)에 머물러 있어 이 변경들은 아직 병합 브랜치에서 반영을 기다리고 있다.

출처: https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main

2.7 MetaX / Moore Threads / MLIR-Ascend: 모니터링 기간 내 푸시 없음 (10-22)

tilelang-metax의 마지막 푸시는 10-16, tilelang-musa는 09-30, tilelang-mlir-ascend는 10-19에 멈췄으며, 세 적응 라인 모두 모니터링 기간 내 커밋이 없다. Ascend, Hygon 두 라인이 이번 모니터링 기간의 모든 백엔드 변경을 담당했다.

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

2.8 산업 측: SCMP, SemiAnalysis 보고서 전재, TileLang은 6개 저장소 중 유일하게 사용 가능 (10-22)

사우스차이나모닝포스트(SCMP)가 10-22에 《Jensen’s nightmare》를 게재하며 SemiAnalysis가 목요일 발표한 화웨이-DeepSeek 보고서를 인용해 전했다. DeepSeek의 최신 소프트웨어 릴리스와 Ascend 950PR에 대한 추가 실험은 미국과의 격차가 좁혀지고 있음을 보여주며, DeepSeek가 2027년 또는 2028년까지 화웨이의 “상당히 성숙한” 기술 스택에 “주로” 의존하게 될 것이라고 판단했다. 이 보고서 원문은 지난 호에 이미 수록된 그 보고서이며, “Ascend 지원을 받는 여섯 개 저장소 중 TileLang만 사용 가능하다”는 판단은 변하지 않았고, Ascend 950 백엔드가 독립 적응 계층에 머무르지 않고 TileLang 메인 저장소에 병합되었다는 점도 언급했다. 동일 출처의 2차 전파에 해당하므로 새로운 기술적 진전을 구성하지는 않지만, 해당 판단이 주류 언론의 보도 범위로 진입하고 있음을 보여준다.

출처:
https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report

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

3.1 TileOPs: FP8 페이지드 warp 특화 kernel 공유 메모리 별칭으로 인한 NaN (#2708)

2개 파일 +80/-30. GQAPagedFP8WSKernel이 “일곱 개의 단일 token 요청이 300 token 요청 앞에 배열되고, 4099행 FP8 cache, 48행 페이지”의 조합에서 NaN을 반환한다. 근본 원인은 공유 메모리 플래너가 warp 역할별 버퍼를 서로 다른 시간대에 활성 상태로 판독하여, split handover를 여전히 TMA가 채우고 있는 key 또는 value 링 주소에 배치한 것이다. handover의 running maxima도 링 플래그 비트로 눌려버렸다. 수정은 이 버퍼들을 공유 메모리에서 명시적으로 분리하고, 24행 회귀 테스트 케이스(qwen3-fp8-ragged 기준)를 추가했다.

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

3.2 TileOPs: dense prefill을 영구 순회로 변경, D64 삼중 소비자 (#2706)

6개 파일 +590/-5, GQADensePersistentWSKernel(외부 명칭 gqa_dense_persistent, gqa_dense보다 우선)을 신규 추가: query tile 수가 SM 수를 초과할 때, 하나의 CTA가 하나의 SM을 점유하여 교대로 자신의 tile을 완료하고, 현재 tile의 epilogue 동안 다음 tile의 query를 로드한다. D64에서 각 tile은 192행 query를 세 개의 64행 소비자 warpgroup에 분배하여 K/V 재읽기가 삼분의 일 감소하고, D128은 128행을 유지한다.

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

3.3 TileOPs: SM90 dense GQA decode가 warp 특화로 전환 (#2694)

GQADecodeWSKernel 신규 추가: TMA 생산자 warp가 두 슬롯 K/V 링을 채우고, 하나의 WGMMA warpgroup이 head 그룹의 온라인 softmax를 유지한다. split 수는 상주 CTA 슬롯을 한 번에 채우는 기준으로 정하며, 각 split이 최소 네 개의 KV tile을 갖도록 보장하고, 단일 split일 때는 CTA 내에서 정규화하여 직접 출력을 쓰고 combine kernel을 거치지 않는다. KV head는 CTA 간에 가장 빠르게 변화한다.

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

3.4 TileOPs: 인과 슬라이딩 윈도우가 하나의 KV head의 두 key tile로 전체 query head 그룹을 서비스 (#2696)

GQADenseSlidingWindowBandKernel 신규 추가: GQADenseFwdOp는 기존에 각 query head마다 동일한 두 개의 key/value tile을 재로드했고, 하나의 CTA는 하나의 (query tile, head)만 처리했다. 새 kernel은 하나의 CTA가 (query tile, KV head)에 대해 band의 두 key/value tile을 한 번만 로드하고, 순차적으로 해당 그룹의 query head에 대해 리덕션을 수행하며, 동시에 다음 head의 query tile을 비행 상태로 유지한다. 슬라이딩 윈도우 행 수가 128개 key를 초과하지 않는 시나리오를 대상으로 한다.

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

3.5 TileOPs: 페이지드 decode 자동 튜닝이 이전에 하나의 block table을 공유 (#2703)

4개 파일 +24/-13. 페이지드 decode의 autotune 공급이 모든 요청에 동일한 block table을 제공하도록 하여, 결과적으로 하나의 요청 페이지만 테스트하게 되고 해당 페이지가 L2에 상주하게 되어, 후보 순위가 실제 서비스 배치가 지불한 트래픽을 반영하게 되었습니다. 1500 MHz H200에서 실측한 결과, MHA decode 튜너는 opt-1.3b-b16에서 block_N 128(콜드 80.0 us)을 선택하고 block_N 64(72.5 us)를 포기하여 flashinfer의 0.994로 떨어졌습니다. 수정은 튜닝이 요청별로 서로 다른 cache 페이지를 읽도록 하는 것입니다.

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

3.6 TileOPs: masked softmax의 FMA 오버플로와 varlen RoPE의 홀수 반폭(#2699 / #2700)

#2699(4개 파일 +54/-6)는 dense GQA 슬라이딩 윈도우 순방향의 NaN 행을 수정합니다: 인과 윈도우 아래의 128x64 튜닝 타일이 각 query tile이 첫 행 윈도우에서 시작하도록 하면서, 마지막 몇 행이 연속으로 두 개의 전면 마스킹된 key tile을 만나게 됩니다. 첫 블록 이후 두 running maxima가 모두 -1e38 클램프에 있으며, fast math가 prev * scale - curr * scale을 FMA로 융합하면서 반올림 오차(약 2의 99제곱)가 NaN으로 오버플로됩니다. 수정 방법은 스케일 재조정 전에 running maxima를 먼저 빼는 것입니다. #2700은 varlen query RoPE가 홀수 반회전 폭(rotary_dim 10, 14 등)일 때의 루프 경계를 수정합니다: 병렬도 rows 곱하기 chunks가 회전 분기의 스레드 수로 나누어떨어지지 않을 때 TileLang이 루프 상한을 잃어버려, warp 특화 경로의 후속 query 행이 잘못 계산되고 일반 경로(dim 256)는 바로 멈춰버립니다.

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

3.7 TileOPs: FP8 grouped GEMM의 단계별 스테이징 scale과 gate/up 핑퐁 warp group(#2704 / #2693)

#2704는 구성 키 stage_scales를 추가합니다: 생산자가 cp.async로 각 단계의 A/B scale을 공유 메모리에 복사하고 착지 시 해당 단계 barrier에 도달하도록 합니다; stage_scales=0은 여전히 기존 kernel을 생성합니다. 배경은 masked FP8 grouped GEMM이 수학 warpgroup에서 각 단계 K의 A scale을 전역에서 읽는데, 링이 깊어지면 L1이 압축되어 이러한 load가 멈추게 되고, 이로 인해 ds-v3-e8-up-masked가 DeepGEMM의 0.979에 그치는 것입니다. #2693은 GemmTemplate에 핑퐁 형태를 추가합니다: 두 수학 warpgroup이 각각 완전한 64행 tile을 실행하고 CTA의 tile을 번갈아 가져가며, 한 그룹의 epilogue가 다른 그룹의 mainloop 아래에 놓이고 회전 barrier로 정렬됩니다.

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

3.8 TileOPs: 술어 브로드캐스트와 불리언 피연산자의 16바이트 벡터 읽기 전환(#2702 / #2697)

#2702는 행 브로드캐스트 술어(Eq/Ne/Gt/Lt/Ge/Le와 LogicalAnd/Or)가 1024열 단위 블록의 stride-1 피연산자를 16바이트 벡터로 load16_evict_last를 통해 읽도록 하며, 일부 블록과 다른 staged body는 fragment 복사를 유지합니다; linear-bias-float32(1024x4096 대 1x4096)의 여섯 가지 비교가 0.99에서 1.06으로, 논리 연산자는 1.04로 향상되었습니다. #2697은 비교와 논리의 이항 kernel에서 전체 크기 피연산자를 16바이트 evict-first 벡터 읽기로 변경하며, float32 행은 이전에 torch.compile에서 0.99에서 1.00에 불과했습니다.

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

3.9 TileOPs: SM89 세 경로 동시 진행(#2705 / #2691 / #2690)

#2705는 SM89의 SSD chunk scan 기본 파이프라인 단계 수를 2단계로 낮춰 scan kernel의 버퍼 요구량을 줄이며, 여섯 가지 형상의 FP16/BF16 실측은 3회, 회당 200회 CUPTI 샘플링의 중앙값을 취합니다. #2691은 기존 플래너가 다중 파티션을 선택할 때 SM89의 paged DSA key 구간을 분할하고, 기존 FP32 부분 출력과 log-sum-exp의 merge를 재사용합니다; 단일 파티션일 때는 BF16 출력과 단일 kernel을 유지하고, 컴파일 대상을 merge 캐시 식별자에 포함하며 각 대상에 독립 JIT builder를 만듭니다. #2690은 SM89에서 query head 그룹 32/64, head dim 64/128의 NSA top-k 선택을 32열 tile로 전환하고, 64행 tile과 동일한 스코어링 정렬 구현을 유지하여 여덟 가지 FP16/BF16 사례의 GPU busy 중앙값을 2.1%에서 15.0%까지 낮춥니다.

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

3.10 TileOPs: 리덕션 폴드 단일 발사 및 빈 DSA 선택 검증 (#2689 / #2698)

#2689는 ReduceFoldKernel을 추가하여 각 보존 인덱스의 lead 세그먼트 연속 접미사를 한 번의 발사로 폴드한다(이전에는 축 접두사와 접미사에 대한 리덕션에 두 번의 발사가 필요했다). 또한 ReduceEdgeKernel과 VectorNormEdgeKernel이 행을 512 스레드 블록이 비행 중에 보유할 수 있는 호출을 모두 이것에 위임하도록 한다. 배경은 Sum이 3d-multidim-float16(0.977)과 mlp-intermediate-float32(0.990)에서 torch-compile에 뒤처진다는 점이다. #2698은 페이징 DSA 검증을 수정한다: query 행에 유효하게 선택된 key가 전혀 없을 때 참조 구현은 0 출력과 음의 무한대 LSE를 반환하는데, 참조 LSE를 빼면 NaN 오차 범위가 되어 올바른 결과를 오류로 판정한다. 수정은 빈 행에 정확한 0 LSE 범위를 부여하고, “모든 행이 빈” 경우 검증기가 비영 출력과 유한 LSE를 거부해야 한다는 제어 케이스를 보강한다.

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

3.11 TileOPs: 벤치마크 참조를 float32에서 계산하도록 변경 (#2692)

ref_program은 이전에 정확성 참조이자 torch 베이스라인 타이밍 대상으로 겸용되었고, 많은 참조가 16비트 입력 dtype으로 직접 torch를 호출하여 결과가 디바이스 측 torch의 각 단계 반올림 방식에 의존했다. CUDA는 이 차이를 가렸지만 CPU torch와 torch_musa는 가리지 않았다 — GroupNorm과 InstanceNorm이 MTT S5000에서 실패했으나 동일 kernel은 float64 참조에 대해서는 통과했다. 수정은 ref_program이 float32 이상에서 계산하도록 하고, torch 베이스라인은 입력 dtype에 따라 타이밍을 측정하도록 한다. 이는 국산 GPU 백엔드에 의해 드러났으나 참조 구현 계층에서 먼저 수정해야 하는 드문 사례다.

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

3.12 TileFoundry: sm_89 아키텍처 및 RTX 6000 Ada 디바이스 설치 (#259)

11개 파일 +392/-22. nvidia_sm89.toml과 nvidia_rtx_6000_ada.toml 두 가지 대상 설명(각각 114행과 124행)을 추가하고, docs/spec/target.md에 42행을 동기화하여 보강했다. 이 변경은 TileFoundry의 디바이스 모델링을 Hopper에서 Ada로 확장하며, 같은 날 #256에서 드러난 “sm_90 대상으로 컴파일된 테스트가 sm_89 단일 카드에서 값만 틀리는” 문제와 동일 사안의 양면이다. 함께 sm_89의 e4m3 warp mma 선언(#261, 심사 중)도 추가되었다.

출처: https://github.com/tile-ai/TileFoundry/pull/259

3.13 TileFoundry: gmem 계산을 풋프린트와 실행 mesh 기준으로 변경하고 피호출자 차원 바인딩 (#258)

당일 변경 범위가 가장 넓은 항목: 74개 파일 +1128/-553, #239, #252, #253 수정. 기존에는 132개 CTA가 40개만 탐색할 때 계산이 “각 창의 각 위치마다 한 번 이동”으로 산정되어 under-fill을 전혀 감지할 수 없었고, 피호출자만 선언할 수 있는 차원은 바인딩할 수 없었으며, analyze의 출력 대상과 함수별 메모리 행의 의미가 불명확했다. 수정은 check와 analyze의 차원 바인딩을 instantiate로 집중시키고, 피호출자 또는 하위 모듈이 선언한 차원을 이제 바인딩할 수 있게 한다. 커밋은 커뮤니티 기여자가 시작하고 메인테이너가 인수하여 재작성했다.

출처: https://github.com/tile-ai/TileFoundry/pull/258

3.14 TileFoundry: 디바이스가 kernel 시작을 거부할 때 반드시 오류를 발생시켜야 함 (#256)

7개 파일 +69/-11. 기존에는 런처가 kernel<<<grid, block, smem, nullptr>>>(...)를 발행한 후 시작 상태를 전혀 읽지 않아, 디바이스가 시작을 거부해도 호출이 정상 반환되고 출력은 이전 내용을 유지하여 오류가 호출자에 도달하지 못했다. sm_89 단일 카드에서 번들된 sm_90 대상으로 컴파일된 테스트는 값이 틀리는 것으로만 나타났고, compute-sanitizer를 통해서야 cudaLaunchKernel이 cudaErrorNoKernelImageForDevice를 반환함이 드러났다. 수정은 CUDA 런처가 cudaGetLastError()를 읽은 후 오류를 발생시키도록 하고, CPU 측 module도 동기화하여 변경하며, 사양 문서와 테스트도 함께 업데이트한다.

출처: https://github.com/tile-ai/TileFoundry/pull/256

3.15 배포판: Sunrise, 0.1.15+sunrise.1.2.0 릴리스 PR 생성 (#9, 심사 중)

10-22 16:48 candidate/sunrise-0.1.15-1.2.0-20261019 브랜치를 생성하고 PR을 열어, Sunrise 배포판을 0.1.15+sunrise.1.2.0으로 업데이트하고 TileLang 0.1.15 및 함께 리뷰된 TileOPs, TileKernels, TVM, tvm-ffi 변경 사항을 포함했습니다. 소스 가져오기와 병합, 배포판 적응을 별도 커밋으로 분리했으며, 최종적으로 TANG 백엔드용 Sunrise wheel을 적응 구성하고 슬라이스 GEMM 회귀를 수정했습니다(두 누산기 플래그 모두 공개 템플릿에서 검증하고, stride, origin 및 offset 단언은 유지). 자체 검증 결과: 로컬 계약 테스트 55개 통과, TANG 파이프라인 codegen 테스트 54개 통과, 두 lint job 및 신뢰 경계 검사 통과. 그러나 S2 wheel 빌드 후 설치 후 import 스모크 검사가 tangGetDeviceCount에서 실패했습니다. 브랜치의 세 커밋은 각각 0.1.15 배포판 적응, CI를 TANG 백엔드로 축소, 슬라이스 GEMM 누산기 플래그 테스트입니다.

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

3.16 채택자: TileKernels와 FlashQLA 모니터링 기간 내 푸시 없음 (10-22)

deepseek-ai/TileKernels의 마지막 푸시는 10-14에 멈춰 있고, QwenLM/FlashQLA도 마찬가지로 10-14에 멈춰 있으며, 두 채택자 모두 이번 모니터링 기간에 커밋이 없습니다.

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

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

4.1 CANN 커뮤니티, ops-tilelang 연산자 저장소 개설 (10-22)

Atlas-Ascend-Fork/ops-tilelang은 10-20에 저장소를 생성했으며, CANN 커뮤니티가 Ascend NPU를 대상으로 하는 TileLang 고성능 연산자 저장소입니다: Kernel 구현, 자동화 테스트 케이스, 표준 워크로드 및 성능 베이스라인을 통합 관리하며, 재사용 가능하고 검증 가능하며 지속적으로 발전 가능한 연산자 자산을 형성하는 것이 목표입니다. 저장소에는 이미 src, tests, scripts, docs 4개 계층과 OAT.xml, SECURITY.md 등 Open Atom 커뮤니티 규범 파일이 있으며, 10-22에 “합법적 상한 시드와 케이스 해시 조합 후 오버플로로 CPU 테스트 실패”를 수정하는 커밋을 푸시했습니다; 더 이전 커밋에는 “Ascend A5 참조를 Ascend 950으로 교체”와 “benchmark 실패 또는 중단된 세션도 성능 베이스라인을 덮어쓰는 문제”가 있습니다. 단일 DSL을 유일한 구현 언어로 하는 커뮤니티급 연산자 저장소는 이전에 없었으며, 이 라인은 주간 단위로 추적할 가치가 있습니다.

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

4.2 tilelang/agent: 연산자 최적화를 agent의 SKILL 워크플로로 작성 (10-22)

tilelang/agent(저장소 생성 08-28, 3 스타)는 10-22에 푸시했으며, README 자체가 실행 가능한 prompt 템플릿입니다: @tilelang-operator-dev/SKILL.md로 도메인 지식을 주입하고, 입력은 연산자 이름, GPU/CPU 참조 구현, shape 및 “tilelang npu 납품, 중국어 납품” 제약뿐입니다; 주석에는 agent가 A5 환경에서 실행되어야 하며, 로드 후 10단계 폐쇄 루프를 따라야 한다고 명시되어 있고, Skill에는 선택 결정 트리, 8-CSV 규칙, L0에서 L4 전략 및 11개 함정이 포함되어 있습니다. 4.1의 연산자 저장소, 3.15의 명령어 수준 스케줄링과 이 저장소를 함께 보면, Ascend 측에서는 “사람이 규칙을 작성하고, agent가 튜닝하며, 저장소에 베이스라인을 축적하는” 분업이 형성되고 있습니다.

출처: https://github.com/tilelang/agent

4.3 커뮤니티 신규 저장소: Sail 래퍼, 원격 자동 튜닝 agent 및 튜토리얼 저장소 (10-22)

flytiger-eco/tilelang-wrapper-for-sail은 10-21에 저장소를 생성하고 10-22에 푸시했습니다(Python, 설명 없음); cc-c122/TileLang-Remote-Autotuning-Agent(Python, 10-22 푸시)와 easy-tilelang/tilelang-devkit, easy-tilelang/tilelang-tutorials는 같은 날 푸시되었습니다. 수가 많고 개별 저장소의 스타 수가 낮아, 튜토리얼 및 도구 계층의 자연스러운 성장에 해당하며 아직 판단할 수 있는 기술 방향은 없습니다.

출처: https://github.com/flytiger-eco/tilelang-wrapper-for-sail

4.4 문서 및 버전: 문서 사이트 세 번 업데이트, VERSION은 여전히 0.1.15 (10-22)

tilelang.github.io는 Update docs를 세 번, TileFoundry.github.io는 render TileFoundry v0.0.3 한 건, TileOPs.github.io는 한 건을 기록했습니다. 메인 저장소의 VERSION은 0.1.15에 멈춰 있고, 최근 릴리스는 여전히 09-30의 v0.1.15입니다; 저장소는 8742 스타, 915 fork, 열린 issue 389건입니다.

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

5. 동향 관찰

첫째, 메인 저장소의 결함 유형이 한 층 바뀌었다. 지난주 메인 저장소에서 정리한 것은 “형상과 기하 검증”이었다. cast 반올림 주석이 조용히 버려지고, arch 없는 CUDA target이 조용히 sm_50으로 후퇴하고, 희소 메타데이터 형상을 검증하지 않던 문제였다. 이번 모니터링 기간의 9건 중 4건은 분석기와 소스 위치 같은 컴파일 시점 상태에 속한다. analyzer가 재작성 후 재사용되고, 분할 후 변수가 재바인딩되고, eager 프런트엔드가 Python 위치를 잃는 문제다. 이런 결함의 공통점은 “생성된 kernel은 여전히 정확할 수 있지만, 컴파일러가 당신에게 말하는 것과 실제로 하는 것이 불일치한다”는 것이다. TileLang이 agent 자동 kernel 작성(4.2)을 접목하고 있는 만큼, 컴파일 시점 상태 일관성이 agent 반복 폐루프의 수렴 여부를 직접 결정하게 된다.

둘째, Ascend 측 진전이 연산자 수량에서 저수준 데이터 경로로 옮겨갔다. #3530은 18개 파일에 걸쳐 NZ 압축 레이아웃의 일부 UB→L1 복사에서의 행 간격 표현을 다룬다. #3581과 #3523은 각각 컴파일 시점 프레디커트 폭발과 SIMT 벡터 지역 변수의 할당 의미를 수정한다. 세 건 모두 연산자를 새로 추가하지 않고, 기존 경로가 더 많은 형상에서 컴파일되고 정확히 실행되도록 만든다. CANN 커뮤니티가 별도로 ops-tilelang 연산자 저장소를 새로 연 것(4.1)과 함께 읽으면 분업이 보인다. 연산자 자산은 커뮤니티 저장소가 담당하고, 메인 저장소는 백엔드 의미와 codegen에 집중한다.

셋째, TileOPs는 이틀 연속 20건을 처리하며 작업 중심이 SM90 특화와 SM89 보완에 있다. 모니터링 기간에 SM90 warp 특화 kernel 세 개(dense decode, 인과 슬라이딩 윈도 band, FP8 varlen 오버랩)를 새로 추가했고, 동시에 SM89의 SSD chunk scan, paged DSA, NSA top-k 세 가지 소규모 경로를 함께 추진했다. 주목할 점은 두 건의 NaN 수정(#2708, #2699)이 모두 warp 특화 경로에서 나타났다는 것이다. 공유 메모리 플래너의 활성 기간 판단과 fast math 하의 FMA 융합은 이런 수작업 편성 kernel에서 가장 오류가 나기 쉬운 두 지점이며, 바로 이것이 TileLang이 사용자를 대신해 관리하겠다고 약속한 부분이다.

넷째, 하나의 역방향 지표는 벤치마크 계층이다. #2692는 참조 구현을 float32 연산으로 바꿨는데, 발단은 MTT S5000에서 GroupNorm과 InstanceNorm이 실패한 것이었다. 문제는 kernel 자체가 아니라 참조 구현이 입력 dtype을 따라간 데 있었다. 같은 날 TileFoundry #256은 “장치가 시작을 거부함”을 조용한 처리에서 예외 발생으로 바꿨다. 두 건 모두 같은 방향이다. “통과한 것처럼 보임”을 “실제로 통과함”으로 바꾸는 것이다. 이 도구들에 의존해 칩 적응을 하는 팀에게 이런 수정의 실질 가치는 몇 가지 성능 수치보다 높다.

다섯째, 배포와 산업 두 축이 가속하고 있다. Sunrise 배포판은 0.1.15와 TANG 백엔드 wheel을 묶어 릴리스 PR로 올렸다. 이는 TileLang에 이미 하나 이상의 다운스트림 배포 형태가 나타났음을 뜻한다. 산업 측에서는 SCMP가 SemiAnalysis의 판단을 전재하며 “여섯 저장소 중 TileLang만 사용 가능”을 주류 매체로 밀어냈다. 기술 리듬상으로는 여전히 조용한 하루였다(메인 저장소 릴리스 없음, breaking change 없음). 그러나 생태계 계층의 조직적 움직임이 코드 계층보다 확연히 많았다.

부록: 자료와 검증 설명

출처 검증표

출처 검증 결과
tile-ai 조직 저장소 29개 저장소, 모니터링 기간 내 10개 푸시, 35건 병합, 71건 심사 중 업데이트
tilelang 메인 저장소 9건 병합, 1건 단일 행 수정, 1건 18개 파일 Ascend 변경, VERSION 0.1.15, 신규 릴리스 없음
TileOPs 20건 병합, 전부 Attention / MoE / Elementwise / Reduction의 kernel과 테스트
TileFoundry 3건 병합, 별도 4건 심사 중, 0.0.3의 10-21 릴리스 설명은 지난 호에 수록됨
tilelang-ascend 2건 병합(1건 예제 추가, 1건 예제 롤백), 11건 심사 중
tilelang-hygon main 커밋 없음, 병합 브랜치 gcy_merge_main에 스위치 정리 1건
tilelang-sunrise candidate 브랜치 신규 생성 후 0.1.15+sunrise.1.2.0 릴리스 PR 개설, 미병합
tilelang-metax / musa / mlir-ascend 모니터링 기간 내 푸시 없음, 각각 10-16 / 09-30 / 10-19에 정지
TileRT / tilescale / DeepStack 모니터링 기간 내 푸시 없음
채택 측 TileKernels / FlashQLA 모두 10-14에 정지, 푸시 없음
Google News(중영문 여러 조합, 프록시 경유) 3건 적중, 2건은 동일한 SemiAnalysis 보고서의 매체 전재, 1건은 DeepSeek 자금 조달로 주제 무관하여 제외
Hacker News TileLang 관련 논의 없음
arXiv TileLang을 제목 또는 초록 키워드로 검색, 모니터링 기간 내 0건
커뮤니티 신규 저장소 신규 또는 활발한 저장소 5개, 모두 개설일과 푸시 시간 확인 완료
회원 또는 벤더 공식 사이트 본 주제에 해당 없음

전체 출처 목록

cale(#2704) — https://github.com/tile-ai/TileOPs/pull/2704