보고 기간: 2026년 11월 2일(월)부터 11월 6일(금)까지, 총 5영업일 자료 출처: 본지의 일일 TileLang 커뮤니티 동향 데일리 리포트 5호(11-02, 11-03, 11-04, 11-05, 11-06)로, 5개 호의 수집 기간은 서로 이어져 있으며 합계 10-30 07:00부터 11-06 07:00까지 총 168시간(월요일 호는 72시간 기간으로 직전 주말을 포함하고, 화요일부터 금요일까지는 24시간)을 포괄하며, 이전 위클리 리포트(10-26~10-30)와 중복되지 않음. 또한 tile-ai 조직의 각 저장소 168시간 기간 보완 점검, 각 저장소 릴리스 페이지 및 태그, TileLang×RISC-V 4개 출처 순회 재검토(상세 내용은 부록 참조)

1. 주간 주요 뉴스

  • 승등 950의 영속화 FlashMLA 커널이 별도 저장소 이식이 아닌 컴파일러가 스케줄링 가능한 방식으로 메인라인에 직접 반영됨(11-02 병합, #3678): 1997줄 추가, 17줄 삭제, 20개 파일, 42개 커밋으로 BF16 prefill과 컴팩트 FP8 decode 두 경로를 포괄하며, 선택적 FP8 및 FP4 추가 KV 캐시, attention sink, 불법 인덱스와 query별 독립 길이를 지원함. 어텐션 루프, 온라인 softmax, 역양자화 및 스케줄링이 전부 TileLang으로 표현되며, T.Stage를 통해 컴파일러가 버퍼 버전과 동기화를 관리하도록 함. QK와 PV는 L0B 저장 링의 서로 다른 두 layout 뷰를 공유하고, L1 배치는 T.annotate_buffer_addresses로 직접 지정함. prefill은 하나의 누산기 할당에서 네 개의 서로소 출력 슬라이스를 나누고, 정규화가 한 조각 끝나면 즉시 한 조각을 비움. 이 방식을 성립시키기 위해 컴파일러 측에서 같은 시기에 세 가지 기능을 보완함: L0C에서 UB로의 복사가 NZ 출력을 유지, 각 AIV 스트림이 자체 파이프라인 이전에 선택됨, 명시적 스케줄링 단계가 물리적 발사 위상이 동일할 때 소스 순서를 유지. 목표 형상은 매우 좁게 작성됨: 헤드 수 64, Dqk와 Dv 모두 512, topk는 64의 배수. 출처: https://github.com/tile-ai/tilelang/pull/3678

  • ThreadSync 일주일 세 번 수정, 먼저 마이그레이션을 위해 양보한 뒤 성능과 증명을 복원(11-01, 11-04, 11-05): 월요일 #3733은 3주 전 병합된 정밀 스레드 카운트(#3291)를 전면 롤백했으며, 그 이유는 매우 직설적으로 작성됨—그 변경은 값비싼 카운트 조회를 가득 찬 블록 펜스 검사 이전으로 옮겨, 모든 스레드가 펜스에 도달하더라도 증명을 실행해야 했고, TVM의 솔버 예산이 고갈된 후 부분 카운트를 반환하여 정당한 remap과 sequence 보조 kernel이 곧바로 “Cannot prove a thread rectangle”로 컴파일 실패했음. 상수 가드, 정밀 참여 스레드 경계, 비직사각형 스레드 집합 세 가지 정확성 개선이 함께 철회되었고, #3732는 마이그레이션 후 재최적화를 위한 추적 티켓으로 남음. 목요일 #3761은 롤백에서 파생된 부작용을 처리함: #3643의 명명된 배리어 중복 번호를 수정하기 위해 동일한 스레드 그룹이 스코프를 넘어서도 ID를 재사용할 수 없게 되었고, TileKernels의 mhc_pre_big_fuse 커널은 블록당 96 스레드이기 때문에 2단계 cp.async 파이프라인에서 ID를 하나 더 받게 되어, 컴파일 산출물이 선언하는 배리어 리소스가 4에서 5로 변경되고, B300에서 히든 차원 1280의 배리어 점유 상한이 16블록/SM에서 12블록/SM으로 떨어짐. 수정 방법은 ThreadSync가 정렬된 스코프 목록을 받아들이도록 하고, 스코프를 넘어 ID를 공유하기 전에 새 분석기로 참여 술어가 항상 성립함을 증명하는 것이며, 수용 기준은 리소스 선언이 4로 복귀하고 점유가 16블록/SM으로 복귀하는 것이고, #3643의 수정은 롤백되지 않음. 금요일 심사 중인 #3787은 증명기의 사각지대를 계속 보완함: 192 스레드에서 tx < 160 and tx // 32 % 2 == 0과 같은 “경계 항과 warp 인덱스 모듈로 항을 and로 결합한” 형태가 오탐을 발생시키며, 근본 원인은 FloorDiv와 FloorMod의 인코딩이 “상한 더하기 모듈로” 조합에서 두 나눗셈의 관계를 확립하지 못하는 것이고, 수정 방법은 threadIdx.x를 warp_index * warp_size + lane으로 다시 쓰는 것임. 출처: https://github.com/tile-ai/tilelang/pull/3733, https://github.com/tile-ai/tilelang/pull/3761, https://github.com/tile-ai/tilelang/pull/3787

  • TIRx 마이그레이션이 ‘구두 근거’에서 일주일 만에 세 가지 확실한 증거로 발전했고, 기반 교체가 공개 준비 단계에 진입했다 (11-03 ~ 11-06): 화요일 #3736는 tilelang/language/ast/, tilelang/language/parser/ 그리고 tilelang/language/tir/entry.py를 함께 삭제했다. 이는 초기 로컬 TVMScript 파서 분기의 잔재로, 1줄 추가, 3384줄 삭제, 10개 파일에 해당한다. 언어 계층을 업스트림 tvm.tirx.script.parser로 변경하고, AST와 명령문 오버라이드 등록을 tilelang.language.overrides에 두었다. 목요일 #3752가 의존하는 tile-ai/tvm#91과 금요일의 #3752는 Z3의 rlimit을 하드코딩된 100,000에서 kernel별로 설정 가능하도록 변경했는데, 작성자의 표현에 따르면 “이로 인해 Z3 증명기와 관련된 많은 문제가 간접적으로 해결되었다”. 금요일 #3803은 전체 마이그레이션 초안(draft)을 제출했다. vendored된 TVM을 독립적인 Apache TIRx로 교체하고, 프런트엔드, IR 변환, 코드 생성, JIT 어댑터 및 패키징을 모두 네이티브 객체 모델로 이전했으며, apache/tvm-tirx의 7e49a148과 apache/tvm-ffi의 3e906bc6를 pin했다. 의존성 트리는 패치를 주입하지 않고 깨끗한 업스트림을 유지한다. 세 가지 흐름이 동일한 지점을 가리킨다. 메인라인의 이 단계 최우선 순위는 기반 교체이며, 나머지 최적화는 모두 이것에 자리를 양보해야 한다. 출처: https://github.com/tile-ai/tilelang/pull/3736, https://github.com/tile-ai/tilelang/pull/3752, https://github.com/tile-ai/tilelang/pull/3803

  • ‘조용히 잘못된 코드 생성’이 이번 주 가장 주요한 결함 유형이 되었으며, 네 개 릴리스 버전에 걸쳐 나타났다 (11-03 ~ 11-05): 가장 심각한 항목은 #3808이다. 조건 관련 스칼라화 이후 공유 메모리 epilogue에서 __syncthreads()가 하나 누락되어, 0.1.12, 0.1.13, 0.1.14, 0.1.15 네 버전 모두 컴파일 시 오류를 전혀 보고하지 않으면서 잘못된 CUDA를 조용히 생성한다. 보고자는 완전히 실행 가능한 재현 스크립트를 제공했다. 같은 계열로 #3800(warp specialization을 끈 후 T.Pipelined가 tcgen05의 barrier 대기 전에 다음 stage를 재로드하여, MMA가 다시 쓰인 operand를 읽는 문제), #3721(sm_100a에서 TMEM에 대한 T.gemm 시 threads/32가 N을 나누지 못하면 PTX에 tcgen05.mma가 하나도 없고 곱이 조용히 버려짐), #3775(프로덕션 환경 TP2 sglang에 nvrtc 백엔드 적용 시 CUDA graph 캡처가 손상되어, 바이트 단위로 동일한 설정을 세 번 재현했을 때 서로 다른 첫 실패 지점이 나타남), #1991과 #1992(승등 910_9362에서 threads=2일 때 AscendVidReduction이 주입된 tmp workspace도 무조건 절반으로 줄여, 간접 주소 지정 tile에 적용된 sin/cos가 3978/4096개 원소에서 잘못 계산됨, 즉 97.1%)가 있다. 이 계열은 모두 오류를 보고하지 않고, 일반 테스트를 통과하며, 오직 표적화된 재현으로만 포착할 수 있다. 출처: https://github.com/tile-ai/tilelang/issues/3808, https://github.com/tile-ai/tilelang/issues/3800, https://github.com/tile-ai/tilelang-ascend/issues/1991

  • TileFoundry, 동기화를 「분석으로 증명 가능」에서 「발행은 반드시 증명과 일치해야 함」으로 진전 (11-02 ~ 11-06): #283은 「공유 메모리 접근을 어떻게 정렬해야 하는가」를 증명 가능한 분석 계층으로 구현했습니다. lowering의 발행 순서에 따라 각 함수를 순회하며, 충돌 접근 이전에 commit, wait 및 스코프 내 배리어를 계획합니다. 83개의 스케줄링 픽스처 중 52개를 계획하고 31개를 거부했으며(24개는 TMA 복사, 6개는 wgmma 사용, 두 방식의 완료 시맨틱은 아직 선언되지 않음), 기본적으로 호출자 요청이 없기 때문에 finalize되는 75개 픽스처의 바이트는 main과 완전히 동일합니다. 목요일에 머지된 #285는 그 실행 측입니다. CUDA lowering이 동기화 계획에 따라 발행하고 항목별로 역검증하며, 불일치 시 LoweringError를 발생시킵니다. 단일 sm_89에서 빌드 가능한 7개 프로그램 중 6개가 참조와 완전히 일치했고 synccheck와 memcheck 모두 깨끗했으며, 동일한 7개를 main 방식으로 lower하면 4개가 결과 오류를 보였고 직접적 증상은 sm_89의 E4M3 GEMM 최대 편차 659였습니다. 금요일에 머지된 #286은 WGMMA의 비동기 완료 시맨틱을 보완하고, 수요일에 머지된 #284는 FSDP와 HSDP의 학습 시맨틱을 HIR 평가기에 구현하여 4카드 sm90에서 실제 2×2 PyTorch HSDP와 정렬했습니다. 작성자는 #286에 대해 명확한 한계를 유지했습니다. 행렬 명령 발행기가 아직 연결되지 않았으며, 이 커밋은 실행 정확성을 증명하지도 않고 성능을 제공하지도 않습니다. 출처: https://github.com/tile-ai/TileFoundry/pull/283, https://github.com/tile-ai/TileFoundry/pull/285, https://github.com/tile-ai/TileFoundry/pull/286

  • TileOPs, 5개 모니터링 기간에 15, 7, 1, 7, 9건 머지 — 경쟁점이 「연산자 수」에서 「선택기와 준비 비용」으로 완전히 이동 (전주): 월요일 모니터링 기간의 15건 중 13건이 11-01 하루에 집중되었고, 거의 모든 건에 H200 실측 표와 「manifest에 측정된 행이 없으면 폴백」 선언이 포함되었습니다. 목요일 7건, 금요일 9건은 리듬을 연휴 이후 가장 밀집한 상태로 되돌렸습니다. 세 가지 핵심 수치가 이번 주 방향을 요약합니다. GQA dense forward는 하나의 실측 규칙으로 decode 또는 prefill 프로그램을 선택하며, 최고 속도 후보 대비 선택 후회가 1.818에서 1.017로 감소(#2808); warp 전용 decode는 softmax와 QK, PV를 중첩하여 영향을 받는 478회 스캔 호출의 기하 평균 지연이 main의 0.952배(#2817); 브로드캐스트 피연산자 인플레이스 읽기로 lerp, where, clamp, masked_fill 5류 연산자의 브로드캐스트 행 소요 시간이 메인라인의 0.13~0.44배로 감소(#2832). 동시에 roofline 가격 책정 재작성(물리적으로 불가능한 세 개의 대역폭 행을 7.48, 7.50, 5.30 TB/s에서 0.62, 0.62, 2.65 TB/s로 수정, #2830)과 3단계 타일 인자 개별 피팅(선택 후회 0.89%에서 0.06%로 감소, #2836)도 있었습니다. 출처: https://github.com/tile-ai/TileOPs/pull/2808, https://github.com/tile-ai/TileOPs/pull/2832, https://github.com/tile-ai/TileOPs/pull/2830

  • 국산 백엔드, 「자체 개발 브랜치」에서 「공개 리뷰」로 — Hygon이 처음으로 공식 PR 개설 (11-05): Hygon 저장소의 메인 브랜치는 11-02의 #32(gfx948을 gfx946 호환으로 처리)에 멈춰 있었고, 이번 주 움직임은 먼저 feat/hcu-scaled-cast 브랜치에 네 건을 쌓았습니다(E8M0 스케일의 FP8/BF8/FP4 변환, 64 bank LDS 읽기 위상, 혼합 LDS wrap의 최대 2-way 증명, 물리적 읽기 폭에 따른 혼합 검사). 금요일에는 LDS 위상 수정 세 건을 추가로 반영하고, 09:08과 19:52에 각각 #33(scaled cast)과 #34(reduction의 XOR shuffle에 DPP와 lane swap 사용) 두 PR을 열었습니다. 이는 국산 백엔드 중 「자체 개발 브랜치에서 공개 리뷰까지」 이 경로를 처음 완주한 사례입니다. mlir-ascend는 같은 기간 BishengIR 툴체인 어댑터를 contrib로 옮겼고(#209), #203, #211, #212 세 개의 작은 PR로 tl.sub/mul/div/max/min, tl.shl, tl.shr을 하나씩 TLIR 방언으로 이전했습니다. Moore Threads는 11-05에 README 문서 커밋 한 건을 반영하여 09-30 이후의 완전한 침묵을 끝냈습니다. 출처: https://github.com/tile-ai/tilelang-hygon/pull/33, https://github.com/tile-ai/tilelang-hygon/pull/34, https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

  • 하드웨어 지원이 양쪽으로 동시에 확장: Blackwell 신규 SKU와 국산 가속기 (전 주간): 한쪽은 새 카드입니다. B300에서 보고된 sm_103a는 두 경로에서 각각 한 번씩 막힙니다. tcgen05.mma.kind::i8은 sm_100a와 sm_110a에만 존재하며, 커널이 lowering을 거친 후 ptxas에서 Feature '.kind::i8' not supported 오류가 발생합니다(#3770). CuTeDSL의 TMA 보조 함수는 아키텍처 화이트리스트를 유지하는데, B300은 DSL에 진입하기 전에 거부됩니다(#3772, sm_103a와 gr100의 sm_107a를 추가한 후 B300에서 8 통과 1 스킵, gr100에서 13 통과 1 스킵). SM120에서 FP8에 fp32 누산을 적용한 GEMM은 계속 SM89의 mma.sync 원자를 사용했는데, ue8m0 단위 스케일링이 있는 block-scaled MMA로 교체한 후 RTX 5090에서 8192 큐브가 463에서 727 TFLOPS로 상승했습니다(심사 중 #3790). 다른 한쪽은 국산 카드입니다. Hygon은 64 bank LDS 위상에서 마이크로아키텍처 수준 튜닝을 수행했고, Ascend는 Cube와 Vector의 코어 비율을 고정 1:2에서 구성 가능하도록 변경했으며(#3783), AutoSchedule의 주기 하한을 조여 Z3 해석 시간을 10.04% 단축했습니다(#3776). 출처: https://github.com/tile-ai/tilelang/pull/3772, https://github.com/tile-ai/tilelang/pull/3790, https://github.com/tile-ai/tilelang/pull/3783

2. 핵심 프로젝트 진행 상황(메인 저장소 tilelang)

버전 및 릴리스

  • 메인 저장소의 최신 태그와 릴리스는 여전히 v0.1.15(09-30 릴리스)이며, 전 주간 신규 태그와 신규 릴리스 페이지 항목이 없어 이번 주 금요일 기준 37일째입니다.
  • 조직 내 두 신규 버전은 모두 지난주에 멈춰 있습니다. TileOPs는 v0.1.0(10-27), TileFoundry는 0.0.4rc1(10-25)이며, 이번 주 모두 정식 전환되지 않았습니다.
  • 적응 라인 태그는 전 주간 변동이 없습니다. tilelang-ascend v0.1.2.000(09-09), tilelang-metax v0.1.15(10-16), tilelang-musa v0.1.15+musa.1(09-30), tilelang-mlir-ascend는 v0.1.2.020 계열에 머물러 있습니다.
  • 머지량과 리듬: 월요일 모니터링 기간은 72시간(10-30 07:00 ~ 11-02 07:00)으로 35건 머지되었습니다. 화요일부터 금요일까지 네 개의 24시간 기간은 각각 18, 8, 12, 18건으로 다섯 기간 합계 91건입니다. 심사 중 업데이트는 순서대로 82, 44, 48, 42, 57건으로 합계 273건이며, 신규 issue는 순서대로 11, 8, 7, 4, 5건으로 합계 35건입니다. 메인 브랜치 커밋 분포: tilelang 7, 5, 5, 3, 5건, TileOPs 15, 7, 1, 7, 9건, TileFoundry 3, 1, 1, 1, 1건, tilelang-ascend 5, 4, 0, 1, 2건, tilelang-mlir-ascend 1, 0, 1, 2, 0건, tilelang-hygon 메인 브랜치 2, 1, 0, 0, 0건(별도 기능 브랜치 커밋 있음), tilelang-lsp와 VS Code 확장 각 1건, tilelang-musa 1건입니다.
  • 리듬 측면에서 이번 주는 ‘주 초 정산, 주 중 저점, 주말 반등’의 형태를 보였습니다. 수요일은 8건 머지로 이번 주 최저점이었고, 금요일에는 18건으로 회복했습니다. 심사 대기열은 주 초 82건에서 42건까지 내려간 후 금요일에 57건으로 반등하여, 주 초에 적체된 리뷰가 이번 주에 대량 소화되었음을 보여줍니다.

컴파일러, 언어 계층 및 실행 경로

  • Ascend 950 FlashMLA 커널 및 이에 대응하는 컴파일러 기능(#3678, 11-02 병합, 20개 파일 +1997/-17): 중요 소식 첫 번째 항목을 참고. 주목할 점은 이번 변경이 ‘명시적 스케줄링’에 필요한 세 가지 저수준 기능을 동시에 메인라인에 보강한 것이지, Ascend 브랜치에서 패치한 것이 아니라는 점이다 — L0C에서 UB로의 복사는 NZ 출력을 유지하고, 각 AIV 스트림은 파이프라인 이전에 선택되며, 물리적 발사 위상이 같을 때 소스 순서를 유지한다.
  • cast 분리형 벡터 폭 회귀 연속 수정 두 건(#3734 11-02 병합, #3759 11-05 병합): #3734는 SM100에서 ViT 포워드의 Cannot convert type float32x16 to CUDA type을 수정했다 — 벡터화 플래너가 256비트 BF16 저장을 위해 16개 lane을 선택하고 cast 제약을 지연시키는데, #3579가 분리형으로 메모리 읽기 가드를 건너뛰게 한 이후 최종 합법화가 여전히 지연 전략을 따르면서 잔여 FP32 표현식이 512비트 벡터로 확장되었다. 수정 방법은 defer_cast_constraints를 IR 속성이 아닌 플래너의 C++ 파라미터로 만드는 것이며, 가드의 FP32→BF16 변환은 8개 lane으로 줄이고 성공적으로 분리된 BF16 복사는 여전히 16개를 유지한다. #3759는 그 뒷정리다: 조건 내 int64 cast가 이전에는 데이터 벡터 폭을 16 lane에서 4 lane으로 압축했는데, 해당 조건이 16개 lane에서 완전히 일치하는 경우에도 그러했다. 수정 방법은 조건 일관성 분석을 cast 폭 면제에 재사용하는 것이며, if 문, tirx.if_then_else와 스칼라 술어의 Select 조건은 별도로 접근하고, 서브절 내의 cast는 여전히 제한된다.
  • ThreadSync의 3단계 진화(#3733, #3761, #3787): 중요 소식 두 번째 항목을 참고. #3733은 명확한 임시 조치로, 원문에 “TIRX 마이그레이션을 위한 길 비우기”라고 명시되어 있으며, 대응하는 TVM fork 최적화 제안 tile-ai/tvm#89는 이미 닫혔다. #3761의 401줄 변경에서 가장 기록할 만한 것은 귀인 실험이다: 저자는 컴파일 산출물에 선언된 배리어 리소스 수만 단독으로 변경하고 실행 가능한 명령은 전혀 건드리지 않은 채 동일한 성능 저하를 재현하여, “성능 회귀”를 느낌에서 지목 가능한 단일 변수로 바꾸었다.
  • 언어 계층과 프런트엔드 감량: #3736은 3384줄의 레거시 TVMScript 파서 분기를 삭제했으며(중요 소식 세 번째 항목 참고), 이를 유지할 위험이 매우 구체적으로 기술되어 있다 — 누군가 tilelang.language.parser.operation을 직접 import하면 tirx.PrimExpr과 tirx.IterVar에 서로 충돌하는 연산자를 등록한다; #3760은 중첩 T.macro에 경계가 있는 프로세스 내 변환 캐시를 도입하여 컨텍스트와 무관한 변환 결과만 저장하고, 매크로 인스턴스는 현재 globals와 클로저 값을 다시 바인딩하여 “kernel은 이미 캐시에 적중했지만 매크로는 여전히 반복 파싱·컴파일”되는 문제를 해결한다; #3802는 T.annotate_host_mapped_buffers를 신설하여 주석된 파라미터의 인덱스를 PrimFunc의 tl.host_mapped_params 속성에 기록하고, 런타임에서 cuMemHostGetFlags로 등록 상태를 검증하며 cuMemHostGetDevicePointer로 디바이스 가시 주소를 해석하고, 일반 페이지 가능 메모리는 여전히 거부된다.
  • 디버깅과 관측성: #3703은 #line 지시자를 “각 TIR 문장이 첫 줄만 앵커링”에서 단계별 발사(0은 끔, 1은 문장 수준, 2 이상은 줄별 재앵커링)로 변경했으며, 구현은 CUDA, HIP, CPU C, C host, Metal, Ascend가 공유하는 C 계열 출력 스트림에 들어가고, Python 측 기본값으로 TILELANG_EMIT_LINE_DIRECTIVES를 신설했다; 저자는 CUDA와 HIP 사용 사례가 로컬에서 건너뛰어졌고 cuda-gdb의 실제 동작은 직접 검증되지 않았다고 명시했다. 같은 날 LSP 저장소가 한 건 따라가며 CUDA 디버깅을 새로운 줄 매핑에 정렬했다; 목요일 VS Code 확장 저장소가 다시 한 건 따라가며 문서를 업데이트하고 산출물 이름을 vsix로 변경했다.
  • 자동 튜닝의 기준: #3546은 set_autotune_inputs의 공급 클로저를 “호출 지점에서 절대 전달되지 않는 파라미터”에서 get_current_device()로 변경했는데, 이전에는 CUDA가 없는 머신에서 폴백 분기가 모든 캡처 텐서를 CPU로 옮겨 매 튜닝 라운드가 런처 디바이스 불일치로 실패했다; 듀얼 카드 Ascend950에서 세 건의 Ascend 자동 튜닝 회귀는 수정 전 전부 실패, 수정 후 전부 통과했다. 동일 계열의 #37

84 AscendC CI에서 디바이스 간 복사 프로브 앞뒤에 각각 한 번씩 디바이스 동기화를 추가하여 비동기 오류가 try ... except RuntimeError의 범위 안으로 되돌아오도록 했다. 이전에는 이 오류가 한 박자 늦게 드러나서 마치 무관한 torch.arange()가 4바이트 memcpy 크기 불일치를 보고한 것처럼 보였다.

  • Z3 해석 리소스가 하드코딩에서 조정 가능으로 변경(#3752, 11-06 머지, tile-ai/tvm#91에 의존): rlimit의 기본값은 여전히 100,000이며, 설정하지 않으면 동작이 변하지 않는다. 이는 목요일 Ascend 측에서 AutoSchedule 주기 하한을 강화한 것(#3776, Z3 해석 시간 10.04% 감소)과 같은 방향의 양 끝단이다. 하나는 검색 예산을 바꾸고, 다른 하나는 검색 시작점을 바꾼다.

정확성 및 결함 수정

  • ThreadSync 및 배리어: #3733의 전체 롤백(요약 참조)과 #3761의 크로스 도메인 재사용(요약 참조) 외에, 심사 중인 #3765는 “단일 스레드 인덱스 경계 가드가 있는 부분 배리어”를 스레드 차원이 워프 전체 수일 때 명시적으로 허용하고, #3767은 다중 인덱스 fragment 접근의 StructuralEqual 내부 단언을 두 인덱스 표현식이 포함된 ValueError로 교체하며, #3786은 T.copy 목적 영역 범위 초과를 컴파일 타임 오류로 변경하고, #3788은 ProveFragmentContains의 빠른 경로를 수정한다——이 경로는 큰 fragment가 완전히 복사될 때 스레드 소속 검사를 건너뛰고 바로 참을 반환하여, “작은 fragment가 스레드 0부터 127까지 소유하고 큰 fragment는 0부터 95까지만 소유”하는 경우를 포함이 성립하는 것으로 오판한다.
  • 조용히 잘못된 코드 계열: #3808, #3800, #3721, #3775, #1991/#1992(요약 네 번째 항목 참조) 외에, 이번 주에는 #3774(16×16 공유 메모리 치환에서 행 이동만 하고 열 오프셋을 놓쳐 256개 요소 중 223개가 틀렸으며, 같은 방식의 32×32는 정상), #3773(균일 while 루프에서 세 번의 공유 메모리 복사와 GEMM 수행, 직렬 대조군은 100회 중 0회 오류, while 버전은 100회 중 21회 오류에 10가지 서로 다른 출력 발생), #3763(중첩 정수 내림 나눗셈 (x // -2) // 4가 int16, 64개 요소, 32 스레드에서 Let cannot bind the same var to two different values 오류 발생)도 있다. 이 세 항목과 #3775는 “소비자용 카드 + WSL + pip 설치” 조합의 집중적인 오류 보고 면을 구성한다.
  • 의존성 분석과 동기화 삽입의 근본 원인: #3749는 이번 주 기술적 깊이가 가장 높은 결함 보고서이다——두 번의 연속 SIMD 로드가 더 이른 GM에서 UB로의 복사에 대한 의존성을 잃었는데, 근본 원인은 MemoryAccessDetector::LogicalPointerRegion이 재구성한 논리 영역이 물리적 로드 범위의 일부를 누락한 데 있다. 최소 재현에서 work.shape는 (2, 128)이고, 생성된 읽기 디스크립터는 물리적 [64, 192)를 커버하지만 재구성된 영역은 [0, 128)이어서, 의존성 분석이 입력 DMA와 로드를 무관하다고 판단하고 AutoSchedule이 입력 복사를 로드 이후로 옮기며 InsertSync도 MTE2에서 V로의 set/wait를 발행하지 않는다. 같은 모니터링 기간에 메인 저장소에는 AscendC와 PTO 두 백엔드의 코드 생성 및 동기화 삽입을 지적하는 issue 8건(#3740부터 #3749)이 추가되었으며, 그중 7건은 동일 신고자가 제출했고 모두 재현 가능한 prim_func와 줄 단위 대조 생성 코드를 첨부했다.
  • 경계 조건이 깨지면 내부 오류 발생: #3738은 “패딩 후 공유 tile보다 작은 전역 영역에서의 T.copy(예: K=100을 BK=112로 복사)가 TMA 전체 블록 복사로 판정되어 강제 실패하는” 문제를 수정하고, #3766은 cta_group::2의 TMA 바이트 카운트를 수정한다——한 쌍의 CTA 바이트가 모두 짝수 CTA의 배리어에 떨어지는데, cluster 배리어가 있는 tma_copy는 원래 rank 0에 대해서만 cluster_size × tile_bytes를 기대하여 (2,1,1)과 (1,2,1)에서는 우연히 맞았고 (4,1,1)과 (2,2,1)에서는 틀렸다.
  • 연산자와 예제의 경계: #3709는 CPU 스칼라 GEMM이 목적 영역만 정리하도록 하고, #3712는 LegalizeNegativeIndex가 도달 불가능한 분기를 건너뛰도록 한다. #3714는 fp32에서 tf32로의 변환에 짝수 반올림(cvt.rn.tf32.f32, 비포화는 SM90 이상, 포화는 SM100 이상 요구)을 추가하고, #3702는 여러 어텐션 예제가 shared_memory_per_block_optin에서 tile을 추론하도록 하여 RTX PRO 6000(sm_120)에서 CUDA 9개 항목이 “2개 통과 7개 실패”에서 9개 전부 통과로 바뀌었으며, H100은 227 KiB 선택 상한이 있어 여전히 원래 tile을 사용한다.
  • Ascend 측의 두 가지 의미 편차: #3762는 T.simd.vcvt(src, "uint8", sat=False)가 int32 벡터에 대해 sat=False를 조용히 무시하고 항상 포화한다고 보고한다——문서에 적힌 랩어라운드 의미(-1이 255, -128이 128, 300이 44)가 실제로는 절단(-1이 0, -128이 0, 300이 255)으로 측정되며, int8 양자화 마무리의 f32에서 int8로의 경로에 영향을 준다. #3789는 PTO 백엔드가 정수 T.atomic_add에 대해 부호 있는 pto.si32/si64를 발행하지만 PTO 원자 인터페이스는 부호 없는 캐리어만 받는다고 보고하며, 당일 #3792와 #3794 두 방향의 수정안이 나왔다.

런타임, CI 및 툴체인

  • CI의 두 가지 안정성 문제 모두 처리됨: #3784는 AscendC CI의 자동 튜닝 공급 케이스를 수정했고(위 참조), #2833은 TileOPs의 RUNNER_VERSION을 2.336.0에서 2.338.2로 올렸다(GitHub가 구 runner를 폐기한 후 gpu-smoke가 전부 대기열에 쌓였고, entrypoint.sh가 --disableupdate를 포함해 자체 업그레이드가 불가능했음). #2827은 TileOPs의 off-CUDA 스위트를 야간 작업으로 옮겼다(원래 preflight에서 이 두 단계가 ubuntu-latest에서 13~18분 걸렸는데, 옮긴 후 preflight에는 수집 검사만 남기고 테스트 보고서와 cuda_only 필수 실패 검사는 제거해 약 3분 소요).
  • 병렬 병합에서 post-merge 게이트 부재의 실증: #2824가 21:06에 병합될 때 grouped GEMM의 roofline 행을 sizes를 반드시 선언하고 그로부터 G를 도출하도록 변경했는데, #2830이 21:08에 병합되며 여전히 G만 있고 sizes가 없었다. 두 커밋은 약 2.5분 간격으로 각자의 CI는 모두 통과했지만, 병합 후 main의 GPU Smoke가 즉시 RowError: row misses ['sizes']를 보고했고 당일 #2837로 보수했다. 이는 이번 주 가장 기록할 만한 프로세스 기록이다: 단일 커밋의 CI 통과는 ‘두 커밋의 의미적 교차’를 커버할 수 없다.
  • 문서와 로컬라이제이션: tilelang.github.io는 이번 주 네 개의 수집 기간에 각각 3, 5, 3, 3건의 Update docs 자동 빌드가 있었고, 내용은 #line 등급 매핑, ThreadSync 스코프 목록, host-mapped 어노테이션, TF32 rounding 및 Z3 설정 등 인터페이스 변경에 대응한다. TileOPs.github.io는 한 번 푸시가 있었고, TileOPs-nightly는 브랜치 푸시만 있었으며 메인 브랜치에는 새 커밋이 없었다.
  • Ascend 측 툴체인 세부 사항: #1956은 AI Core 예외 덤프를 CANN 9.2.0과 호환되도록 했다(acldumpTensorInfo의 필수 필드 보충, 빈 주소와 크기 0 텐서 건너뛰기, 등록 인터페이스를 폐기된 aclrtSetExceptionInfoCallback에서 aclrtExceptionInfoCallbackRegister로 교체, 버전 기준을 9.3에서 9.2로 하향). #1932는 A2/A3 성능 튜닝 스킬을 Ascend950 워크플로 구조와 정렬하고 칩 유형 식별 및 칩 정보 검색 두 가지 스킬을 새로 추가했다.

리뷰 대기열 관찰

  • 가장 큰 규모의 항목은 여전히 Ascend 전체 저장소 마이그레이션(#1975)이다. 11-04 00:34 업데이트 시점에 추가 607317줄, 삭제 97595줄, 2887개 파일에 달했으며, TileLang-Ascend를 메인라인 멀티 백엔드 아키텍처로 옮기고 기존 Ascend 구성 요소를 ascend/a2a3 아래로 이동시켰으며 arch="dav-2201"과 메인라인의 A5 경로로 분기한다. 작성자는 설명에서 할 일을 매우 구체적으로 적었다: 세 디렉터리의 파일과 원본 파일 간 차이 수렴 대기, 정적 검사 및 포맷팅 CI 복구 대기, 문서와 Agent 관련 마이그레이션 대기.
  • TIRx 마이그레이션 초안(#3803)은 이번 주 가장 비중 있는 심사 항목이다. 이는 이후 TVM 측 수정이 어떤 리듬으로 메인라인에 들어올지 결정하고, Ascend, Hygon 및 mlir-ascend의 글루 레이어를 함께 변경해야 할지도 결정한다.
  • TileOPs의 심사 대기열은 선택기와 레이아웃에 집중: #2836은 coop2_swap의 3단계 타일을 위해 별도 피팅 계수(1.36)를 적용하고, #2838은 ‘K를 연속 차원으로 간주해 TN 레이아웃 K=1의 합법적 호출을 거부하던 문제’를 수정하며, #2839는 FP8 그룹 GEMM이 TMA로 A의 스케일을 로드하도록 하고 K ≤ 1536인 비융합 긴밀 그룹에 64×256 ping-pong 타일을 적용하며, #2840은 SM89의 prefill sinks를 보충하고 sink 유무를 컴파일 인스턴스 아이덴티티에 포함시킨다.
  • 심사 대기 수 추이: 82, 44, 48, 42, 57건. 주 초반의 높은 수치는 지난주 적체에서 비롯되었고, 주 중반 42건으로 떨어진 후 금요일에 반등했으며, 병합량의 주말 반등과 동기화되어 정상적인 정산—신규 등록 리듬으로 보인다.

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

Ascend: ‘적응 대상’에서 ‘최초 출시 대상’으로

  • 최초 릴리스 측면: FlashMLA 커널(#3678)은 Ascend 950을 대상으로 메인라인에 직접 작성되었으며, 별도의 저장소를 만들어 포팅하는 대신 T.Stage, T.annotate_buffer_addresses와 같은 명시적 스케줄링을 위한 저수준 기능을 사용합니다. 목표 형상은 헤드 수 64, Dqk와 Dv 512, topk는 64의 배수입니다. 이에 수반되는 배치 전송 기능은 11-03에 보완되었습니다: #3689는 [B, M, N] 영역을 한 번의 T.copy로 작성하고, lowering이 배치 ND2NZ 디스크립터를 생성하며, 소스의 배치 및 행 stride를 유지하고, 전치 배치는 DN2NZ로 처리하며, GM 경계를 벗어난 소스 배치는 건너뜁니다. 검증은 Ascend950DT와 CANN 9.2에서 수행되었으며, 50개의 AscendC 예제 테스트와 1733개의 Ascend 단위 테스트가 통과했고, 그중 28개의 BF16/FP32 디바이스 케이스가 배치 및 M/N 테일 블록, 전체 배치 경계 초과, 비연속 소스 stride, 목적지 저장소가 소스보다 큰 경우 및 직사각형 전치를 커버합니다.
  • 동기화 정확성: #1981은 copy_pipe_to_l1의 동기화 어노테이션을 수정했습니다——A2/A3의 PTO 크로스코어 V→C 통신은 GM을 경유하며, 세 단계 전송이 각각 AIV의 MTE3, AIC의 MTE2, AIC의 MTE1에 속하는데, 이것이 operation_config.h에서 PIPE_MTE3로 표기되어 있었습니다. 따라서 자동 동기화는 AIC 자체 코어의 MTE3를 기다리게 되지만 해당 큐는 비어 있어 이벤트가 즉시 설정되며, MTE1의 L1 읽기가 MTE2 쓰기 완료 전에 허용될 수 있습니다. 재현 증거는 두 가지입니다: 생성된 소스 코드에서 PIPE_MTE3를 PIPE_MTE2로 변경하면 10/10 모두 정확하게 복구되고, 로컬 최신 메인라인에서는 10/10회 실행 오류가 발생하며 단일 실행당 최대 16384/16384개 요소가 불일치합니다. 수정 방법은 이름을 분리하는 것이지 직접 교체하는 것이 아닙니다: copy_pipe_to_l1은 PIPE_MTE2로 표기 변경하고, 새로 추가된 copy_pipe_to_l1_bind는 PIPE_MTE1(A5 직결 경로)로 표기하여 A2/A3를 고치면서 A5에서 오류가 발생하지 않도록 합니다.
  • 메인라인으로의 수렴: #1971의 RFC는 동기화 구현의 교체 경로를 제시합니다——메인라인 동기화 코어와 독립 Z3 증명 서비스를 이식하고, 구 TIR 어댑터 계층을 통해 기존 컴파일 파이프라인에 연결하며, TVM, 스케줄링, 메모리 계획, AscendC/PTO codegen 및 런타임 ABI를 유지하고, 검증 통과 후 기본 구현으로 전환하고 구 알고리즘을 삭제합니다. 이에 수반되는 #1965는 메인라인의 z3 버전 insert sync를 Ascend 저장소로 마이그레이션합니다; #1975는 저장소 수준 골격의 마이그레이션입니다. 세 건은 동일한 경로의 세 가지 계층입니다. 같은 시기 #3752는 Z3의 rlimit을 kernel별로 구성할 수 있게 하고, #3735는 thread sync에서 모듈러 인덱스의 서로소성을 증명하는 것을 제안하여, 바로 “구성 해결 자원”과 “증명해야 할 형상 감소” 두 측면에서 지난주 롤백이 남긴 증명 예산 문제를 처리합니다.
  • 호스트 오버헤드 측면: #1948은 CythonKernelWrapper.forward()에서 매 호출마다 무조건 생성되던 TVM Analyzer()를 지연 생성 및 필요 시 재사용으로 변경했습니다. 910B3에서 vec add 단일 블록은 442~467마이크로초에서 105.1마이크로초로, 32 블록은 441~493에서 108.0으로, kpool 압축은 447~450에서 111.3으로 감소했으며, 순수 ctypes 호출은 약 12.5마이크로초로 거의 변하지 않았습니다. 이어지는 #1967은 cython wrapper에서 심볼 형상의 파싱 결과를 캐시하여 심볼 경로의 비용을 계속 줄입니다. 같은 계열로 심사 중인 #3757이 있습니다(출력 이름 대신 Var 객체를 키로 사용하여 동일 이름의 다른 변수가 같은 심볼로 취급되는 것을 방지).
  • 연산자 공급: #1827은 WY Fast 전방을 추가하고, #1828은 wy_fast_bwd_split의 분할 역방향을 마이그레이션 및 보완했으며(두 건은 11-02 02:33 전후로 1분 이내에 병합), 선형 어텐션의 이 연산자 체인을 Ascend에서 학습 가능한 수준까지 보완했습니다. #1866은 Chunk O 전방(chunk 게이트 선형 어텐션)을 마이그레이션했으며, 개발자 모드와 네 가지 pass 구성을 사용하고, 단일 @tilelang.jit 커널과 고정 20코어의 grid-stride 루프로 온칩 융합, GM workspace 제로, 호스트 ATen 연산자 제로를 달성했습니다. #1879는 MHA 전방 예제를 추가했으며(BSHD 레이아웃, Expert 모드의 삼분배 MMA 파이프라인, wrapper에 T.* 코드 없음), cann-bench 20개 케이스 정밀도 전부 통과, torch-npu 베이스라인 대비 평균 가속비 0.97배, 종합 점수 73.13(로컬 910B2 기준)——0.97배는 아직 순이익이 없음을 의미하지만 골격은 제시되었습니다. #1966은 심사 중인 GLM-5.3 kpool fp8_mqa_logits 예제로, 이중 kernel과 AIV/Cube 분업이며 헤드 수 32, 차원 128, 페이지 크기 64입니다.
  • 엔지니어링 기준: #1954는 T.tile.bitwise_rshift가 int64/uint64에서 컴파일 타임 세그멘테이션 오류를 발생시키던 것을 실제 dtype과 지원 목록을 포함한 ValueError로 변경했습니다(int16/uint16/int32/uint32만 허용). “64비트를 새로 추가하지 않는다”는 점을 명확히 했습니다.

오른쪽 시프트 지원」; #3783은 Cube와 Vector의 코어 비율을 명시적으로 구성하거나 장치에 따라 추론할 수 있게 하고(vector_core_num / cube_core_num, 값을 얻지 못하면 2로 복귀), 선택된 그룹 수를 sid 영역, 작업 할당, 동기화 및 코드 생성까지 일관되게 전달한다; #3776은 AutoSchedule의 시작 간격에 단일 작업 경계, 자원 그룹 경계, 자기 의존 경계의 세 가지 필요 조건을 보강하고, 세 가지의 최댓값으로 이분 탐색을 초기화하며, Ascend 950PR 및 CANN 9.3.0 주간 업데이트 버전에서 모든 예제가 컴파일을 통과하고 스케줄링 결과가 변하지 않으며, Z3 해결 단계의 누적 소요 시간이 10.04% 감소했다.

  • 회귀 상태: 월요일 모니터링 기간 일일 테스트 2970건 전부 통과, 11-01에 워크플로우 실패 1건 발생(#1976); 수요일과 목요일의 일일 테스트 기록은 3030건 전부 통과, 통과율 100%(#1980, #1988). 테스트 케이스 규모가 일주일 사이 2970에서 3030으로 증가하여 커버리지가 여전히 확장 중임을 보여주지만, #1991/#1992와 #3762 두 가지는 수동 보고를 통해서만 발견된 문제로, ‘통과율’을 정확성의 대체 지표로 삼을 수 없음을 보여준다.

Hygon: 자체 개발 브랜치에서 공개 리뷰로 진입

  • 메인 브랜치: #31이 10-30에 병합되어 동적 MLS 경계 추론과 MatrixStore의 stride 처리를 개선하고, MMAC과 blockscaled GEMM의 clear_accum 하향을 수정하며, 지속 WASP 경로에서 M 파티션 컨슈머와 독립적인 A/B 펜스를 사용하고 slot과 phase 라운드 로빈을 유지한다; #32가 11-02 10:46에 병합되어 gfx948을 gfx946 호환으로 처리하고, 기존 MLS atoms, MMAC, barrier 및 컴파일 옵션을 재사용하며, 목적은 gather load를 통합하기 전에 기존 builtins이 새 모델에서 먼저 동작하도록 하는 것이다. 이는 전형적인 2단계 접근으로, 먼저 새 아키텍처 번호를 기존 경로에 연결하여 테스트 가능한 베이스라인을 확보한 후, 실제로 새 명령이 필요한 부분을 보충하는 것이다.
  • 기능 브랜치: feat/hcu-scaled-cast가 11-04의 8시간 내에 네 건을 커밋했으며, 각각 E8M0 스케일링의 FP8/BF8/FP4 변환(스칼라와 패킹 두 가지 lowering, 223줄의 HCU scaled_cast.h와 lower_scaled_cast.cc 신규 추가, 메인 변환 파일 650줄 변경), bank-word 지오메트리에 따른 plain-copy의 swizzle 유도 및 좁은 b16 tile을 위한 명시적 역변환 제공(12개 파일 +1536/-112, 테스트 파일 5개 신규 추가), 혼합 LDS wrap의 ‘최대 2경로’ 증명, 혼합 wrap 검사를 물리적 읽기 폭으로 변경; 11-05에 다시 세 건 커밋(기본 swizzle을 64 bank에 적응, MN 슬라이스 컨슈머 측 swizzle 계획 복원, K 파티션의 레이아웃 인식 읽기).
  • 리뷰 상태: 11-05 09:08에 #33(scaled cast, 즉 브랜치 본체) 오픈, 19:52에 #34(reduction의 XOR shuffle을 DPP와 lane swap으로 최적화, 추가로 ‘레이아웃 추론이 알려진 제약을 먼저 전파한 후 자유 후보를 할당’하는 수정) 오픈. ‘기능 브랜치에 네 건을 쌓고 PR 미오픈’에서 두 개의 정식 PR로 전환한 것은 국산 백엔드 중 처음으로 이 경로를 완주한 사례이며, 주제가 공유 메모리 bank 충돌과 위상 정렬에 집중되어 있어 해당 백엔드가 이미 ‘컴파일 가능’에서 ‘마이크로아키텍처 기반 튜닝’ 단계로 진입했음을 보여준다.
  • 별도로 리뷰 중인 #19(제목이 Distributed만 있고 설명이 비어 있으며, 11-03 11:06 업데이트)가 있으며, 본 보고서는 구현 세부 사항에 대해 추론하지 않는다.

Moore Threads(MUSA)와 MetaX: 하나는 문서 유지보수 재개, 다른 하나는 여전히 업스트림 동기화만

  • Moore Threads: 11-05 19:46에 메인 브랜치에 docs: update README installation and operator development references가 커밋되어, 해당 저장소에서 09-30 이후 첫 번째 커밋이다; 같은 시기 6개의 잔여 PR(의존성 업그레이드, TileOPs 서브모듈, MUSA 통합 증거 임계값 문서)의 타임스탬프가 함께 갱신되었으나 내용은 여전히 04월에서 08월의 오래된 제안이다. 판단은 보수적으로 해야 한다: 문서 커밋은 저장소에 유지보수 인력이 있음을 보여주지만 백엔드 개발 재개를 증명하기에는 부족하며, 진정한 신호는 이후 연산자나 lowering 커밋이 나타나는지 여부를 봐야 한다.
  • MetaX: 데일리 리포트는 다섯 개 모니터링 기간 모두 ‘푸시 없음’으로 기록했으나, 168시간 모니터링 기간 재검토 시 해당 메인 브랜치에 이번 주 확실히 커밋이 있었음을 확인했으며, 건별로 보면 전부 업스트림 tilelang 동기화(#3760, #3752, #3759, #3802, #3714, #3784, #3772, #3534, #3761, #3689, #3546, #3736, #3758, #3748, #3678, #3734, #3702, #3703 등)이며, 작성자는 모두 업스트림 유지보수자로 MetaX 측 자체 커밋은 보이지 않는다. 따라서 정확한 표현은: 메인 브랜치는 업스트림 동기화 중이며, 자체 개발은 여전히 침묵기로 연속 세 번째 모니터링 기간이다. dev 브랜치 최신은 여전히 10-16에 멈춰 있다.

mlir-ascend와 TLIR: 의미 체계를 한 층씩 방언에 내재화

  • 구조 정리: #206은 10-31에 병합되어 CI에서 내부 캐시 서비스를 가리키던 내용을 제거했으며, 이는 재현성 정리에 해당한다. #209는 11-04 16:35에 병합되어 BishengIR 툴체인 적응 코드를 contrib로 옮겼고, 「업스트림 툴체인의 진화에 따라 빈번한 수정이 필요한 글루 레이어」를 핵심 lowering과 분리하여 이후 BishengIR 업그레이드 시 변경 범위가 contrib 내로 수렴되도록 했다.
  • Agent 워크플로: #208은 harness 브랜치에 병합되었으며(메인라인이 아니므로 인용 시 「메인라인에 병합」으로 표기하지 않도록 주의), 2715줄, 66개 파일을 추가하여 연산자 harness에 첫 번째 Codex 워크플로 세트를 도입했다. 프로젝트 초기화 시 반복 질의를 사용하고, 기존 연산자 저장소를 연결하며, 영속화된 컨텍스트에서 개발을 복원하고, 튜닝 또는 분석을 수행한다. 세 가지 스킬은 .workspace/ 아래의 편집 가능한 Markdown 기록을 공유하며, 작업 인계 시 「현재 후보」와 「검증된 최적 후보」를 구분하고 실패 및 롤백된 실험을 보존한다. 작성자는 이번에 실제 NPU 실측이 없었음을 명확히 했으며, NPU 실행 적응과 npm 배포는 후속 작업으로 명시했다.
  • 방언 연산자: #203은 TLIR에 tl.sub/mul/div/max/min을 추가하고 TL에서 HIVM으로의 lowering을 보완했으며, #211은 tl.shl, #212는 tl.shr을 추가했다. 세 건은 11-05의 01:42, 19:44, 20:58에 순차적으로 진행되었다. 이 라인은 속도는 느리지만 방향은 단일하다. TileLang의 시맨틱을 한 층씩 MLIR 방언에 내리는 것이다.

4. 생태계와 채택 측

TileOPs: 0.1.0 이후 세 번째 주, 튜닝에서 선택기 엔지니어링으로 전환

  • 머지 리듬: 월요일 모니터링 기간 15건(13건은 11-01), 화요일 7건, 수요일 1건, 목요일 7건, 금요일 9건으로 주간 합계 39건입니다. 화요일 배치에서는 테스트 인프라(#2816의 tileops_test 타깃, 1281줄, 59개 파일로 tests/ops의 스모크 테스트가 더 이상 CUDA에 의존하지 않게 함)와 CI 계약(#2818)이라는 두 가지 비성능 변경이 나타났는데, 이는 연산자 라이브러리가 다중 백엔드와 유지보수성을 위해 구조적 준비를 시작했음을 보여줍니다. 목요일과 금요일은 순수 성능 작업으로 돌아갔습니다.
  • 어텐션: #2808은 GQADenseFwdOp와 GQADenseWithLSEFwdOp의 프로그램 선택을 하나의 규칙 AttentionCall.dense_decode_program으로 통합했습니다. 입력은 query와 key 길이, batch, 헤드 그룹, head dim, RoPE, SM 수이며, 2~64개의 query 위치는 warp 전용 decode를 사용합니다. 가장 빠른 사용 가능 프로그램 대비 선택 후회(기하 평균 및 최댓값)는 1.818과 25.2에서 1.017과 1.81로 감소했고, 메인라인의 “score scale에 log2(e)를 곱한 뒤 FP32에서 사라져 NaN이 발생”하는 문제와 “D256 다중 위치 호출이 SM90에서 공유 메모리 부족으로 시작 실패”하는 두 가지 실제 오류를 수정했습니다. #2817은 warp 전용 decode의 softmax를 QK, PV와 겹치게 했습니다(확률은 레지스터에 유지, K와 V는 독립 배리어 사용, 공유 3슬롯 링 버퍼). H200에서 (16, 8, 4096, 64, 8, 128)의 float16은 0.074610밀리초 대 fa3의 0.077490(1.039배, 3.654 TB/s), bfloat16은 0.074260 대 0.077220(1.040배, 3.671 TB/s)입니다. #2807은 페이지드 MHA decode를 “하나의 블록 공유 메모리에 들어가는” 모든 head dim으로 확장했고, #2797은 GQAPagedFwdOp에서 기존 페이지드 decode kernel을 삭제했습니다.
  • MoE와 GEMM: #2815는 tight 레이아웃 MoE 입력을 토큰별로 수집하고 동일한 hidden 행을 여러 라우팅 간에 재사용하게 했습니다. (4096, 256, 8, 7168, 2048)에서 bfloat16은 7.4417밀리초 대 flashinfer-cutlass의 7.6982(1.034배, 387.844 TFLOPS), float16은 7.5040 대 7.9869(1.064배, 384.624 TFLOPS)입니다. #2805는 SM의 최소 4분의 3을 채우는 부분 1웨이브 W4A8 prefill 그리드를 stream-K 대신 CTA당 타일 하나로 변경했습니다. llama-8b-up-b33은 23.52마이크로초에서 17.98로 감소했고(vLLM 20.42), 목록 외 M 33~128의 44개 N/K 스캔에서 스케줄 형태 변경이 5%~36% 빠르며 중앙값 18%, 느려진 경우는 없었습니다. #2834는 작은 M의 NT GEMM을 위해 1웨이브를 채울 수 있는 너비의 네이티브 폭 타일을 만들었습니다. decode-down(128×7168×2048)은 메인라인에서 132개 SM 중 112개만 사용했는데, 변경 후 가장 빠른 비교 항목보다 3% 앞서고 gate-up-b32는 메인라인보다 1.4% 빠릅니다.
  • 원소별 연산과 리덕션: #2814는 bfloat16의 Sigmoid와 16비트 SiLUAndMul이 4분의 1 스케일 역수를 사용해 범위 검사가 있는 나눗셈을 우회하게 했고, #2813은 8바이트 피연산자가 긴 시작에서 스레드당 벡터 하나로 처리되게 했습니다(먼저 TileFoundry의 roofline 분석으로 이러한 비트 OR 연산자가 모두 메모리 바운드임을 확인). #2832는 lerp, where, clamp, masked_fill 다섯 연산자가 stride 0 브로드캐스트 텐서를 출력 형태로 먼저 확장하지 않게 했고, 브로드캐스트 행 소요 시간이 메인라인의 0.13~0.44배로 감소했습니다. #2829는 “본문에 최대 두 개의 저렴한 명령, 읽기/쓰기 총량이 L2의 4분의 3 이하”인 단항 원소별 연산을 256 스레드의 1웨이브로 축소하고, ReLU가 NaN 입력에 0을 반환하던 오래된 문제도 수정했습니다. #2828은 행 단위 블록 ArgMax/ArgMin에서 “1웨이브 상주” 전제로 스레드당 8개의 16바이트 벡터를 갖게 했고, (2048, 4096)의 bf16 ArgMin이 8.704마이크로초에서 7.393으로 감소했습니다(FlagGems 8.543). #2835는 리덕션 kernel의 축 치환과 출력 형태를 인스턴스당 한 번만 계산하게 했고, 여섯组 FP32/BF16 VarMeanFwdOp 구성의 전체 핫 호출 소요 시간이 4.33%~6.17% 감소했으며 생성된 CUDA와 호스트 소스는 완전히 동일합니다. 절약된 것은 호출 준비 비용이지 GPU 시간이 아닙니다.
  • 정밀도류 결함이 정면으로 수정됨: #2811은 행으로 구성된 각 fp32 량에 불변식 “행의 확산도에만 제약되고 오프셋에는 제약되지 않음”을 추가했습니다. 구현상 행 내의 어떤 원소를 기준으로 삼아 뺄셈하며(warp 주도 커널은 warp 첫 원소, 분할 커널은 행의 첫 원소), 1e36 근처의 오버플로를 제거합니다. #2821은 분산 병합을 빈 누산기가 아닌 첫 실제 부분량에서 초기화하게 했고, 상수 입력의 분산이 더 이상 NaN을 내지 않습니다(tests/ops/test_reduce.py가 155에서 158로 증가). 이

‘공 누산기 + 큰 오프셋’ 같은 수치 함정은 유지 축 앞뒤에서 동시에 리덕션하는 형태에서만 드러나며, 일반적인 전체 축 리덕션으로는 측정되지 않는다.

  • 측정 기준이 제품 특성이 되다: #2830은 roofline을 ‘이번 호출이 실제로 얼마나 읽었는가’ 기준으로 다시 작성했다. GQA의 dense, varlen, paged forward는 ‘query가 볼 수 있는 key 행’에 대해서만 K/V를 계산하고, key가 실제로 도달한 절단 위치에 대해 RoPE 테이블을 계산한다. GroupedGEMM은 새로운 grouped_gemm_roofline을 통해 batch_sizes에 따라 특정 행의 그룹 가격을 매긴다. 수정 후 물리적으로 불가능했던 세 행이 7.48, 7.50, 5.30 TB/s에서 0.62, 0.62, 2.65 TB/s로 바뀌었다. #2822는 bench_kernel의 L2 플러시를 ‘int8 버퍼에 zero_() 수행’에서 ‘보고된 L2 크기의 두 배인 전용 int32 버퍼 읽기’로 변경했다. 기존 flush는 약 120 MB의 더티 행을 남겨 그 writeback이 다음 타이밍 호출에 영향을 주었기 때문이다(2048×11008 f32의 Sum이 27.81에서 21.38 마이크로초로 감소).
  • 선택기와 비용 모델: #2825는 W4A16의 _select_config를 _legal_configs에서 _config_cost의 argmin을 취하는 방식으로 변경하고, tile 우선, 가장 넓은 K의 동점 허용 범위, joint_switch_margin 세 가지 예외 규칙을 삭제하고 stage_fill 항목을 추가했다. #2826은 dense 선택기의 보정 계약을 문서 문자열에 명시했다(스캔, 피팅, 검사 3단계, 소프트 어닐링과 홀드아웃 검증 포함). #2836은 coop2_swap의 3단계 타일에 별도 피팅 인자(1.36)를 적용하여 regret을 기하 평균 0.89%에서 0.06%로, 최대 5.04%에서 0.99%로 낮췄다.
  • 릴리스 형태와 테스트 기준: #2816의 tileops_test는 스모크 테스트를 CUDA에서 분리했으며, 작성자는 경계를 명확히 했다. 공유 참조와의 비교는 백엔드 독립적인 계약과 이식성만 증명할 뿐 실제 백엔드의 수치 정확성을 증명하지 못하며, 그 부분은 여전히 #2611에 걸려 있다. #2827은 off-CUDA 스위트를 야간으로 옮겼고, #2833은 runner를 업그레이드했다.

TileFoundry: 동기 분석, 학습 시맨틱, 명명 거버넌스

  • 동기: #283의 분석 레이어와 #285의 lowering 실행단(요점 5번 항목 참조) 외에, #286이 11-06 13:40에 머지되어 각 피연산자 효과에 대한 완료 조건을 선언하고, WGMMA의 fence, commit, wait를 필요한 HIR 노드에 연결했다. lowering은 명령의 발행 지점에서 이러한 단계를 소비하며, 후속 스코프가 집단 대기를 수용할 수 없으면 발행 스코프에서 완료한다. 동시에 레지스터 Reshard의 별칭을 수정하여 요소가 사본을 보유한 모든 스레드에서 동일한 레지스터에 유지되도록 했다. 작성자는 두 가지 위험 선언을 유지했다. 행렬 명령 발행기가 아직 연결되지 않았고, 이 커밋은 실행 정확성을 증명하지도 성능을 제공하지도 않는다. TMA의 완료 시맨틱은 별도로 처리한다. 11-06 20:12에 #288을 새로 열어 ‘완료된 부분 리덕션’의 분석 내 가격 산정을 처리한다.
  • 학습 시맨틱: #284가 11-05 16:36에 머지되었다. 기존 평가기는 reshard를 전역 텐서로 취급하여 서로 다른 참여자의 기여 차이를 놓쳤다. 이 커밋은 참여자별로 지역 값, GPU 좌표, 호출과 루프, gather/scatter 및 축 방향 Partial(sum) 리덕션을 계산하여 Reshard의 별칭 증명이 소유권에 민감하도록 했다. 검증: 코어 스위트 775개, 백엔드 스위트 778개(2카드 sm90에서 13개 스킵), 10개 HIR 구성이 CPU에서 통과, 두 가지 상주 전략이 4카드 sm90에서 실제 2×2 PyTorch HSDP와 두 업데이트 스텝을 실행하여 로컬 그래디언트, 파라미터, 모멘텀이 모두 일치했다.
  • 스케줄링과 릴리스: #280과 #282는 ‘각 호출이 무엇을 발행하는가’, ‘어떤 mesh가 발행하는가’, ‘윈도를 어떻게 자르는가’에 대한 결정을 analysis/emission.py 한 곳에 모았다. 검증은 바이트 단위 대조로, 83개 스케줄링 픽스처 중 75개의 finalize 출력이 메인라인과 완전히 일치하고 나머지 8개는 메인라인에서도 동일하게 실패한다. 사이트 저장소는 10-30에 0.3 버전 스케줄링 글 ‘Scheduling, Designed for Agents’를 게시하여 ‘에이전트를 위한 스케줄링’을 코드와 릴리스 노트에서 공개 서사로 옮겼다.
  • 명명 거버넌스: #287은 문서 지향 이슈로, 코드베이스에서 ‘하나의 개념에 다섯 개의 이름’이 있음을 지적한다. mesh의 한 멤버가 placement 133곳, program 110곳, instance 92곳, participant 44곳으로 불리며, unit이라는 단어는 빈 타입, 하드웨어 연산 유닛, mesh 멤버를 동시에 지칭하고, Python과 CUDA 런타임의 local_layout 시그니처도 일치하지 않는다. 이런 이슈는 보통 즉시 코드를 바꾸지 않지만, 컴파일러 프로젝트가 ‘작성 가능’에서 ‘장기적으로 읽힐 수 있음’으로 가는 데 반드시 거쳐야 할 단계다.

TileSight: 성능 모델이 논문에서 저장소로

  • tile-ai/TileSight(24 stars)는 07-24 arXiv 논문 《TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters》의 오픈소스 구현이다. 타일 형상, 데이터 이동, 소프트웨어 파이프라인 및 하드웨어 프로파일로부터 CPU에서 연산자 비용, 캐시 트래픽과 실행 시간을 추정한다. 11-05에 저장소의 첫 PR(#1)이 등장했다. H20은 Hopper/GH100이며, requires_arch=("wgmma",)가 포함된 계획은 원래 아키텍처 프로파일이 WGMMA 지원을 선언하지 않아 UnsupportedError를 던졌는데, 이번 커밋은 support_wgmma=True로 설정하고 support_utcmma=False를 명시적으로 유지하여 이러한 계획이 수용되고 tcgen05와 Tensor Memory 요구는 여전히 거부된다. 검증은 CPU에서 847개 항목 전부 통과이며, 작성자는 명확히 “GPU 성능 검증이 아니다”라고 선언했다.
  • 이 라인의 의의는 TileOPs의 roofline 원가 계산, TileFoundry의 분석 계층과 삼각 구도를 형성한다는 점이다. 하나는 해석 모델을, 하나는 라인별 원가 계산을, 하나는 증명 가능한 동기화 계획을 제공한다. 세 가지 모두 이번 주에 실질적인 움직임이 있었다.

채택자: TileKernels와 FlashQLA

  • 두 채택자 저장소는 주중 내내 커밋이 없었다. deepseek-ai/TileKernels는 10-14 05:49에 멈춰 있고, QwenLM/FlashQLA는 10-30 03:02에 멈춰 있다(최신 릴리스는 여전히 09-30의 v0.1.3).
  • 그러나 채택자가 푸시하지 않는다고 해서 영향을 받지 않는 것은 아니다. 목요일 #3761에서 수정된 배리어 점유 회귀의 발화점은 바로 TileKernels의 mhc_pre_big_fuse 커널이었다. 동일한 커널이 B300, 숨김 차원 1280에서 배리어 점유 예산의 25%를 잃었다. 이는 채택자의 커널이 이미 컴파일러 측의 회귀 프로브가 되었음을 보여준다. TileKernels가 코드를 푸시하지 않은 몇 주 동안에도 그 커널은 메인라인의 변경을 통해 수동적으로 “테스트되고” 있다. 생태계를 추적하는 사람에게 이는 주목할 만한 역방향 신호다.

커뮤니티 브랜치와 서드파티 저장소

  • RISC-V 연구 체인 형성: 커뮤니티 연구자 zishu0131-png가 10-31 같은 날 세 개 저장소를 정리했다. tilelang-riscv는 총 진입점(연구 목표, 두 단계 컴파일 체인과 버전 조합 설명), numba2tilelang은 Numba CUDA Python을 정적 파싱한 후 TileLang 소스로 변환, tilelang2rvv는 지원되는 실제 tvm.tirx.PrimFunc를 RVV intrinsic C로 변환한 후 Clang을 거쳐 RV64GCV ELF로 컴파일하여 QEMU virt에서 단일 hart로 실행한다. 작성자는 경계를 명확히 밝혔다. TileLang 공식 등록 백엔드가 아니며, threadIdx에 의존하는 주소 지정, hart 간 동기화와 지원되지 않는 IR은 일괄 거부하고, block은 순차 실행하며, 연구 환경은 TileLang 0.1.15, Clang 23.1.3, QEMU 11.1.0이다. 168시간 모니터링 기간의 4개 출처 점검에서 이번 주 새로운 변화는 없었다.
  • PTO-ISA/tilelangGym 독립 저장소화: 이 민감도 및 회귀 테스트 세트는 PTO-ISA/tilelang의 examples/ascend/t_parallel_st_suite에서 분리되었다. 이유는 테스트 세트가 TileLang과 PTOAS 버전을 독립적으로 고정할 수 있게 하기 위해서다. 현재 세 가지 주요 경로는 Ascend SoftLib SimtVF(target=ascend, cython, 주요 회귀 게이트), VectorizeParallelToPTO와 VL의 vload/vstore, 그리고 benches 아래 @pto.jit로 작성된 수동 SIMD DSL 대조 구현이며, 각 테스트 케이스 디렉터리에는 config.toml과 VERDICT.md가 있다.
  • 서드파티 소규모 저장소: sablin39/tilelang-cuda-skills(149 stars, TileLang 작성과 CUDA 툴체인 디버깅을 위한 스킬 세트 제공), xwhzz/tilelang-hexagon(2 stars, TileLang을 Qualcomm Hexagon DSP로 이식), Index4NekoMeteor/Tilelang-AttnRes(1 stars, TileLang으로 Kimi K3의 어텐션 리지듀얼 구조 구현), Tianchen-ai/tilelang-weft-bench(0 stars, CPU 측 연산자 벤치마크 통합 엔진), flytiger-eco/tilelang-wrapper-for-sail(0 stars, 설명 없음). 처음 세 가지 방향은 흥미롭다. 하나는 툴체인 측 개발 경험을, 둘은 새로운 대상 아키텍처와 최전선 모델 구조로의 이식을 다룬다. 나머지 둘은 매우 초기 단계로 우선 관찰 목록에 기록하되 기술 진전으로 작성하기는 적절하지 않다.

5. 커뮤니티, 튜토리얼 및 이벤트

  • 문제 보고가 설치 형태에 집중됨: 수요일 하루에 올라온 7개의 새 이슈 중 네 건이 pip 패키지를 사용하고 WSL에서 실행하며 소비자용 sm_89 카드(RTX 4060 Laptop)를 쓰는 동일한 외부 보고자로부터 나왔고, 모두 재현 가능한 스크립트와 전체 출력을 첨부했습니다. 기술적 깊이가 가장 높은 #3775는 프로덕션 환경의 «pip 패키지 + 듀얼 H200 NVLink + 다중 CUDA 컨텍스트 + nvrtc 백엔드 + CUDA graph 캡처» 조합이었으며, 보고자는 세 번 실행의 대조 조건을 모두 기록했습니다. 바이트 단위로 동일한 구성이 세 번 재현되었지만 첫 실패 지점은 세 번 모두 달랐고, /proc/environ과 깨끗한 대조군의 차이는 TILELANG_EXECUTION_BACKEND=nvrtc 환경 변수 하나뿐이었습니다. 이는 컴파일러의 심층 결함은 아니지만 새 사용자가 처음 시도할 때의 경험을 결정합니다.
  • Ascend 쪽에서 «보고자가 곧 리뷰어»인 폐루프 발생: #1991/#1992의 보고자는 #1897의 리뷰어였는데, #1897의 수정 방식은 count 측 의미(HasModifiedBuffer가 workspace 슬롯을 건너뜀)만 커버했고 workspace 측은 여전히 처리되지 않았습니다. 동일한 레이아웃 로직의 다른 측면이 2주 후에야 지적된 것입니다. 이는 Ascend 백엔드의 레이아웃 불변 조건이 아직 검사 가능한 문서화된 계약으로 형성되지 않았음을 보여줍니다.
  • 문서 및 현지화 체인 정상: tilelang.github.io에는 주중 여러 건의 Update docs가 자동 동기화되었고, TileOPs.github.io는 한 번 푸시되었습니다. LSP 저장소는 CUDA 디버깅을 새로운 행 매핑에 맞췄고(71e762e0), VS Code 확장 저장소는 문서를 업데이트하고 산출물 이름을 vsix로 변경했습니다. 이는 11-02의 «CUDA 디버깅을 TileLang 행 매핑에 맞춤» 이후의 작업으로, 지난 호 메인라인에 머지된 행 단위 앵커링이 생성 코드의 행 번호 출처를 바꾸었기 때문에 확장 측 디버그 구성도 따라 변경해야 했습니다.
  • 튜토리얼 및 지식 자산: 이번 주에는 새로운 튜토리얼 저장소가 나타나지 않았고, 커뮤니티 투자는 기존 저장소 유지보수와 위에서 언급한 두 개의 RISC-V 방향 소규모 저장소에 집중되었습니다. PTO-ISA/tilelangGym의 독립 저장소 분리는 테스트 자산이 처음으로 별도 버전 관리된 사례입니다.
  • 뉴스 측 전 주간 실질적 적중 제로: 다섯 개 수집 창의 Google News 중영문 쿼리(정확한 용어, GPU kernel, 중국어 연산자와 컴파일러, Ascend와 MetaX/Hygon/Moore Threads 포함)에서 가장 최신 관련 보도는 여전히 10-21의 SemiAnalysis «The Huawei-DeepSeek Symbiosis»였고, 나머지는 09-30의 «DeepSeek, Ascend版 TileLang·DeepGEMM·DeepEP 오픈소스» 보도 물결에 집중되어 모두 이번 주기 이전의 것입니다. Hacker News에서 tilelang 검색 24시간 결과는 0건이었고, arXiv 최신 적중은 11-03의 «Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving»(TileLang을 인용한 작업, 창 외부)과 07-24의 TileSight 본체였습니다. 따라서 보고서 본문은 전부 코드 측에서 나왔으며, 이는 이미 여러 호 연속 안정적인 형태입니다.
  • 산업 측의口径 충돌은 기록해 둘 만함: 수요일 창에서 Wccftech는 DeepSeek가 Ascend 950DT로 구성된 Atlas 950 SuperPoD에서 V5를 처음부터 학습할 것이라고 보도한 반면, 같은 시기 The Information의 보도는 «이르면 2026년 말부터 950DT에서 사전 학습 시작»이었고, Bloomberg와 Financial Times의 초기 보도는 950DT의 16만 장 주문이 주로 추론용이며 이전에 Ascend에서 학습이 좌절된 후 Nvidia로 회귀했다고 전했습니다. 세 가지说法은 «학습에 사용되는지 그리고 언제 사용되는지»에 대해 일치하지 않으며 모두 Huawei나 DeepSeek 공식 확인을 받지 못했습니다. 본 보고서는 배경 기록으로만 남기며, 이를 근거로 Ascend 측 투자의 실행 리듬을 추론하지 않습니다.

6. 트렌드 관찰

첫째, TIRx 마이그레이션은 이미 메인라인의 사실상 최우선 순위이며, 그것이 «뺄셈»의 형태로 나타나기 시작했습니다. 일주일 동안 세 가지 증거가 있었습니다. 월요일에는 마이그레이션을 위해 ThreadSync 정확성 개선 전체를 되돌렸고(#3733, 동시에 TVM fork 측 최적화 제안 #89 종료), 화요일에는 3384줄의 레거시 파서 분기를 삭제했으며(#3736), 금요일에는 전체 마이그레이션 초안을 제출했습니다(#3803, 업스트림 커밋 두 개를 pin하고 패치를 주입하지 않음을 명시). 중간에는 Z3 리소스를 하드코딩에서 구성 가능하게 바꾼 작업(#3752와 tvm #91)도 있었습니다. 이 일련의 조치의 공통점은 복잡성을 크리티컬 패스에서 철수시키는 것입니다. 일시적으로 보수적 구현으로 되돌아가더라도 vendored 의존성의 내부 구조가 계속 언어 계층에 자라도록 두지 않겠다는 것입니다. 대가는 메인라인이 마이그레이션 창에서 «알려진 비최적» 경로를 일부 갖게 된다는 것이고, 이득은 기반 교체가 완료된 후 모든 다운스트림 백엔드가 안정적인 IR 경계에 맞춰 적응할 수 있다는 것입니다.

둘째, 정확성 리스크는 이미 형태가 바뀌었다: 「컴파일 불가」에서 「컴파일 통과, 실행 시 오류 없음, 결과 오류」로 바뀌었고, 게다가 버전 간에 잠복한다. 이번 주에 최소 일곱 건이 이 계열에 속한다: #3808은 0.1.12부터 0.1.15까지 네 버전에 걸쳐 __syncthreads() 하나가 빠져 있고, #3800은 Blackwell의 tensor memory 경로에서 MMA가 덮어쓴 operand를 읽게 만들며, #3721은 sm_100a에서 전체 GEMM 곱을 조용히 버리고, #3775는 프로덕션 환경에서 CUDA graph 캡처를 손상시키고, #3749는 의존성 분석이 물리적 로드 범위를 보지 못하게 하며, #1991/#1992는昇腾 threads=2에서 sin/cos가 97.1%의 원소를 잘못 계산하게 만들고, #3774와 #3773은 소비자용 카드에서 잘못된 결과를 낸다. 이 계열은 사용자에게 가장 큰 피해를 준다. 왜냐하면 「컴파일러가 맞는가」라는 전제를 검증이 필요한 가정으로 바꿔버리기 때문이다. 이번 주의 대응 방향은 정확성을 거부 가능한 증명으로 만드는 것이다—TileFoundry의 #283/#285는 먼저 증명하고 발사하며, 발사 후 역검증하고, #3786과 #3787은 조용히 잘못된 코드를 생성하는 것을 컴파일 시점 오류나 명시적 거부로 바꾸는 것이다—단위 테스트 수를 계속 늘리는 것이 아니다.

셋째, 昇腾의 역할이 「적응 대상」에서 「최초 출시 대상」으로 바뀌었고, 엔지니어링이 두 축에서 동시에 추진되고 있다. 한 축은 아키텍처다: FlashMLA 같은 DeepSeek 자체 디코딩 경로의 핵심 커널이 昇腾 950을 대상으로 메인라인에 직접 작성되었으며, 사용한 것은 여전히 T.Stage, annotate_buffer_addresses 같은 명시적 스케줄링을 위한 저수준 기능이다; 컴파일러 측의 세 가지配套 변경과 함께 昇腾 경로는 「돌아간다」에서 「스케줄 의도대로 돌아간다」로 전환 중이다. 다른 축은 골격이다: #1975는 2887개 파일, 60만 줄 추가로 전체 TileLang-Ascend를 메인라인 다중 백엔드 프레임워크로 옮겼고, arch="dav-2201"로 백엔드 분기를 한다; mlir-ascend의 #208은 Codex의 초기화, 연동, 복구 세 가지 워크플로와 작업 인계를 .workspace/에 기록했다. 「마이그레이션 저장소」와 「에이전트로 마이그레이션 저장소를 옮기기」를 같은 주에 놓고 보면, 昇腾 측의 병목이 이미 「연산자 부족」에서 「일관된 엔지니어링 골격과 재사용 가능한 마이그레이션 프로세스 부족」으로 전환되었음을 알 수 있다.

넷째, 연산자 라이브러리의 경쟁점이 완전히 선택기 품질과 준비 비용으로 옮겨갔고, 성능 격차의 잔여 공간은 「어떻게 계산하는가」에 있지 않다. 금요일 TileOPs의 아홉 건 머지 중 GPU kernel을 실제로 바꾼 것은 브로드캐스트 인플레이스 읽기와 네이티브 폭 타일 두 건뿐이다; reduction 레이아웃 캐시가 절약한 것은 호출 준비이며(4.33%에서 6.17%, 생성 소스 코드는 불변), 3단계 타일 팩터가 바꾼 것은 상수 하나다(regret 0.89%에서 0.06%로 하락). 여기에 roofline 가격 책정 재작성과 L2 플러시를 읽기 전용으로 바꾼 것을 더하면 명확한 노선이 보인다: 먼저 각 형상의 합법적 구성을 열거하고, 통일된 비용 모델로 정렬하며, 캘리브레이션 과정을 재현 가능한 계약으로 작성한다(#2826은 이미 「스캔, 피팅, 검사」 세 단계를 문서 문자열에 기록했다). 이 길을 제대로 하면 성능은 더 이상 개별 커널의 수동 튜닝에 의존하지 않는다; 제대로 하지 못하면 「특정 형상은 특히 빠르고 다른 형상은 운에 맡기는」 상태로 퇴화한다.

다섯째, 국산 백엔드가 「한 번에 큰 백엔드 머지」에서 「리뷰 가능한 최소 단위를 공용 저장소에 푸시」로 전환했다. 海光은 자체 개발 브랜치에서 #33과 #34 두 개의 정식 PR로 나아갔고, 摩尔线程은 README 한 건으로 두 달의 침묵을 끝냈으며, mlir-ascend는 세 개의 작은 PR로 tl.sub/mul, tl.shl, tl.shr을 하나씩 TLIR 방언으로 옮겼고, 昇腾은 핵심 비율, 예외 덤프, 성능 기능 같은 엔지니어링 세부 사항을 항목별로 보완했다. 세 노선의 공통점은 더 이상 큰 머지를 기다리지 않는다는 것이다: 이는 업스트림에 좋은 일이다. 리뷰어가 3만 줄짜리 일괄 머지 때가 아니라 일찍부터 배치와 lowering 설계에 의견을 낼 수 있기 때문이다; 다운스트림에도 대가는 있다—메인라인이 이런 모니터링 기간에는 「알려진 몇 가지가 작동하지 않는」 상태에 놓이므로, 더 강력한 출시 전 회귀 규율이 필요하다.

마지막 리스크 영역: 기반 전환이 모든 글루 레이어의 비용을 한꺼번에 노출시킨다. 만약 #3803이 추진되면, 3rdparty/tvm의 vendored 구조가 두 개의 pin 업스트림 커밋으로 대체되며, 그때 모든 TVM 내부 API에 의존하는 백엔드(昇腾, 海光, mlir-ascend의 글루 레이어)가 함께 수정해야 한다. 초안에는 「의존성 패치를 주입하지 않는다」고 명시되어 있는데, 이는 약속이자 제약이다—이번 주 mlir-ascend가 BishengIR 어댑터를 contrib로 옮기고, 昇腾이 성능 기능을 칩 세대별로 분기한 것은 모두 이런 기반 전환을 위해预留한 교체 가능한 면으로 볼 수 있다. 이 노선이 건강한지 판단하는 기준은 마이그레이션 초안이 머지된 후 2주 이내에 이 몇 개 백엔드가 핵심 lowering을 건드리지 않고 적응을 완료할 수 있는지다.

부록: 자료 및 검증 설명

출처 검증표

출처 검증 결과
GitHub tile-ai 조직 저장소별(29개 저장소) 5개 수집 기간에 각각 10, 9, 9, 8, 10개 저장소가 푸시됨; 168시간 기간 재검증 결과 14개 저장소에 푸시 있음
병합 확정 검색(is:pr is:merged) 5개 기간에 35, 18, 8, 12, 18건, 합계 91건; 첫 기간은 72시간(10-30 07:00 ~ 11-02 07:00)
심사 중 확정 검색(is:pr is:open updated) 5개 기간에 82, 44, 48, 42, 57건, 합계 273건
신규 issue 확인 5개 기간에 11, 8, 7, 4, 5건, 합계 35건
메인 브랜치 커밋 단위 검증 tilelang 7, 5, 5, 3, 5건; TileOPs 15, 7, 1, 7, 9건; TileFoundry 3, 1, 1, 1, 1건; tilelang-ascend 5, 4, 0, 1, 2건; tilelang-mlir-ascend 1, 0, 1, 2, 0건; tilelang-hygon 메인 브랜치 2, 1, 0, 0, 0건; tilelang-lsp와 VS Code 확장 각 1건; tilelang-musa 1건
릴리스 페이지 및 태그 tilelang 최신 v0.1.15(09-30), 주간 신규 버전 없음; TileOPs 최신 v0.1.0(10-27); TileFoundry 최신 0.0.4rc1(10-25); tilelang-ascend v0.1.2.000(09-09); tilelang-metax v0.1.15(10-16)
채택처 TileKernels 10-14에 정체; FlashQLA 10-30에 정체, 최신 릴리스는 여전히 09-30의 v0.1.3
TileLang×RISC-V 4개 출처 점검 168시간 기간 내 변화 없음
서드파티 및 생태계 저장소 연구 체인 3개 저장소, tilelangGym, cuda-skills, hexagon, AttnRes, weft-bench, wrapper-for-sail, TileSight 모두 푸시 또는 신규 생성 있음
Google News(중영문 4~6개 조합) 5개 기간 모두 기간 내 실질적 적중 없음, 최신 관련 보도는 10-21의 SemiAnalysis 기사
Hacker News tilelang로 24시간 검색 결과 0건, 수집된 것은 주제와 무관한 Show HN
arXiv 최신 적중은 11-03의 KV 압축 논문(TileLang 인용, 기간 외)과 07-24의 TileSight 본체

주요 인용

com/tile-ai/tilelang/pull/3783>

설명

  • 본 보고서의 사실, 버전 번호 및 날짜는 모두 모니터링 기간 내 5회차 데일리 리포트 원문 및 그 일차 출처에서 가져온 것이며, 추측성 보충은 하지 않았다. 다섯 개 수집 기간은 서로 연결되어 있으며(10-30 07:00부터 11-06 07:00까지 총 168시간), 이전 위클리 리포트와 겹치지 않으므로 다섯 기간의 머지, 심사 중 및 신규 issue 수는 그대로 합산하였다.
  • 머지 수, 심사 중 수 및 issue 수는 GitHub 검색의 실측 출력에서 가져왔으며, 메인 브랜치 커밋 수는 저장소별 atom 검증 결과이다. 두 기준은 서로 다르므로(한 번의 머지가 여러 커밋에 대응될 수 있고, 기본 브랜치가 아닌 커밋은 머지 수에 포함되지 않음) 부록에서는 두 범주를 각각 나열하고 합산하지 않았다.
  • TileOPs의 “15건”은 월요일의 72시간 기간에 나타난 것으로, 그중 13건이 11-01에 집중되어 있으므로 인용 시 “월요일 하루 15건”으로 표기해서는 안 된다.
  • #1991과 #1992는 제목과 본문이 일치하여 “동일 제목 2건”으로 기록하고 하나로 합치지 않았다. Hygon 기능 브랜치의 7건과 Moore Threads의 문서 커밋은 브랜치 수준 동향으로 기술한 것이며, 메인라인에 이미 머지되었음을 의미하지 않는다.
  • MetaX는 다섯 기간 모두 “푸시 없음”으로 기록되었으며, 168시간 재검증 결과 메인 브랜치에 커밋이 있었지만 건별 확인 결과 전부 업스트림 tilelang의 동기화(작성자는 업스트림 유지관리자)였으므로 “메인 브랜치는 업스트림 동기화 상태이며 자체 개발은 여전히 정체 기간”으로 기술하였다. TileSight와 TileFoundry의 WGMMA 변경은 모두 작성자가 capability 메타데이터 또는 동기화 계획으로 한정하여 서술하였으므로, 보고서는 그 한정을 그대로 따랐다.