TileLang 데일리 리포트 (2026-11-04)
모니터링 기간: 2026-11-03 07:00 ~ 2026-11-04 07:00 (베이징 시간, 24시간; 이전 기간 기준점은 11-03 07:00).
출처: GitHub tile-ai 조직 저장소별 점검(29개 저장소, 모니터링 기간 내 푸시 9건) + 머지 및 심사 중 실측 검색(머지 8건, 심사 중 업데이트 48건) + 신규 issue 전수 확인(7건) + 메인 브랜치 커밋별 점검(tilelang 5건, TileOPs 1건, TileFoundry 1건, VS Code 확장 1건, 문서 사이트 3건) + 릴리스 페이지 및 tag 점검(5개 저장소, 모니터링 기간 내 신규 버전 없음) + 채택 기관 2개 저장소 + Google News 중영문 검색(프록시 경유) + Hacker News + arXiv.
이번 호 인덱스
- 오늘의 핵심: ThreadSync가 저장소 도메인 간 부분 배리어 재사용을 허용하여 #3643 정확성 수정 시 도입된 성능 회귀를 복구(#3761); 3384줄의 레거시 TVM 파서 브랜치 정리(#3736); 자동 튜닝 입력 공급을 현재 디바이스 기준으로 변경(#3546); TileFoundry 동기화 계획이 공식적으로 CUDA lowering에 반영(#285); 24시간 내 8건 머지, 48건 리뷰 중 업데이트, 신규 issue 7건
-
- 핵심 프로젝트 진행 상황
- 1.1 메인라인: ThreadSync가 저장소 도메인 간 부분 배리어 재사용을 허용, B300에서 배리어 점유가 12에서 16블록/SM으로 복귀(#3761)(11-04)
- 1.2 메인라인: 레거시 TVMScript AST 및 파서 브랜치 제거, 한 번에 3384줄 삭제(#3736)(11-03)
- 1.3 메인라인: 자동 튜닝의 입력 공급을 현재 디바이스로 변경, Ascend 및 MPS에서 튜닝 사용 가능 복구(#3546)(11-03)
- 1.4 리뷰 중: 조건 내 스칼라 cast가 더 이상 데이터 벡터 폭을 압축하지 않음, 어제 #3734 후속(#3759)(11-03)
- 1.5 리뷰 중: 두 곳의 사용자 노출 오류를 명확한 에러로 변경(#3765 / #3767)(11-03)
- 1.6 리뷰 중: padding이 있는 공유 tile의 전체 복사 불가 시 폴백, 2-CTA TMA 바이트를 peer leader 기준으로 계산(#3738 / #3766)(11-03)
- 1.7 리뷰 중: B300의 sm_103a 두 라인 적용, tcgen05 int8 선택 및 CuTeDSL TMA 아키텍처 게이트(#3770 / #3772)(11-03)
- 핵심 프로젝트 진행 상황
-
- 멀티 백엔드 적응
- 2.1 Ascend: 3D GM에서 L1으로의 배치 ND2NZ 복사 머지, 단일 디스크립터로 전체 배치 커버(#3689)(11-03)
- 2.2 Ascend 리뷰 중: TileLang-Ascend 전체 저장소를 메인라인 멀티 백엔드 아키텍처로 이전, 2887개 파일(#1975)(11-03)
- 2.3 Ascend 리뷰 중: copy_pipe_to_l1의 동기화 주석을 실행 엔진별로 분리, A2/A3에서 10/10 재현되는 오류 수정(#1981)(11-03)
- 2.4 Ascend 리뷰 중: 첫 Parallel 축 브로드캐스트를 스칼라 읽기로, Chunk O 포워드 마이그레이션, MHA 예제(#1979 / #1866 / #1879)(11-03)
- 2.5 mlir-ascend: Codex 워크플로 및 작업 인계를 harness 브랜치에 머지(#208)(11-03)
- 2.6 Hygon: Distributed 리뷰 중 PR 업데이트 있음, 메인 브랜치 신규 커밋 없음(#19)(11-03)
- 2.7 MetaX, Moore Threads 모니터링 기간 내 푸시 없음(11-03)
- 멀티 백엔드 적응
-
- 생태계 및 채택사
- 3.1 TileFoundry: 동기화 계획을 CUDA lowering으로 완성 및 역검증, 메인라인 머지(#285)(11-04)
- 3.2 TileFoundry 리뷰 중: wgmma 동기화 계획, FSDP 및 HSDP 훈련 시맨틱을 HIR 평가기에 추가(#286 / #284)(11-03)
- 3.3 TileOPs: 분산 병합을 첫 번째 부분량부터 계산, 상수 입력의 분산에서 더 이상 NaN 발생 없음(#2821)(11-03)
- 3.4 TileOPs 리뷰 중: GEMM에 선택적 bias 추가, FP8 그룹 GEMM 연속 레이아웃 지원(#2823 / #2824)(11-03)
- 3.5 TileOPs 리뷰 중: W4A16 및 dense GEMM 선택기를 단일 비용 모델로 재적합(#2825 / #2826)(11-03)
- 3.6 TileOPs 리뷰 중: 짧은 varlen을 KV 그룹별 패킹, L2 flush를 읽기 전용으로 변경(#2819 / #2822)(11-03)
- 3.7 채택사: FlashQLA 및 TileKernels 푸시 없음, 단 #3761은 TileKernels의 fused 커널 수정(11-04)
- 생태계 및 채택사
-
- 커뮤니티, 튜토리얼 및 이벤트
- 4.1 하루 7건 신규 issue: 프로덕션급 CUDA graph 문제 1건, 동일 외부 보고자로부터 4건(11-03)
- 4.2 Ascend 950의 vcvt 포화 시맨틱 및 mlir-ascend 일일 테스트(#3762 / #1980)(11-03)
- 4.3 문서 사이트 3회 업데이트, VS Code 확장 vsix로 개명(11-03)
- 4.4 뉴스 측: 3개 채널 모니터링 기간 내 해당 없음, 외신의 Ascend 훈련 보도 논조 상이(11-03)
- 커뮤니티, 튜토리얼 및 이벤트
-
- 트렌드 관찰
- 부록. 자료 및 검증 설명
오늘의 핵심
이날의 주된 흐름은 「성능을 되찾는」 수정이었다. 이전 모니터링 기간에 병합된 #3643은 실제 정확성 문제를 해결했다: 각각 호출되는 ThreadSync("shared")와 ThreadSync("shared.dyn")가 서로 겹치지 않는 스레드 그룹에 동일한 명명된 배리어 ID를 할당할 수 있었다. 수정 방식은 두 번째 호출이 첫 번째 호출에서 이미 사용한 ID를 보존하도록 하는 것이었지만, 이 보존 규칙이 동시에 다른 한 가지를 막았다 — 동일한 스레드 그룹이 두 스코프 사이에서 같은 ID를 재사용할 수 있었던 것이다. 그 대가는 곧바로 TileKernels의 mhc_pre_big_fuse 커널에 떨어졌다: 이 커널은 블록당 96개의 스레드를 가지며, 스레드 32부터 95까지가 2단계 cp.async 파이프라인을 실행하는데, 비동기 대기 이후의 동기화가 첫 번째 패스에서 ID 3을 받고, 동일한 64개 스레드 배치의 버퍼 의존 동기화가 두 번째 패스에서 ID 4를 받아, 추가된 그 ID 때문에 컴파일 산출물이 선언하는 배리어 자원이 4에서 5로 늘어났다. B300에서 히든 차원 1280을 실행할 때 배리어 점유 상한이 이로 인해 16블록/SM에서 12블록/SM로 떨어졌으며, 레지스터, 공유 메모리, 동기화 명령어 수에는 변화가 없었다. #3761의 처리 방식은 ThreadSync가 정렬된 스코프 목록을 받아들이도록 하고, CUDA 파이프라인을 ThreadSync(["shared", "shared.dyn"])로 바꾸는 것이었다. 동기화는 여전히 원래의 스코프 순서대로 계획되고 삽입되지만, 부분 배리어의 ID는 두 스코프가 모두 처리된 후에야 통일적으로 할당된다. 스코프 간에 ID를 공유하기 전에, 새로운 분석기로 참여 프레디케이트가 전체 후보 스레드 범위에서 성립함을 증명하며, 분기와 Bind 컨텍스트는 ConstrVisitor가 제공하고, 스레드 관련 프레디케이트는 기본적으로 성립하는 것으로 간주하지 않고 증명 의무로 처리한다. 증명이 완료된 밀집 그룹은 바운딩 박스 부피를 직접 사용하고 부분 스레드 수 열거를 건너뛰며, 증명되지 않은 그룹은 스코프를 키에 남기고 참여자 카운트도 함께 키에 넣는다. 명시적으로 번호가 매겨진 배리어(더 이른 ThreadSync가 할당한 ID 포함)는 계속 보존된다. 변경 규모는 401줄, 4개 파일에 걸쳐 있다.
작성자는 설명에서 매우 깔끔한 귀인 실험을 하나 제시했다: 컴파일 산출물에 선언된 배리어 자원 수만 단독으로 바꾸고 실행 가능한 명령어는 전혀 건드리지 않으면 동일한 성능 저하를 재현할 수 있다는 것이다. 이런 실험은 기록해 둘 가치가 있다 — 「성능 회귀」를 막연한 느낌에서 지목 가능한 단일 변수로 바꾸어 주었고, 수정의 합격 기준도 명확하게 만들었다: 포괄적으로 점수가 빨라지는 것이 아니라, 자원 선언이 4로 돌아가고 점유가 16블록/SM로 돌아오는 것이다.
두 번째는 공급 측의 감산과 기준 통일이다. #3736은 tilelang/language/ast/, tilelang/language/parser/와 tilelang/language/tir/entry.py를 함께 삭제했다. 이들은 초기 로컬 TVMScript 파서 분기에서 남은 모듈로 약 3380줄의 Python이며, 저장소 내에 더 이상 호출자가 없다: 언어 계층은 이제 tilelang.language.tir.common을 통해 업스트림의 tvm.tirx.script.parser 인터페이스를 사용하고, AST와 명령문 오버라이드 등록은 tilelang.language.overrides에 있으며, T.prim_func와 T.macro는 tilelang.language.eager에서 내보낸다. 실제 변경은 1줄 추가, 3384줄 삭제, 10개 파일이며, 동시에 docs/conf.py의 *language/ast*에 대한 무시 규칙을 제거했다. 공개 언어 계층 내보내기는 방언 테스트와 프런트엔드 테스트가 뒷받침한다. #3546은 자동 튜닝의 디바이스 기준을 수정했다: set_autotune_inputs가 설치하는 공급 클로저가 호출 지점에서 전혀 전달되지 않는 파라미터에서 기준 디바이스를 해석했는데, 실제로 전달되는 것은 커널 파라미터 목록이었고, 그 결과 CUDA 또는 CPU 폴백 분기가 CUDA가 없는 머신에서 캡처된 모든 텐서를 CPU로 옮겨, 매 튜닝 라운드가 런처 디바이스 불일치로 실패했다; 이제 클로저는 get_current_device()를 사용하며, 기본 텐서 공급과 동일한 전략을 취한다. CUDA 동작은 변하지 않고, Ascend와 MPS는 현재 디바이스에서 공급하도록 바뀌었으며, 나머지 백엔드는 여전히 CPU로 폴백한다. 검증은 듀얼 카드 Ascend950에서 이루어졌다: 세 건의 Ascend 자동 튜닝 회귀가 수정 전 전부 실패, 수정 후 전부 통과했고, 메인 저장소 testing/python/autotune/의 12건 통과, 28건 스킵, 그 외 3건은 해당 머신에 CUDA가 없어 실패했으며, 작성자는 이를 명시적으로 CUDA CI 검증에 남겼다.
세 번째는 Ascend 측이다: GM에서 L1으로의 복사가 「하나의 디스크립터에 하나의 행렬」에서 전체 배치로 확장되었다. #3689는 [B, M, N] 영역을 하나의 T.copy로 작성할 수 있게 허용하며, lowering은 배치 ND2NZ 디스크립터를 생성하고, 소스 측의 배치와 행 stride를 보존하며, 목적 측의 배치 간격은 물리적 레이아웃에 따라 추론되고, 전치 배치는 DN2NZ로 처리된다. GM 경계를 벗어난 소스 배치는 그대로 건너뛰고 목적 저장소는 그대로 둔다. M/N의 테일 제로 패딩은 유효한 배치에만 적용된다. 기존 2차원 및 단일 행렬 동작은 변하지 않으며, 구현은 AscendC 백엔드만을 대상으로 하고 L1 목적 측이 정규 행렬 레이아웃일 것을 요구한다. PTO는 이번 범위에 포함되지 않는다. 검증은 Ascend950DT와 CANN 9.2에서 이루어졌다: 50개의 AscendC 예제 테스트와 1733개의 Ascend 단위 테스트가 통과했으며, 그중 28개의 디바이스 케이스가 배치와 M/N 테일 블록, 전체 배치 범위 초과, 비연속 소스 stride, 목적 저장소가 더 큰 경우와 직사각형 전치를 커버한다. CUDA와 PTO 케이스는 제외되었고, CI 이미지가 사용하는 CANN 9.3은 로컬에서 테스트되지 않았으며, 이 점은 작성자가 설명에 적어 두었다. 변경 규모는 579줄, 10개 파일이다.
네 번째는 TileFoundry의 숨은 작업 하나가 마무리된 건입니다. 지난 리포트에서는 「심사 중」으로 걸려 있던 것이 이번 모니터링 기간에 머지되었습니다: #285는 sync 분석이 계획한 커밋, 대기, 펜스를 실제로 CUDA lowering이 발행하도록 하고, 실제 발행과 계획이 일치하는지 항목별로 역검증하여 불일치 시 LoweringError를 던집니다. 코퍼스에서의 결과는 기본 흐름 14개 프로그램에 동기화를 더하고 31개를 거부하는 것으로, 이 14개의 TIR은 메인라인과 차이가 신규 문장뿐이며, 별도로 30개는 메인라인과 완전히 동일합니다; 단일 sm_89에서 빌드 가능하고 동기화가 포함된 프로그램 7개를 실행한 결과, 6개가 참조와 완전히 일치하고 synccheck와 memcheck 모두 깨끗했으며, 동일한 7개를 메인라인 방식으로 lower하면 4개가 결과가 틀렸습니다. 작성자가 동기에서 제시한 직접적 증상은 sm_89의 E4M3 GEMM 최대 편차 659입니다. 이번 변경은 788줄 추가, 214줄 삭제, 25개 파일로, 「분석이 증명할 수 있다」에서 「발행은 반드시 증명과 일치해야 한다」로 나아간 것입니다.
1. 핵심 프로젝트 진행 상황
1.1 메인라인: ThreadSync가 저장소 도메인 간 부분 배리어 재사용을 허용, B300에서 배리어 점유가 12에서 16블록/SM으로 회복 (#3761)
11-04 00:50에 머지되었으며, 401줄 추가, 72줄 삭제, 4개 파일입니다. 배경은 지난 모니터링 기간의 #3643입니다: ThreadSync("shared")와 ThreadSync("shared.dyn")을 분리 호출할 때 서로 겹치지 않는 스레드 그룹에 동일한 명명 배리어 ID를 할당할 수 있었는데, 이는 실제 정확성 문제이며, 수정 방식은 더 이른 호출이 이미 사용한 ID를 보존하는 것이었습니다. 부작용은 동일 스레드 그룹이 스코프를 넘어 재사용될 수 없다는 것으로, TileKernels의 mhc_pre_big_fuse 커널에서는 블록당 96스레드에 한 번의 불필요한 할당이 발생했습니다: 스레드 32~95의 2단계 cp.async 파이프라인이 첫 번째 패스에서 ID 3을, 두 번째 패스에서 ID 4를 받아, 컴파일 산출물이 선언하는 배리어 리소스가 4에서 5로 늘고, B300 숨김 차원 1280의 배리어 점유 상한이 16블록/SM에서 12블록/SM으로 떨어졌으며, 레지스터, 공유 메모리, 동기화 명령 수는 전혀 변하지 않았습니다. 이번 변경은 ThreadSync가 정렬된 스코프 목록을 받도록 하여 CUDA 파이프라인을 ThreadSync(["shared", "shared.dyn"])으로 바꿉니다: 동기화는 여전히 원래 순서대로 계획 삽입되고, 부분 배리어 ID는 두 스코프가 모두 처리된 후 일괄 할당됩니다; 스코프 간 ID 공유 전에 새 분석기로 참여 술어가 후보 스레드 범위 내에서 항상 성립함을 증명하며, 분기와 Bind 컨텍스트는 ConstrVisitor가 제공하고, 스레드 관련 술어는 가정이 아닌 증명 의무로 취급하며, 증명된 밀집 그룹은 경계 상자 부피로 계산하고 일부 스레드 수 열거를 건너뛰고, 미증명 그룹은 스코프를 키에 남기고 참여자 수를 키에 포함시키며, 명시적으로 번호가 부여된 배리어(더 이른 호출이 할당한 ID 포함)는 계속 보존되고, 겹치지 않는 스레드 그룹은 여전히 격리됩니다—#3643의 수정은 되돌려지지 않았습니다. 문자열 형태의 기존 호출은 호환성을 유지하며, CuTeDSL도 동일한 CUDA pass 파이프라인을 탑니다.
출처: https://github.com/tile-ai/tilelang/pull/3761
1.2 메인라인: 잔존 TVMScript AST와 파서 분기 제거, 한 번에 3384줄 삭제 (#3736)
11-03 13:05에 머지되었으며, 1줄 추가, 3384줄 삭제, 10개 파일입니다. 삭제된 것은 tilelang/language/ast/, tilelang/language/parser/, tilelang/language/tir/entry.py로—초기 로컬 TVMScript 파서 포크의 잔재이며, 저장소 내에 호출자가 전혀 없습니다: 언어 계층은 tilelang.language.tir.common을 통해 업스트림의 tvm.tirx.script.parser 인터페이스를 사용하고, TileLang 자체의 AST와 문장 오버라이드 등록은 tilelang.language.overrides에 있으며, T.prim_func와 T.macro는 tilelang.language.eager에서 내보내지고, tilelang/language/tir/ 아래에서는 common, exports, ir, op 네 모듈만 소비됩니다. 이 고립된 파서 포크를 유지할 때의 위험은 구체적으로 기술되어 있습니다: 누군가 tilelang.language.parser.operation을 직접 import하면 tirx.PrimExpr와 tirx.IterVar에 서로 충돌하는 연산자를 등록하게 됩니다. 이번 변경은 동시에 tilelang/language/tir/__init__.py에서 prim_func에 대한 무용한 재내보내기와 docs/conf.py의 *language/ast* 무시 규칙을 삭제합니다(*parser*는 유지되는데, tilelang/language/overrides/parser.py가 아직 사용 중이기 때문입니다). 수용 기준은 공개 내보내기가 변하지 않는 것이며, test_tilelang_language_dialect.py와 언어 프런트엔드 테스트로 커버됩니다.
출처: https://github.com/tile-ai/tilelang/pull/3736
1.3 메인라인: 자동 튜닝의 입력 공급이 현재 디바이스를 사용하도록 변경, Ascend와 MPS에서 튜닝 사용 가능 복구 (#3546)
11-03 13:19에 병합되어 153줄 추가, 12줄 삭제, 파일 3개, #3545 수정. set_autotune_inputs가 설치한 동결 공급 클로저가 「호출 지점에서 절대 전달되지 않는」 특정 매개변수로부터 기준 디바이스를 해석하는데, 공급 프로그램의 계약상 실제로 전달되는 것은 커널 매개변수 목록이다. 그래서 CUDA가 없는 머신에서 CUDA 또는 CPU 폴백 분기가 캡처된 모든 텐서를 CPU로 옮기고, 매 튜닝 라운드마다 런처 디바이스 불일치로 실패했다. 이제 클로저는 get_current_device()를 사용해 기본 텐서 공급과 동일한 전략을 취한다. CUDA에서는 동작이 동일하고, 멀티 카드 기준 worker는 여전히 카드별로 복사한다. Ascend와 MPS에서는 캡처된 텐서가 각자의 현재 디바이스에서 공급되어 튜닝이 가능해진다. 나머지 백엔드는 여전히 CPU로 폴백한다. 새로 추가된 Ascend 테스트 케이스는 testing/ascend/autotune/에 배치되어 데코레이터 경로, AutoTuner의 기준 공급 디바이스, 교차 디바이스 캡처 입력의 세 가지 상황을 커버하며, 듀얼 카드 Ascend950에서 수정 전에는 전부 실패하고 수정 후에는 세 항목 모두 통과했다. 메인 저장소의 testing/python/autotune/은 12개 통과, 28개 건너뛰기이며, 별도로 3개(eager 모드, 스칼라 입력, 입력 포함)는 해당 머신에 CUDA가 없어 실패했으며, 작성자는 CI의 CUDA runner 검증을 위해 남겨두었다고 명시했다.
출처: https://github.com/tile-ai/tilelang/pull/3546
1.4 검토 중: 조건 내 스칼라 cast가 더 이상 데이터 벡터 폭을 압축하지 않음, 어제 #3734 후속 (#3759)
11-03 14:09 업데이트, 148줄 추가, 29줄 삭제, 파일 2개. 이는 어제 병합된 #3734의 마무리 작업이다. #3734는 DecoupleTypeCast 이후 데이터 cast의 폭 제한을 올바르게 강화했지만, 「스칼라 조건 내의 cast」가 데이터 벡터 폭을 제한하기 시작하기도 했다. b5cb02633에서 가드된 BF16 복사가 조건 내 int64 cast 때문에 16개 lane에서 4개로 압축될 수 있었는데, 해당 조건이 16개 lane에서 완전히 동일함에도 그러했다. 수정 방법은 플래너에 이미 있는 조건 일관성 분석을 cast 폭 면제에 재사용하여, if 문, tirx.if_then_else, 스칼라 술어의 Select 조건을 그 절에서 분리해 별도로 방문하는 것이다. 조건 내의 cast는 더 이상 데이터 타입 폭 제한을 적용하지 않고, 절 내의 cast는 여전히 적용하며, lane별 벡터 술어가 있는 대상에서는 조건 cast의 제약을 유지한다. #3579의 메모리 가드 동작과 #3734에서 수정한 가드 FP32→BF16 제한은 모두 그대로 유지된다. 회귀 테스트는 8 lane과 16 lane의 일관 가드, 세 가지 조건 형식, 절 내 FP32 cast, 벡터 술어와 반올림 비트 그룹화를 커버하며, 9개 int64 가드 케이스가 수정 전 실패했고, 최종적으로 법적화와 cast 반올림 두 그룹 스위트에서 총 74개가 통과했다.
출처: https://github.com/tile-ai/tilelang/pull/3759
1.5 검토 중: 두 가지 사용자 가시적 오류를 명확한 오류 보고로 변경 (#3765 / #3767)
두 PR 모두 11-03 17:10 전후에 업데이트되었으며 동일한 보고자가 제출했다. #3765는 #3653에 대응한다. 단일 스레드 인덱스 경계 가드에 의한 부분 배리어에서, 해당 경계가 warp 크기의 정수 배일 때 전체 warp의 스레드가 조건 판단에 일치한다. 그러나 일반 산술 증명이 외부 분기가 비트 연산 &를 사용할 때의 thread_binding < 448을 놓쳐 ThreadSync가 바로 중단되었다. 이 PR은 스레드 차원이 정수 개의 warp일 때 이 상황을 명시적으로 허용하며(경계가 해당 차원과 같거나 초과하는 경우도 포함), 나머지 조건은 기존 증명 경로를 따른다. #3767은 #3005에 대응한다. fragment의 교차 스레드 분포에서 동일한 T.Parallel 본문 내에서 구조가 다른 두 인덱스로 같은 fragment를 읽고 쓰는 경우(예: buf[i]와 buf[(i + 1) % N] 같은 교차 lane 패턴)는 교차 스레드 교환 없이는 lower할 수 없다. 이전에는 ParallelOpNode::RecordBufferAccess가 StructuralEqual의 ICHECK에 부딪혀 컴파일러 내부 단언으로 노출되었다. 이제 ValueError로 변경되어 두 인덱스 표현식을 모두 오류 메시지에 기록하고 fragment를 이렇게 접근할 수 없음을 설명한다. 두 PR 모두 「내부 단언을 사용자가 읽을 수 있는 진단으로 교체」하는 성격이며, 사용자 체감 영향이 변경 규모보다 크다.
출처: https://github.com/tile-ai/tilelang/pull/3765, https://github.com/tile-ai/tilelang/pull/3767
1.6 검토 중: padding이 있는 공유 tile을 전체 복사할 수 없을 때 폴백, 2-CTA의 TMA 바이트를 피어 leader 기준으로 계산 (#3738 / #3766)
두 건 모두 11-03 13:47에서 13:31 사이에 업데이트되었습니다. #3738은 #3704를 수정합니다: padding 후 공유 tile보다 작은 전역 영역에서 T.copy를 수행할 때(예: K=100을 BK=112로 복사) ClassifyWarpSpecializedCopy와 SelectCopyInstForLowering이 이를 TMA 전체 블록 복사로 판정한 뒤 lowering 단계에서 s_extent == g_extent (112 vs. 100) 오류로 강제 실패합니다. 원인은 선택 로직이 마지막 차원 정렬, dtype 쌍, 전역 stride만 검사했을 뿐 padding 후의 범위가 하나의 전체 블록 복사를 구성할 수 있는지 확인하지 않았기 때문입니다. 이번 커밋은 이 판정 기준을 추가하여, 충족하지 않을 경우 일반 복사 경로로 폴백합니다. #3766은 cluster 시나리오에서의 TMA 바이트 카운트를 수정합니다: cta_group::2의 TMA load는 배리어 주소의 패리티 비트를 지우므로, 한 쌍의 CTA 바이트가 모두 짝수 CTA의 배리어에 떨어지는데, cluster 배리어를 사용하는 tma_copy는 원래 rank 0에 대해서만 cluster_size × tile_bytes를 기대했습니다. (2,1,1)과 (1,2,1)에서는 우연히 맞았지만, (4,1,1)과 (2,2,1)에서는 rank 0이 2개만 받고 4개를 기대하며, 추가 바이트는 rank 2의 배리어에 떨어집니다. 두 건 모두 신규 기능이 아니라 「경계 조건이 깨지면 내부 오류를 보고하는」 유형의 결함입니다.
출처: https://github.com/tile-ai/tilelang/pull/3738, https://github.com/tile-ai/tilelang/pull/3766
1.7 검토 중: B300의 sm_103a 적응 두 갈래, tcgen05의 int8 선택과 CuTeDSL의 TMA 아키텍처 게이트 (#3770 / #3772)
B300이 보고하는 대상 아키텍처는 sm_103a이며, 두 건의 검토 중 PR이 각각 이로 인해 발생하는 두 가지 걸림돌을 처리합니다. #3770은 #3769에 대응합니다: tcgen05.mma.kind::i8은 아키텍처 전용인 sm_100a와 sm_110a에만 존재하며, sm_103a(B300)에는 없고 패밀리 이식 가능한 sm_10xf에도 속하지 않습니다. 그런데 AllowTcgen5Mma는 TargetIsSm100(arch 100~110)에 대해 int8/uint8 피연산자를 일괄 허용하므로, 커널이 lowering을 순조롭게 통과한 뒤 ptxas에서 Feature '.kind::i8' not supported on .target 'sm_103a'를 보고합니다. 이번 커밋은 TileLang이 이러한 대상에서 더 이상 kind::i8을 선택하지 않도록 합니다. #3772는 CuTeDSL의 TMA 아키텍처 게이트를 처리합니다: B300의 sm_103a는 DSL에 진입하기 전에 거부되며, 오류 메시지는 sm_90/90a/100a/110/120/120a/121/121a를 나열할 뿐 103a는 포함하지 않습니다. 메인라인 fa69f7fa에 nvidia-cutlass-dsl 4.7.0과 CUDA 13.3 조합에서는 B300에서 여섯 개의 CuTeDSL 테스트가 실패하며, tma_load의 여러 JIT 케이스가 포함됩니다. 같은 배치에는 #3768도 있어 SM100 flash attention 역방향 예제를 수정합니다: T.tcgen05_gemm이 이제 mbar를 필수 키워드 전용 인자로 규정하고 더 이상 암묵적으로 대기하지 않는데, 예제는 여전히 동기 GEMM 방식으로 호출하여 두 스크립트가 컴파일 전에 missing 1 required keyword-only argument를 보고합니다. 작성자는 동시에 examples/flash_attention_sm100/ 디렉터리에 test_*.py가 없어서 이 문제가 계속 발견되지 않았다고 지적합니다.
출처: https://github.com/tile-ai/tilelang/pull/3770, https://github.com/tile-ai/tilelang/pull/3772, https://github.com/tile-ai/tilelang/pull/3768
2. 멀티 백엔드 적응 (Ascend / Hygon / Moore Threads / MetaX)
2.1 Ascend: 3D GM에서 L1으로의 배치 ND2NZ 복사 병합, 하나의 디스크립터로 전체 배치 커버 (#3689)
11-03 17:07에 병합되어 579줄 추가, 39줄 삭제, 10개 파일을 변경했다. 기존에는 Ascend의 GM에서 L1로의 복사에서 DMA 디스크립터 하나가 행렬 하나만 처리했기 때문에, 여러 개를 복사하려면 for b in T.serial(B)로 배치별 호출을 작성해야 했다. 이번 커밋은 [B, M, N] 영역을 한 번의 T.copy로 작성할 수 있게 하며, lowering은 배치 ND2NZ 디스크립터를 생성하고 소스의 배치 및 행 stride를 유지하며, 목적 측의 배치 간격은 물리 레이아웃에 따라 추론하고, 전치 배치는 DN2NZ로 처리하며, 기존의 행렬 축 경계 처리를 재사용한다. GM 경계를 벗어난 소스 배치는 건너뛰고 목적 저장소는 그대로 두며, M/N의 꼬리 부분 제로 패딩은 유효한 배치에만 적용된다. 기존 2차원 및 단일 행렬 동작은 변경되지 않으며, 구현은 AscendC 백엔드만 대상으로 하고 L1 목적 측이 정규 행렬 레이아웃일 것을 요구하며, PTO 지원은 이번 변경에 포함되지 않는다. 검증 환경은 Ascend950DT와 CANN 9.2이며, 50개의 AscendC 예제 테스트와 1733개의 Ascend 단위 테스트가 통과했고, 그중 28개의 BF16/FP32 디바이스 케이스가 배치 및 M/N 꼬리 블록, 전체 배치 범위 초과, 비연속 소스 stride, 소스보다 큰 목적 저장소, 직사각형 전치 복사를 커버한다. 작성자는 CUDA와 PTO 케이스가 제외되었으며 CI 이미지가 사용하는 CANN 9.3은 로컬에서 검증되지 않았다고 명시했다. 이 항목은 지난 호에 병합된 Ascend 950 FlashMLA 커널과 같은 방향의 짝을 이룬다. 어텐션과 GEMM의 배치 이동 모두 디스크립터 하나가 전체 배치를 커버해야 하며, 배치별 호출은 프로덕션 형상에서 상당한 발사 오버헤드를 유발한다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/3689
2.2 Ascend 심사 중: TileLang-Ascend 전체 저장소를 메인라인 멀티 백엔드 아키텍처로 이전, 2887개 파일 (#1975)
11-03 16:34 업데이트, 변경 규모는 607317줄 추가, 97595줄 삭제, 2887개 파일에 달하며, 이번 호에서 가장 큰 규모이자 전체 조직의 심사 중 PR 가운데 가장 주목할 만한 항목이다. 목표는 제안 #1911에 따라 TileLang-Ascend를 메인라인의 멀티 백엔드 아키텍처로 옮기는 것이다. 기존 Ascend 구성 요소는 메인라인의 해당 디렉터리 아래 ascend/a2a3로 이동하며(tilelang/는 tilelang/ascend/a2a3, src/는 src/ascend/a2a3), 기존 Ascend C와 PTO A2A3 백엔드는 메인라인에서 독립적으로 등록되어 메인라인의 Ascend, PTO A5와 직교하며, pass와 codegen은 서로 간섭하지 않고 소량의 인프라만 재사용한다. 양쪽은 target="ascend"와 "pto"를 공유하지만 A2A3 백엔드는 arch="dav-2201" 식별자를 동반하며(NPU 이름으로 자동 인식) 이에 따라 분기한다. A2A3 프런트엔드 원시 요소는 import tilelang.ascend.a2a3.language as T로 바뀌어 A5의 import tilelang.ascend.language as T와 공존하며, 연산자 작성 방식은 원래 저장소와 동일하다. pass와 codegen은 TileLang-Ascend에서 이전한 버전을 직접 사용해 원래 컴파일 흐름을 유지하고, 실행 백엔드는 메인라인 cython을 재사용하면서 workspace 인덱스 등 Ascend 고유 적응을 보완하며, examples와 testing은 각각 examples/ascend_a2a3와 testing/ascend_a2a3에 위치하고, CI는 당분간 원래 저장소의 ci_cd.yml을 사용하며 새 경로에 적응한 후에도 정상 동작함을 확인했다. 작성자는 설명에서 할 일을 매우 솔직하게 적었다. 세 디렉터리 아래 파일과 원본 파일의 차이는 아직 정리가 필요하고, 정적 검사와 포맷 CI는 복구가 필요하며, 문서와 Agent 관련 사항은 이전이 필요하고, “관련 파일이 많아 Agent 일괄 작업에 누락이 있기 쉽다”고 명시하며 검토 협조를 요청했다. 이처럼 이전 중의 불확실성을 PR 설명에 직접 드러내는 방식은 “완료됨”이라는 결론을 기다리는 것보다 리뷰에 더 유리하다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1975
2.3 Ascend 심사 중: copy_pipe_to_l1의 동기화 주석을 실행 엔진별로 분리, A2/A3에서 10/10 재현되는 오류 결과 수정 (#1981)
11-03 12:36 업데이트, #1780에 해당. 근본 원인이 명확히 기술되어 있다: A2/A3의 PTO 크로스코어 V 대 C 통신은 GM을 경유하며, 세 구간의 이송이 각각 세 엔진에 속한다—AIV의 MTE3가 UB에서 GM으로 전송, AIC의 MTE2가 GM에서 L1으로 수신(TPOP 내부 TLOAD), AIC의 MTE1이 L1에서 L0A로 이송; 그런데 copy_pipe_to_l1이 operation_config.h에서 PIPE_MTE3로 표기되어 있어, 자동 동기화 pass가 이와 copy_l1_to_l0a 사이에 set_flag(PIPE_MTE3, PIPE_MTE1, …)와 대응하는 wait_flag를 생성했지만—기다리는 것은 AIC 자체 코어의 MTE3이며, AIC의 MTE3 큐는 비어 있고 이벤트가 즉시 설정되어, MTE1의 L1 읽기가 MTE2의 쓰기 완료 전에 허용될 수 있어 L0A가 이전 데이터를 받아 MMA 결과가 틀리게 된다. 잘못된 표기는 #1238에서 도입되었으며, 같은 표에서 copy_pipe_to_ub와 copy_l0c_to_pipe는 모두 실행 측 엔진 기준으로 표기했지만 이 항목만 송신 측 관점을 따랐다. 재현 증거 두 가지: 이슈 작성자가 생성된 소스 코드에서 PIPE_MTE3를 PIPE_MTE2로 변경한 후 10/10 전부 정확하게 복구되었고, 로컬에서 최신 메인라인 기준 10/10회 실행 오류, 단일 실행 최대 16384/16384개 요소 불일치. 수정 방법은 이슈의 제안대로 직접 교체가 아니라 이름을 분리하는 것: copy_pipe_to_l1 표기를 PIPE_MTE2로 변경(A2/A3의 GM에서 L1 로드), 새로 copy_pipe_to_l1_bind를 추가하여 PIPE_MTE1로 표기(A5 직결 경로에서 L1 준비를 기다리고 주소를 바인딩), 기존 copy_pipe_to_ub와 copy_pipe_to_ub_V의 선례를 따라 A2/A3를 고치면서 A5가 틀리지 않도록 한다. 동일 근원 문제는 WAR 방향도 포괄한다: MTE1이 아직 읽고 있을 때 MTE2의 TLOAD가 동일한 L1 버퍼를 덮어쓸 수 있다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1981
2.4 Ascend 심사 중: 최초 Parallel 축 브로드캐스트 스칼라 읽기 적용, Chunk O 전방향 이관, MHA 예제 (#1979 / #1866 / #1879)
세 건 모두 11-03 오전에 집중 업데이트되었으며, Ascend 연산자 마일스톤급 배치에 속한다. #1979는 #1885에 해당: 실패군은 하나의 패턴으로 귀결될 수 있다—첫 번째 T.Parallel 축(또는 외부 serial 변수)으로 인덱싱된 1차원 UB 피연산자가 이항 표현식 좌측에 나타나는 경우이며, 이슈의 축소 패턴 dk_i = sum_j v[j] * A[j, i]가 바로 이 형태로, 21점 이분 매트릭스에서 5개의 FAIL이 전부 이 군에 속하며 감축 방향, 브로드캐스트 표기와 무관하다; 근본 원인은 AscendLowerParallelToVector가 2차원 병렬 중첩을 행 루프로 낮출 때, 진입점 HandleSimpleCase가 우측 피연산자에만 스칼라 접근 검사를 수행하고, 좌측은 벡터 경로로 진행하여 UB 정렬 함정에 빠진 것이다; 이번 건은 첫 축의 브로드캐스트 피연산자도 스칼라 읽기를 사용하도록 한다. #1866은 Chunk O 전방향(chunk 기반 게이트 선형 어텐션 O 전방향, GDN)을 Ascend NPU로 이관하며, 개발자 모드에 네 가지 pass 설정(자동 C/V 병합, 자동 C/V 동기화, 자동 동기화, 메모리 계획)을 추가하고, 단일 @tilelang.jit 커널에 20코어 고정 grid-stride 루프를 결합하여 온칩 융합, 제로 GM workspace, 제로 호스트 ATen 연산자를 달성했다. #1879는 MHA 전방향 예제를 추가: BSHD 레이아웃, Expert 모드의 삼분파 MMA 파이프라인으로, wrapper, 공유 레이어, 세 형태 모듈로 나누어 5개의 Python 파일로 구성되며, wrapper에는 T.* 코드가 전혀 없고 메타데이터 검증과 JIT 디스패치만 수행한다; is_causal은 우하단 정렬 마스크를 사용하며, 전체 마스크 행은 sumexp + 1e-30 보호를 통해 출력 0이 된다; cann-bench 20개 테스트 케이스 정밀도 전부 통과, torch-npu 베이스라인 대비 평균 가속비 0.97배, 종합 점수 73.13(로컬 910B2 실측 기준). 0.97배는 아직 순이익이 없다는 의미이지만, 예제로서 “Ascend에서 완전한 어텐션 연산자를 어떻게 작성하는가”의 골격을 제시했다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1979, https://github.com/tile-ai/tilelang-ascend/pull/1866, https://github.com/tile-ai/tilelang-ascend/pull/1879
2.5 mlir-ascend: Codex 워크플로와 작업 인계를 harness 브랜치에 병합 (#208)
11-03 17:50 기준으로 2,715줄이 추가되고 39줄이 삭제되었으며 66개 파일이 변경되었고, 병합 대상 브랜치는 main이 아니라 harness입니다(메인라인의 최신 커밋은 여전히 10-31의 ci: remove internal cache-service references에 머물러 있습니다). 이 점을 인용할 때 “메인라인에 병합”이라고 쓰지 마십시오. 내용은 연산자 harness에 첫 번째 Codex 워크플로를 추가하는 것입니다: 프로젝트 초기화 시 반복 질문을 사용하고, 기존 연산자 저장소를 연결하며, 영속화된 컨텍스트에서 개발, 튜닝 또는 분석을 재개합니다. 세 가지 스킬은 .workspace/ 아래의 편집 가능한 Markdown 기록을 공유하며, 작업 인계 구간은 “현재 후보”와 “검증된 최적 후보”를 구분하고, 실패 및 롤백된 실험을 보존하여 서브 에이전트가 집중된 컨텍스트를 받고 메인 에이전트가 공유 상태를 유지하도록 합니다. 로컬 설치 프로그램은 스킬 등록을 담당하고 프로젝트 항목을 AGENTS.md에 추가합니다(기존 내용 유지, .workspace/ 무시, Git 프로젝트를 worktree 루트로 해석). README, 의존성, 재사용 설명 및 로드맵은 모두 TileLang 선행, Ascend를 범위로 하는 설정을 설명합니다. 워크플로는 GSD를 참고했으며 출처 개정과 MIT 고지를 유지하고, NPU 실행 어댑터와 npm 배포는 여전히 할 일로 남아 있습니다. 이 커밋은 메인 저장소와 Ascend 저장소의 Agent 관련 디렉터리와 함께 동일한 사안의 세 가지 측면입니다: “에이전트가 연산자를 작성하는” 프로세스를 저장소 내의 재사용 가능한 자산으로 정착시키는 것입니다.
출처: https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
2.6 Hygon: Distributed 검토 중 PR 업데이트 있음, 메인 브랜치 신규 커밋 없음 (#19)
tilelang-hygon은 모니터링 기간 내에 푸시 기록이 있지만(11-03 11:06 베이징 시간), 메인 브랜치에는 신규 커밋이 없으며 최신 커밋은 여전히 11-02의 #32입니다(gfx948을 gfx946 호환으로 처리하여 기존 명령을 먼저 테스트 가능하게 함). 푸시는 검토 중인 #19에서 비롯되었으며(제목은 Distributed만 있고 설명은 비어 있음), 11-03 11:06에 업데이트되었습니다. 시간과 제목으로 보아 분산 관련 적응을 진행 중인 것으로 보이지만 PR 설명에 내용이 없으므로 본 보고서는 구현 세부 사항에 대해 추론하지 않습니다. 해당 저장소의 최근 일반적인 리듬은 이틀에서 사흘에 한 번 HCU(Hygon Compute Unit) 측 수정을 하는 것입니다: 10-30의 #31은 MLS 경계와 영속 WASP 동기화를 수정하고, #30은 일반 복사의 LDS swizzle을 보완하며, 10-27의 #29는 scale 버퍼의 bank 분할을 수행합니다.
출처: https://github.com/tile-ai/tilelang-hygon/pull/19
2.7 MetaX, Moore Threads 모니터링 기간 내 푸시 없음 (11-03)
tilelang-musa(Moore Threads)의 마지막 푸시는 09-30 09:50에 멈춰 있고, tilelang-metax(MetaX)는 10-26 07:08에 멈춰 있으며, 두 저장소 모두 이번 모니터링 기간에 커밋이 없고 검토 중인 PR 업데이트도 없습니다. 이번 호의 백엔드 동향은 따라서 Ascend와 Hygon 양측에 집중됩니다: 병합된 8건 중 1건은 Ascend(#3689)에, 1건은 mlir-ascend(#208, harness 브랜치에 병합)에 속하며, 검토 중인 48건의 업데이트 중 Ascend 저장소가 15건, 메인 저장소가 23건을 차지합니다.
출처: https://github.com/tile-ai/tilelang-musa, https://github.com/tile-ai/tilelang-metax
3. 생태계 및 채택 주체 (TileOPs / TileFoundry / TileKernels / FlashQLA)
3.1 TileFoundry: 동기화 계획으로 CUDA lowering 완성 및 역방향 검증, 메인라인에 병합 (#285)
11-04 00:19에 병합되었으며, 788줄 추가, 214줄 삭제, 25개 파일입니다. sync 분석은 공유 메모리 접근에 필요한 커밋, 대기, 펜스를 계획하지만, finalize는 이전에 하나도 발행하지 않았기 때문에 공유 메모리를 경유하는 스케줄러에 경쟁이 존재했고, 카드에서는 sm_89의 E4M3 GEMM이 torch 대비 최대 659 편차로 나타났습니다. 이번 커밋은 분석 단계를 그것들이 선행하는 HIR 노드에 연결하고, SyncMetadata로 레이블별 인덱싱합니다: 노드 자체 문장의 레이블은 비어 있고, InsertSlice가 추가한 복사는 base와 update를 사용하며, 루프 전달 복사는 ("entry", 슬롯)과 ("backedge", 슬롯)을 사용합니다. 핵심 lowering은 어떤 노드의 문장을 lowering하기 전에 새 훅 before(owner, ctx, tag)를 호출하고, CudaLowering이 그곳에서 해당 노드의 단계를 발행합니다. schedule finalize는 sync 분석을 실행합니다. 계약은 네 개의 사양 문서에 기록되었으며, 그중 분석 사양 §1.2.5는 SyncMetadata, 레이블, 규칙 그리고 분석이 여전히 거부하는 상황을 설명합니다. 저자는 위험 요소도 명시했습니다: 코퍼스에서 메인라인에서 finalize할 수 있는 모든 프로그램은 여전히 finalize할 수 있고, 문장만 추가됩니다. 예외는 hir/staged_dynamic_inner로, 그 lane이 “순회 횟수가 lane과 같은” 루프 안에서 동일한 공유 tile을 읽고 쓰며, 이를 합류시킬 펜스가 없습니다. 또한 공유 메모리를 건드리는 예약 호출은 “알 수 없는 바이트에 대한 한 번의 접근”으로 처리되며, 그러한 두 호출이 스레드 그리드가 없는 영역에 있을 때 TIR이 배치할 수 없는 펜스가 필요해 분석이 이들을 거부합니다. 바이트가 서로 겹치지 않더라도 마찬가지입니다. CopyAsyncTensor, CopyReduceAsyncTensor와 wgmma에는 아직 이미터가 없습니다.
출처: https://github.com/tile-ai/TileFoundry/pull/285
3.2 TileFoundry 심사 중: wgmma 동기화 계획, FSDP와 HSDP 학습 의미론이 HIR 평가기에 진입 (#286 / #284)
#286은 11-04 02:28에 업데이트되었으며, 1907줄 추가, 310줄 삭제, 52개 파일로, #285의 직접적인 후속 커밋입니다: WGMMA는 발행 이후에도 공유 메모리와 레지스터를 계속 읽을 수 있으므로, CUDA 동기화 계획은 후속 접근을 이러한 효과 뒤에 배치해야 합니다. 이번 커밋은 각 피연산자 효과에 대해 완료 의미론을 선언하고, WGMMA의 펜스, 커밋, 대기를 그것들이 필요한 HIR 노드에 연결하며, lowering은 명령의 발행 지점에서 이러한 단계를 소비합니다. 후속 스코프가 집합 대기를 담을 수 없으면 발행 스코프에서 해당 연산을 완료합니다. 동시에 레지스터 Reshard 별칭을 수정하여, 요소가 복사본을 보유한 각 스레드에서 동일한 레지스터에 남도록 하며, 이 저장 신원으로 레지스터 의존성을 추적할 계획입니다. 저자는 이 위험을 명확히 기록했습니다: 행렬 명령 이미터는 여전히 사용할 수 없으며, 이번 커밋은 WGMMA의 실행 정확성이나 성능을 확립하지 않고, TMA 완료 의미론은 여전히 별도로 처리됩니다. #284는 11-03 14:34에 업데이트되었으며, 학습 의미론을 다룹니다: 이전 평가기는 reshard를 전역 텐서로 취급하여 서로 다른 참여자의 기여를 놓쳤습니다. 이번 커밋은 참여자 로컬 값, GPU 좌표, 호출과 루프, gather/scatter 그리고 축별 Partial(sum) 해석을 평가하고, Reshard의 별칭 증명에 소유권 의미론을 부여하여 #260과 #170의 구조화된 수명 주기를 확장합니다. 함께 modded-nanogpt에서 영감을 받은 두 블록 GQA 연구를 추가했으며, 명시적 전방/역방향, 모멘텀 SGD와 독립적인 PyTorch 참조를 포함하고, FSDP와 2×2, 4×4 두 가지 HSDP 상주 전략을 다루며, 집합 통신 그룹, 페이로드와 저장을 보고합니다. 수용 기준 수치는 리팩터링 후 핵심 스위트 775개, 백엔드 스위트 778개(두 카드 sm90에서 13개 건너뜀)이며, 열 개의 HIR 구성이 CPU에서 통과하고, 두 전략이 네 카드 sm90에서 실제 2×2 PyTorch HSDP와 두 업데이트 스텝을 실행하여 로컬 그래디언트, 파라미터와 모멘텀이 모두 정렬됩니다.
출처: https://github.com/tile-ai/TileFoundry/pull/286, https://github.com/tile-ai/TileFoundry/pull/284
3.3 TileOPs: 분산 병합이 첫 번째 부분량부터 시작, 상수 입력의 분산이 더 이상 NaN을 내지 않음 (#2821)
11-03 18:04에 병합되었으며, 43줄 추가, 4줄 삭제, 3개 파일입니다. 결함은 분산 병합이 빈 누산기에서 시작하는 것이었습니다: 유한한 큰 평균을 빈 누산기에 병합할 때, 제곱 차이가 0 가중치를 곱하기 전에 이미 오버플로하여 상수 입력의 분산이 NaN이 되었습니다. 영향을 받은 것은 유지 축 앞뒤에 각각 한 단락의 리덕션이 있을 때 var, std, var_mean 호출입니다. 수정은 병합을 첫 번째 실제 부분량에서 초기화하고, 후속 부분량은 기존 Chan 점화식을 따르도록 하는 것입니다. 행 리덕션, 분배, 출력 레이아웃과 공개 API는 모두 변경되지 않습니다. 수용: 세 개의 새 회귀가 원본 커널에서 실패하고 수정 후 통과하며, 후보는 155개 회귀 테스트와 24개 추가 수치 검사를 통과하고, tests/ops/test_reduce.py는 155에서 158로 증가합니다. 세 개의 자동 튜닝 테스트는 제외되었는데, 이전 베이스라인 시도가 60초 테스트 제한을 초과했기 때문입니다. 이러한 “빈 누산기 + 큰 오프셋” 수치 함정은 리덕션 패밀리에서 드물지 않으며, 기록할 만한 점은 유지 축 앞뒤에서 동시에 리덕션하는 형태에서만 드러나고 일반적인 전체 축 리덕션으로는 검출되지 않는다는 것입니다.
출처: https://github.com/tile-ai/TileOPs/pull/2821
3.4 TileOPs 검토 중: GEMM에 선택적 bias 추가, FP8 그룹 GEMM 연속 레이아웃 지원 (#2823 / #2824)
#2823은 GEMMFwdOp.forward에 선택적 bias[N]을 추가했다. bias는 float32 누적 결과에 더해지고 출력은 한 번만 반올림되어 torch.nn.functional.linear와 기준이 일치한다. 각 밀집 구조는 자체 에필로그에서 bias를 추가한다. warp 전용(TMA 및 행 저장), coop2(176/192 폭 및 좁은 tile 포함), coop2_swap, ping-pong, coop2s, simple, swap_ab, GEMV는 각각 cp.async 커널을 가지며, split-K 경로는 리덕션에서 한 번만 추가한다. #2824는 MoEGroupedGEMMFP8FwdOp가 마스크 분할 외에 연속 전문가 행(컴팩트 패킹 또는 정렬 패킹, 세그먼트 끝 오프셋 physical_psum 또는 행별 전문가 id 포함)도 허용하도록 하며, 의미론은 deep_gemm.m_grouped_fp8_gemm_nt_contiguous에 정렬된다. SM90 커널은 세그먼트 끝별로 각 전문가의 실행 세그먼트를 순회하고, 컴팩트한 행별 id는 생산자와 소비자 분할 전에 이진 탐색으로 세그먼트 끝으로 변환되며, 정렬된 행별 id는 행 tile별로 블록 단위로 진행하여 첫 행에서 해당 tile의 전문가를 읽는다. 컴팩트 전문가의 마지막 블록이 block_m 행 미만일 때는 행별로 저장하며 절대 다음 전문가의 행에 쓰지 않고, SM89는 연속 레이아웃을 거부한다. 마스크 호출은 여전히 이전과 동일한 프로그램을 생성하며, 35개의 마스크로 생성된 TIR은 메인라인과 바이트 단위로 일치한다. 벤치마크 측에는 Qwen3-30B-A3B의 EP8 형상 세 줄이 추가되었으며, 실제 라우팅에서 측정된 그룹 크기를 사용한다(해당 모델은 128 전문가, top-8이며, 48개 레이어 중 편향이 중간 정도인 44번째 레이어를 선택하고, EP8에서 행 수가 가장 많은 rank의 16개 전문가를 선택). prefill 행은 4096개 토큰을 라우팅하고(gate/up 및 down), decode 행은 128개 시퀀스의 마지막 토큰을 사용한다.
출처: https://github.com/tile-ai/TileOPs/pull/2823, https://github.com/tile-ai/TileOPs/pull/2824
3.5 TileOPs 검토 중: W4A16 및 밀집 GEMM 선택기를 단일 비용 모델로 재적합 (#2825 / #2826)
두 건 모두 #2133의 하위 작업 2에 속하며, 각각 11-03 20:54와 20:54에 업데이트되었고, 하나는 W4A16을, 다른 하나는 밀집을 담당하며 후자는 문서만 수정한다. #2825는 _select_config를 _legal_configs상에서 _config_cost의 argmin을 취하도록 변경하고, tile 우선, 최대 K 동률 윈도 및 joint_switch_margin 세 가지 예외 규칙을 삭제했다. _config_cost에는 stage_fill 항목이 추가되었다. 더 깊은 링은 각 K tile의 지연을 더 많이 숨길 수 있지만 채우기 자체가 더 느리다. 이 항목이 없으면 K=4096에서 모델은 256 폭 4단계를 512 폭 2단계보다 앞에 배치하지만, 후자가 실측으로 최대 11% 더 빠르다. mma_big, iter_ws, launch_ms는 데이터가 이들을 고정하기에 불충분하여 삭제되었다. 상수는 벤치마크의 SM 클럭에서 바쁜 시간 기준으로 재적합되며, 적합 형상은 목록 밖의 형상을 사용하고 목록 자체는 홀드아웃 검증으로 남겨둔다. _Calibration의 문서 문자열에는 스캔, 적합, 검사 세 단계를 명시하며, 중립 상수에서 다시 실행하면 출하 값을 재현할 수 있다. 커널 측에도 실질적인 메모리 절약이 있다. split-K의 CTA는 전체 K가 아닌 자신의 슬라이스에 해당하는 scale과 zero 그룹만 임시 저장하며, 슬라이스 정렬 시 벡터화 로드를 사용한다. _smem_bytes와 _legal_configs는 split별로 공유 메모리를 계산하지만, 메인라인에서는 각 split의 CTA가 전체 K의 그룹을 임시 저장한다. 또 다른 세부 사항은 _ConfigSpace.smem_bytes가 선택적 상한보다 1 KB 낮은 위치에 있다는 점이다. 정확히 상한에 맞출 때 스캔된 17개 구성이 시작 실패를 겪는데, 할당자가 동적 공유 메모리를 올림하고 배리어는 정적이기 때문이다. #2826은 밀집 선택기의 상수를 어떻게 재적합하는지를 문서 문자열에 작성하며(스캔, 적합, 검사 세 단계, 소프트 어닐링 및 홀드아웃 검증 포함), 코드는 변경하지 않는다.
출처: https://github.com/tile-ai/TileOPs/pull/2825, https://github.com/tile-ai/TileOPs/pull/2826
3.6 TileOPs 검토 중: 짧은 varlen을 KV 그룹별로 패킹, L2 flush를 읽기 전용으로 변경 (#2819 / #2822)
#2819은 11-03 23:37에 업데이트되었으며, #2496의 일부로, SM90의 GQAVarlenFwdOp에 패킹된 warp 전용 프로그램을 추가했습니다. 하나의 요청이 KV 헤드 그룹으로 tile을 비트 단위로 채우는 것이 tile을 더 절약할 때, 키와 값은 KV 헤드별로 한 번만 읽습니다. 패킹 커널의 거부 로직은 비용 모델로 호출의 배치, 패킹 후 query 총수와 그룹 크기를 측정하며, 형태가 보장하는 절감이 실측된 9% 단일 헤드 tile 오버헤드를 초과할 때만 인수합니다. 넓은 항목과 좁은 항목 두 종류가 같은 발사에서 교대로 할당되며, 호스트 재읽기 없이 추가 동기화도 없습니다. 디바이스에서 집계된 분할되지 않은 항목이 SM 수보다 적을 때, 단일 tile 패킹 요청은 키 스캔을 여러 세그먼트로 나누고 fp32 작업 공간을 통해 병합합니다. 창 측에는 두 가지 실측 이익도 있습니다. 창 행은 1%에서 9% 절약, 인터리브된 RoPE 행은 3%에서 6% 절약입니다. 동시에 메인라인의 세 가지 문제를 수정했습니다. sink가 있고 score scale이 약 2.4보다 클 때 가시 키가 없는 행이 NaN을 반환하는 문제(단일 헤드 warp 전용과 FP8 warp 전용 두 프로그램), 백업 query 회전이 요청의 첫 키 이전 행에서 RoPE 테이블을 먼저 읽는 문제, 그리고 거부된 음성 제어 사례가 호출자 프레임을 참조 순환에 남겨 한 테스트의 CUDA graph가 다른 테스트의 캡처 중에 파괴되는 문제입니다. 테스트 노드는 173에서 197로 증가했습니다. #2822는 11-03 19:01에 업데이트되었으며, bench_kernel의 L2 flush를 「int8 버퍼에 zero_() 수행」에서 「보고된 L2의 두 배 크기인 private int32 버퍼 읽기」로 변경했습니다. 버퍼는 고정 시드 CPU 생성기로 한 번에 추출됩니다. CUDA에서는 Triton 커널로 이 읽기를 수행합니다(요소당 한 번 로드, 16 KB 블록마다 한 번 atomic_max). torch 자체 reducer를 사용하지 않는 이유는 이를 flush로 취급하면 다음 Marlin w4a16 호출이 1.3에서 1.5마이크로초 느려지기 때문입니다. 이익은 H200, 1500 MHz, CUPTI 커널 시간(3회 실행 각 30회 호출의 중앙값) 기준: 2048×11008 f32의 Sum이 27.81에서 21.38마이크로초로, 1024×11008 f32의 Ge(torch-compile)가 31.31에서 25.67로, 1024×4096 f32가 6.30에서 5.47로 감소했습니다. 기존 flush는 약 120 MB의 더티 행을 남겨 그 쓰기 되돌림이 다음 타이밍 호출에 들어갔지만, 새 flush는 자체적으로 71에서 445개의 L2 섹터(2에서 14 KB)만 씁니다. 부수적으로 메인라인 버그 하나를 수정했습니다. L2 크기를 알 수 없을 때 로그에 256 MB를 쓰지만 512 MB를 할당하던 문제입니다.
출처: https://github.com/tile-ai/TileOPs/pull/2819, https://github.com/tile-ai/TileOPs/pull/2822
3.7 채택자: FlashQLA와 TileKernels는 푸시 없음, 그러나 #3761은 TileKernels의 fused 커널 수정 (11-04)
QwenLM/FlashQLA의 마지막 푸시는 10-30 03:02에 멈췄고, deepseek-ai/TileKernels는 10-14 05:49에 멈췄으며, 두 저장소 모두 이 모니터링 기간에 커밋이 없어 채택자 측에 새 코드가 없습니다. 다만 이번 기간의 메인라인 헤드라인은 채택자와 직접 관련이 있습니다. #3761이 수정한 성능 회귀의 발동 지점이 바로 TileKernels의 mhc_pre_big_fuse 커널이었습니다. 동일한 커널이 B300, 은닉 차원 1280에서 배리어 점유 예산의 25%를 잃었습니다. 이는 채택자의 커널이 이미 컴파일러 측의 회귀 탐침이 되었음을 보여줍니다. TileKernels가 코드를 푸시하지 않은 이 몇 주 동안에도 그 커널은 메인라인 변경을 통해 수동적으로 「테스트되고」 있습니다. 생태계를 추적하는 사람에게 이는 주목할 만한 역방향 신호입니다. 채택자의 침묵 기간에도 컴파일러 변경은 그들의 형태에서 문제를 일으킬 수 있습니다.
출처: https://github.com/QwenLM/FlashQLA, https://github.com/deepseek-ai/TileKernels, https://github.com/tile-ai/tilelang/pull/3761
4. 커뮤니티, 튜토리얼 및 활동
4.1 하루 7건의 새 issue: 프로덕션급 CUDA graph 문제 하나, 동일 외부 보고자로부터 네 건 (11-03)
모니터링 기간 동안 메인 저장소에 7개의 issue가 새로 추가되었습니다(번호 #3762부터 #3775까지, 그리고 Ascend 저장소의 자동 일일 테스트 보고서 하나). 기술적 비중이 가장 높은 것은 #3775로, 보고자는 프로덕션 워크로드를 운용하는 팀입니다. 환경은 pip로 설치한 tilelang 0.1.12, Python 3.12.3, H200 두 장에 NVLink, 드라이버 580.178.04, CUDA 13 컨테이너, torch 2.14.1+cu130이며, 서비스 측은 단일 프로세스 TP2(장치마다 워커 스레드 하나, 각자 CUDA 컨텍스트 하나) 구성의 sglang이고, 모델은 GLM 계열 MoE에 DSA 어텐션을 더한 것입니다. 경로상의 TileLang 커널은 mhc_pre_big_fuse_with_norm, mhc_post 그리고 DSA prefill 커널입니다. 실행 백엔드를 TILELANG_EXECUTION_BACKEND=nvrtc로 강제한 뒤(기본값 auto는 tvm_ffi로 해석됨), tilelang의 JIT 커널이 CUDA graph 캡처 중에 실행되면 캡처가 손상되며, 그 손상이 tilelang 실행 이후의 다운스트림 무관 커널에서 나타납니다. 그리고 바이트 단위로 동일한 설정을 세 번 재현했을 때 첫 실패 지점이 세 번 모두 달랐습니다. 한 번은 cudaError700(유효하지 않은 TMA 디스크립터 핸들)이 TMA 디스크립터 초기화 지점에서 발생한 뒤 무관한 실행 지점에서 울림, 한 번은 cublasGemmEx가 CUBLAS_STATUS_EXECUTION_FAILED를 보고(순수 BF16 입출력), 한 번은 호스트 프레임워크가 JIT로 생성한 원소별 커널에서의 잘못된 메모리 접근이었습니다. 세 번 모두 decode의 CUDA graph 캡처 중(배치 크기 버킷 8)에 발생했고, 모두 서비스 시작 후 약 100초 시점이었으며, /proc/environ과 깨끗한 대조군의 차이는 오직 그 환경 변수 하나뿐이었습니다. 이는 가장 위치를 찾기 어려운 유형에 속합니다. 현상은 무작위적이고, 실패 지점은 표류하며, 유인은 환경 변수 하나뿐이기 때문입니다. 보고자는 세 번의 실행에 대한 대조 조건을 빠짐없이 기술했습니다.
나머지 네 건은 같은 외부 보고자로부터 나왔으며, 환경은 일괄적으로 메인 브랜치 또는 0.1.15 pip 패키지, WSL 하의 Linux, Python 3.13.5, RTX 4060 Laptop(sm_89), 드라이버 566.26, CUDA 12.6 런타임이고, 모두 재현 스크립트와 전체 출력을 첨부했습니다. #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 설치’라는 조합의 집중적인 장애 보고 면을 구성합니다. 남은 두 건은 다음과 같습니다. #3769는 B300에서 tcgen05의 int8 GEMM이 ptxas에 의해 거부된 문제(이미 심사 중인 #3770 수정이 있음)이고, #3762는 Ascend 950에서 T.simd.vcvt가 sat=False일 때도 여전히 포화되는 문제로, 자세한 내용은 다음 항목을 참조하십시오.
출처: https://github.com/tile-ai/tilelang/issues/3775, https://github.com/tile-ai/tilelang/issues/3774, https://github.com/tile-ai/tilelang/issues/3773
4.2 Ascend 950의 vcvt 포화 시맨틱과 mlir-ascend 일일 테스트(#3762 / #1980)
#3762는 Ascend 측 사용자가 제출했으며, 환경은 Ascend950PR(NpuArch 3510 / dav-3510), CANN 9.2.0-beta.2, 드라이버 25.7.rc1, torch_npu 2.12.0, 소스에서 설치한 TileLang(10-21의 e84ce3de)입니다. 문제는 T.simd.vcvt(src, "uint8", sat=False)가 int32 벡터에 대해 sat=False를 조용히 무시하고 항상 포화된다는 것입니다. 음수 값은 0이 되고, 255를 넘는 값은 255가 되며, 오류도 경고도 없습니다. 반면 호출자는 자신이 랩어라운드 시맨틱을 쓰고 있다고 생각합니다. 문서 문자열은 이런 상황을 명확히 기술합니다. sat=False일 때 범위를 벗어난 값은 랩어라운드된다는 것입니다. 기대값은 -1이 255, -128이 128, 300이 44가 되는 것이지만, 실측값은 -1이 0, -128이 0, 300이 255가 됩니다. 영향은 int8 양자화 마무리 단계의 f32에서 int8로 변환하는 경로에 떨어집니다. 음의 중간값이 한 번이라도 나타나면 시맨틱이 랩어라운드에서 0으로의 절단으로 바뀝니다. #1980은 Ascend 저장소의 자동 일일 테스트로, 11-03 09:35 라운드의 3030개 항목이 전부 통과하여 통과율 100%였고, 워크플로 실행 번호는 41553942797입니다. 이 두 숫자를 나란히 보면 흥미롭습니다. Ascend 저장소의 일일 테스트는 전부 초록색이지만, 일일 테스트가 커버하지 못하는 시맨틱 차이(예를 들어 sat 파라미터의 실제 동작)는 여전히 사람의 보고에 의존해야 합니다.
출처: https://github.com/tile-ai/tilelang/issues/3762, https://github.com/tile-ai/tilelang-ascend/issues/1980
4.3 문서 사이트 세 차례 업데이트, VS Code 확장 이름 변경 vsix (11-03)
문서 사이트 tilelang.github.io는 모니터링 기간 내에 세 차례의 Update docs 커밋이 있었습니다(11-03 13:11, 13:28, 11-04 00:59 베이징 시간). 이는 메인라인 병합에 따라 자동 빌드되는 정기 업데이트로, 내용은 이번 모니터링 기간에 병합된 몇 건—#line 계층 매핑, ThreadSync 스코프 목록, Ascend 일괄 복사의 문서 해당 위치—에 대응합니다. TileOPs.github.io는 이번 모니터링 기간에 푸시가 없었고, 최신은 여전히 11-01의 “TileOPs 메인라인 따르기: 중첩과 정렬 보장, 빌드 설명”입니다. VS Code 확장 저장소는 11-04 02:19에 “문서 업데이트 및 vsix로 이름 변경”이라는 커밋이 있었으며, 11-02의 “CUDA 디버깅을 TileLang 행 매핑에 정렬”에 이어집니다: 지난 기간 메인라인에 병합된 행별 앵커링이 생성 코드의 행 번호 출처를 바꿨기 때문에 확장 측 디버깅 설정도 함께 바뀌어야 했고, 이 커밋이 산출물 이름과 문서를 함께 조정했습니다.
출처: https://github.com/tile-ai/tilelang.github.io, https://github.com/tile-ai/tilelang-vscode-edition
4.4 뉴스 측: 세 채널 모두 모니터링 기간 내 적중 없음, 외신의 Ascend 훈련 보도 엇갈려 (11-03)
Google News의 중영문 조합 검색, Hacker News, arXiv 세 채널 모두 이번 모니터링 기간에 TileLang 관련 실질적 적중이 없었습니다: 후보 목록 15건 중 12건은 GitHub 저장소 발견에서 나왔고, 2건은 주제와 무관한 Show HN(탐정 게임과 하드웨어 설계 프롬프트 도구), 1건은 Wccftech의 DeepSeek와 Ascend 관련 보도였습니다. arXiv의 all:tilelang 최신 적중은 여전히 2026-07-24의 TileSight(tile 중심 GPU 성능 분석 모델)이며, TileLang과의 결합 검색에서 모니터링 기간 내 새 논문은 없었습니다. 기록할 만한 유일한 업계 측 소식은 보도 간 충돌입니다: Wccftech는 DeepSeek가 Ascend 950DT로 구성된 Atlas 950 SuperPoD에서 V5를 처음부터 훈련할 것이라고 보도한 반면, 같은 시기 The Information의 보도는 “이르면 2026년 말부터 950DT에서 사전 훈련 시작”이었고, Bloomberg와 파이낸셜타임스의 초기 보도는 950DT의 16만 장 주문이 주로 추론용이며, 이전에 Ascend에서 훈련이 좌절된 후 Nvidia로 회귀했다고 전했습니다. 세 가지 설명은 “훈련에 사용되는지, 그리고 언제 사용되는지”에 대해 일치하지 않으며, 모두 Huawei나 DeepSeek 공식 확인을 받지 못했습니다. 여기에 둔 것은 배경 설명일 뿐입니다: 훈련이 실제로 이뤄지든 아니든, Ascend 측의 컴파일러와 연산자 투자(이번 기간 #3689, #1975, #1981)는 모두 이 경로를 위한 토대를 닦는 것입니다.
출처: https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods, https://arxiv.org/
5. 트렌드 관찰
첫째, 성능 회귀가 일급 시민으로 취급되기 시작했으며, 수정 방식이 결론보다 방법론적으로 더 기록할 가치가 있습니다. #3643이 정확성 수정을 위해 도입한 ID 보존 규칙은 그 자체로 필요하고 올바릅니다; #3761은 이를 롤백하지 않고 보존과 재사용이 서로 다른 계층에서 공존하도록 했습니다—스코프를 넘어 ID를 공유하기 전에 참여 술어가 항상 성립함을 먼저 증명하고, 증명할 수 없으면 스코프가 있는 카운팅으로 회귀합니다. 작성자의 “선언된 배리어 리소스 수만 변경”이라는 격리 실험과 함께, 이 수정 체인은 “정확성”과 “리소스 선언” 두 가지를 분리했습니다: 후자는 단독으로 재현하고 단독으로 검수할 수 있는 변수입니다. 이러한 분리 능력이야말로 컴파일러 프로젝트가 “돌아가는” 단계에서 “유지보수 가능한” 단계로 나아가는 표지입니다.
둘째, Ascend의 엔지니어링이 두 축에서 동시에 추진되고 있습니다. 하나는 아키텍처입니다: #1975는 2887개 파일, 60만 줄 추가로 TileLang-Ascend 전체를 메인라인 멀티 백엔드 프레임워크로 옮겼고, arch="dav-2201" 같은 식별자로 백엔드 분기를 수행하여 A2A3와 A5가 같은 저장소에서 직교적으로 공존하는 것을 목표로 합니다; 다른 하나는 파이프라인입니다: mlir-ascend의 #208은 Codex의 초기화, 연결, 복구 세 가지 워크플로와 작업 인계를 .workspace/에 기록했고, 메인 저장소와 Ascend 저장소도 각각 Agent 디렉터리를 가지고 있습니다. “저장소 마이그레이션”과 “에이전트로 저장소 마이그레이션”을 같은 날 함께 보면, Ascend 측의 병목이 이미 “연산자 부족”에서 “일관된 엔지니어링 골격과 재사용 가능한 마이그레이션 프로세스 부족”으로 옮겨갔음을 알 수 있습니다.
셋째, 연산자 라이브러리의 자기 증명 문화가 심화되고 있다. TileOPs 이번 기간 심사 중인 몇 건은 거의 모두 ‘증거’를 PR에 적어 두었다: 실제 라우팅으로 측정한 Qwen3-30B-A3B EP8 그룹 크기(48개 레이어 중 편향이 중앙에 가까운 44번째 레이어 선택), 35개 마스크로 구축한 TIR이 메인라인과 바이트 단위로 일치, L2 flush 변경에 46줄 구현 대조표를 첨부하고 기존 flush가 120 MB의 더티 라인을 남긴다고 설명, W4A16의 비용 모델에는 “어떤 항목이 데이터로 고정되지 않아 삭제되었는지”를 명시적으로 기재. TileFoundry는 다른 경로를 택했다—동기화를 증명 가능한 분석으로 만들고, lowering이 역으로 검증하게 하며, 분석이 거부한 상황을 직접 나열한다(TMA와 wgmma의 완료 의미 미선언, 스레드 경계를 넘어 동일하지 않은 루프). 두 경로를 합치면 같은 지향점을 보여준다: 차라리 “증명하지 못했다”를 인정할지언정 “돌아갔다”를 “맞다”로 여기지 않는다.
넷째, 리스크 면이 설치 형태로 집중되고 있다. 하루 7건의 새 issue 중 네 건은 pip 패키지를 사용하고 WSL에서 실행하며 소비자용 sm_89 카드를 쓰는 동일한 외부 보고자로부터 나왔고, 모두 재현 가능한 스크립트를 첨부했다. #3775는 프로덕션 환경의 “pip 패키지 + 다중 CUDA 컨텍스트 + nvrtc 백엔드 + CUDA graph 캡처” 조합이다. 이는 컴파일러의 심층 결함은 아니지만 신규 사용자가 처음 시도할 때의 경험을 결정한다. Ascend 저장소의 일일 테스트 3030건 전부 통과에도 sat=False인 의미 편차가 여전히 수동 보고에 의존한다는 점과 대조하면, 테스트 커버리지의 형태와 실제 사용 형태 사이에는 아직 거리가 있음을 알 수 있다.
마지막 관찰: 이번 기간 뉴스 측은 여전히 공백이며, 15개 후보 중 12개가 저장소 발견이다. TileLang의 대외적 입지는 여전히 “DeepSeek과 Ascend” 같은 이벤트에 묶여 있고, 일상적 진전은 커밋에서만 일어난다. 이 프로젝트를 추적하는 사람에게 이는 데일리 리포트의 본체가 보도가 아니라 PR과 issue일 수밖에 없음을 의미한다—이는 이미 여러 기간 연속된 안정적 형태다.
부록, 자료와 검증 설명
출처 검증표
| 출처 | 검증 결과 |
|---|---|
| tile-ai 조직(29개 저장소) | 모니터링 기간 내 9개 저장소 푸시 |
| 메인 저장소 tilelang | 메인 브랜치 5건 머지, 신규 issue 6건, 심사 중 업데이트 23건, 문서 사이트 3건 자동 빌드 |
| TileOPs | 메인 브랜치 1건 머지, 심사 중 업데이트 7건 |
| TileFoundry | 메인 브랜치 1건 머지, 심사 중 업데이트 2건 |
| tilelang-ascend | 기본 브랜치 ascendc_pto 모니터링 기간 내 신규 커밋 없음, 심사 중 업데이트 15건, 일일 테스트 3030건 전부 통과 |
| tilelang-mlir-ascend | harness 브랜치에 1건 머지(비주요 라인) |
| tilelang-hygon | 메인 브랜치 신규 커밋 없음, 심사 중 #19 업데이트 있음 |
| tilelang-vscode-edition | 1건 머지(문서 및 vsix 명칭 변경) |
| tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / tilelang-puzzles | 모니터링 기간 내 푸시 없음 |
| 릴리스 페이지(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) | 모니터링 기간 내 새 버전 없음, tilelang 최신은 여전히 09-30의 v0.1.15 |
| FlashQLA / TileKernels | 모니터링 기간 내 푸시 없음 |
| Google News(중영문 6개 쿼리) / Hacker News / arXiv | 모니터링 기간 내 실질적 적중 없음 |
전체 출처 목록
- [1] tilelang ThreadSync, 여러 저장소 영역에 걸친 부분 배리어 재사용 (#3761) — https://github.com/tile-ai/tilelang/pull/3761
- [2] tilelang, 레거시 TVMScript AST 및 파서 제거 (#3736) — https://github.com/tile-ai/tilelang/pull/3736
- [3] tilelang, 자동 튜닝 입력 공급 디바이스 (#3546) — https://github.com/tile-ai/tilelang/pull/3546
- [4] tilelang, 조건부 cast에서 벡터 폭 유지 (#3759) — https://github.com/tile-ai/tilelang/pull/3759
- [5] tilelang, warp 정렬 스레드 경계의 부분 배리어 (#3765) — https://github.com/tile-ai/tilelang/pull/3765
- [6] tilelang, 다중 인덱스 fragment 접근 오류 보고 (#3767) — https://github.com/tile-ai/tilelang/pull/3767
- [7] tilelang, padded 공유 범위를 통째로 복사할 수 없을 때 폴백 (#3738) — https://github.com/tile-ai/tilelang/pull/3738
- [8] tilelang, 2-CTA TMA 바이트를 피어 leader 기준으로 집계 (#3766) — https://github.com/tile-ai/tilelang/pull/3766
- [9] tilelang, tcgen05 kind::i8을 해당 기능이 없는 타깃에서 사용하지 않음 (#3770) — https://github.com/tile-ai/tilelang/pull/3770
- [10] tilelang CuTeDSL, sm_103a 및 sm_107a 지원 (#3772) — https://github.com/tile-ai/tilelang/pull/3772
- [11] tilelang SM100 flash attention 역방향 예제 (#3768) — https://github.com/tile-ai/tilelang/pull/3768
- [12] tilelang CuTeDSL 예제 누락분 보완 (#3685) — https://github.com/tile-ai/tilelang/pull/3685
- [13] tilelang nvrtc 백엔드가 CUDA graph 캡처를 손상시킴 (#3775) — https://github.com/tile-ai/tilelang/issues/3775
- [14] tilelang 공유 메모리 치환에서 열 오프셋 유실 (#3774) — https://github.com/tile-ai/tilelang/issues/3774
- [15] tilelang while 루프 내 공유 메모리 GEMM 불안정 (#3773) — https://github.com/tile-ai/tilelang/issues/3773
- [16] tilelang 중첩 int16 내림 Let 충돌 (#3763) — https://github.com/tile-ai/tilelang/issues/3763
- [17] tilelang B300에서 tcgen05 int8이 ptxas에 의해 거부됨 (#3769) — https://github.com/tile-ai/tilelang/issues/3769
- [18] tilelang Ascend 950에서 vcvt가 sat=False를 무시함 (#3762) — https://github.com/tile-ai/tilelang/issues/3762
- [19] tilelang-ascend 3D GM에서 L1로의 배치 ND2NZ 복사 (#3689) — https://github.com/tile-ai/tilelang-ascend/pull/3689
- [20] tilelang-ascend 전체 저장소를 메인라인 멀티 백엔드 아키텍처로 이전 (#1975) — https://github.com/tile-ai/tilelang-ascend/pull/1975
- [21] tilelang-ascend copy_pipe_to_l1 동기화 어노테이션 분리 (#1981) — https://github.com/tile-ai/tilelang-ascend/pull/1981
-
[22] tilelang-ascend 첫 Parallel 축 브로드캐스트를 스칼라 읽기로 처리 (#1979) — https://github.com/tile-ai/tilelang-ascend/pull/1979
- [23] tilelang-ascend Chunk O 순방향 마이그레이션(#1866) — https://github.com/tile-ai/tilelang-ascend/pull/1866
- [24] tilelang-ascend MHA 순방향 예제(#1879) — https://github.com/tile-ai/tilelang-ascend/pull/1879
- [25] tilelang-ascend 데일리 테스트 3030건 전부 통과(#1980) — https://github.com/tile-ai/tilelang-ascend/issues/1980
- [26] tilelang-mlir-ascend Codex 워크플로 및 작업 인계(#208) — https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
- [27] tilelang-hygon 분산 검토 중 PR(#19) — https://github.com/tile-ai/tilelang-hygon/pull/19
- [28] TileFoundry CUDA lowering 발사 동기화 계획(#285) — https://github.com/tile-ai/TileFoundry/pull/285
- [29] TileFoundry wgmma 동기화 계획(#286) — https://github.com/tile-ai/TileFoundry/pull/286
- [30] TileFoundry FSDP 및 HSDP 학습 의미론(#284) — https://github.com/tile-ai/TileFoundry/pull/284
- [31] TileOPs 분산 병합을 첫 번째 부분량부터 계산(#2821) — https://github.com/tile-ai/TileOPs/pull/2821
- [32] TileOPs GEMM 선택적 bias(#2823) — https://github.com/tile-ai/TileOPs/pull/2823
- [33] TileOPs FP8 그룹 GEMM 연속 레이아웃(#2824) — https://github.com/tile-ai/TileOPs/pull/2824
- [34] TileOPs W4A16 선택기 비용 모델(#2825) — https://github.com/tile-ai/TileOPs/pull/2825
- [35] TileOPs 밀집 GEMM 선택기 재적합 문서(#2826) — https://github.com/tile-ai/TileOPs/pull/2826
- [36] TileOPs 짧은 varlen을 KV 그룹별로 패킹(#2819) — https://github.com/tile-ai/TileOPs/pull/2819
- [37] TileOPs L2 flush를 읽기 전용으로 변경(#2822) — https://github.com/tile-ai/TileOPs/pull/2822
- [38] Wccftech의 DeepSeek 및 Ascend 950DT 관련 보도 — https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods