모니터링 기간: 2026-10-30 07:00 ~ 2026-11-02 07:00 (베이징 시간, 72시간, 주말 포함; 이전 회차 기준 시점은 10-30 07:00).

출처: GitHub tile-ai 조직 저장소별 점검(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / TileOPs-nightly / 문서 사이트) + 머지 및 심사 중 실값 검색(머지 35건, 심사 중 업데이트 82건) + 신규 issue 전수 확인(11건) + 메인 브랜치 커밋별 점검(tilelang 7, TileOPs 15, TileFoundry 3, tilelang-ascend 5, tilelang-hygon 2, tilelang-mlir-ascend 1, 문서 사이트 4) + 릴리스 페이지 및 tag 점검(6개 저장소) + 서드파티 생태계 저장소 점검(FlashQLA, TileKernels, RISC-V 연구 체인 3개 저장소, PTO-ISA/tilelangGym) + Google News 중영문 검색(프록시 경유) + Hacker News + arXiv.

이번 회차 인덱스

  • 오늘의 핵심: 메인라인은 TIRX 마이그레이션을 위해 ThreadSync의 정확한 스레드 카운트를 되돌리는 것(#3733)이며, 동시에 정확성 개선도 함께 철회함을 인정하는 추적 이슈를 새로 열었습니다. TileOPs는 하루에 13건의 성능 병합을 기록했습니다. Ascend 호스트 고정 오버헤드가 약 450마이크로초에서 약 106마이크로초로 줄었습니다. 72시간 동안 35건 병합, 82건 리뷰 중 업데이트, 신규 issue 11건
    1. 핵심 프로젝트 진행 상황
      • 1.1 메인라인: ThreadSync의 정확한 스레드 카운트 롤백, TIRX 마이그레이션을 위한 양보(#3733)(11-01)
      • 1.2 메인라인: CUDA 크로스 warp 전체 리덕션을 단일 공유 메모리 교환으로 병합, 256 스레드 기준 6개 배리어에서 2개로 감소(#3715)(10-30)
      • 1.3 메인라인: PTO 백엔드의 T.Parallel 벡터화 첫 마일스톤 착수, 원소별 연산자는 VMI 명령 사용(#3512)(10-30)
      • 1.4 메인라인: Ascend DeepGEMM의 scale 버퍼를 쓰기 카운트 다중 버퍼로 변경, 데이터 로드를 조건부 로드 이전으로 이동 가능(#3629)(10-30)
      • 1.5 메인라인: 두 건의 수정 병합, CPU 스칼라 GEMM은 대상 영역만 초기화, 음수 인덱스 합법화는 도달 불가 분기 건너뜀(#3709 / #3712)(10-30)
      • 1.6 메인라인: 언어 계층에서 proxy 내 미사용 Ref 플레이스홀더 두 개 제거(#3722)(10-31)
      • 1.7 신규 결함 및 리뷰 중: TMEM 결과 조용히 유실, ROCm 심볼 미발견, TF32 최근접 짝수 반올림(#3721 / #3711 / #3714)(10-30 ~ 10-31)
    1. 멀티 백엔드 적응
      • 2.1 Ascend: Cython wrapper에서 TVM Analyzer 지연 생성, 910B3 단일 호출 약 450마이크로초에서 약 106마이크로초로 감소(#1948)(10-30)
      • 2.2 Ascend: RFC에서 동기화 구현을 메인라인 코어 + 독립 Z3 증명 서비스로 교체 제안(#1971)(10-30)
      • 2.3 Ascend 리뷰 중: TileLang-Ascend 전체를 메인라인 멀티 백엔드 아키텍처로 이전(#1975)(10-31)
      • 2.4 Ascend 리뷰 중: GLM-5.3 kpool의 fp8_mqa_logits 예제, 듀얼 kernel + AIV/Cube 분업(#1966)(10-31)
      • 2.5 Ascend: 문서 및 API 사실 검증 4건 병합, PTO에서 clamp 실효화(#1899 / #1893 / #1620 / #1600)(10-30)
      • 2.6 Ascend: 일일 테스트 2970개 케이스 전원 통과, 11-01 workflow 1회 실패(#1977 / #1976)(10-31 ~ 11-02)
      • 2.7 Hygon: MLS 경계 및 지속 WASP 동기화 수정 병합, gfx948 호환 리뷰 중(#31 / #32)(10-30)
      • 2.8 mlir-ascend: CI에서 내부 캐시 서비스 참조 제거, Codex 워크플로 리뷰 중(#206 / #208)(10-31)
      • 2.9 Moore Threads와 MetaX는 모니터링 기간 내 푸시 없음(11-02)
    1. 생태계 및 채택 측
      • 3.1 TileOPs: 72시간 동안 15건 병합, 13건이 11-01에 집중, 거의 모든 건에 실측 표와 체크리스트 회귀 기준 포함(11-01)
      • 3.2 TileOPs: GQA 밀집 순방향에서 실측 규칙으로 decode 또는 prefill 프로그램 선택, 선택 리그렛 1.818에서 1.017로 감소(#2808)(11-01)
      • 3.3 TileOPs: 페이지드 MHA decode를 공유 메모리에 들어가는 모든 head dim으로 확장, GQA 페이지드 순방향에서 구 decode 커널 제거(#2807 / #2797)(11-01)
      • 3.4 TileOPs: coop2 프로듀서를 단일 warp로 축소, 192 폭 전치 블록 추가, W4A8 부분 1웨이브를 1 CTA당 1 tile로 변경(#2806 / #2805)(11-01)
      • 3.5 TileOPs: 컴파일 캐시를 빌드 타깃 기준으로 키잉, 연산자 계층에서 출력과 입력 겹침 통합 검사(#2802 / #2801)(11-01)
      • 3.6 TileOPs 수정 중 결함: Norm 4개 연산자가 0에서 먼 행에서 정밀도 손실, 약 1e36에서 오버플로되어 NaN 발생(#2812)(11-01)
      • 3.7 TileFoundry: lowering에서 issuing mesh와 윈도 분할을 한 곳에서 결정, 83개 스케줄 픽스처 바이트 단위 일치(#280 / #282)(10-31 ~ 11-01)
      • 3.8 TileFoundry: 사이트에 0.3 버전 스케줄링 글 《Scheduling, Designed for Agents》 게시(#4 / #281)(10-31 ~ 11-01)
      • 3.9 채택 측: FlashQLA에 SM120 Dual-CP 순/역방향 병합, RTX 5090 2장에서 32K~128K 시퀀스로 Inter-CP와 비교(#51)(10-30)
      • 3.10 채택 측: TileKernels는 10-14 이후 푸시 없음(11-02)
    1. 커뮤니티, 튜토리얼 및 이벤트
      • 4.1 RISC-V 연구 체인 형성: 세 개 저장소에서 Numba CUDA Python을 TileLang으로 변환하고, 제한된 PrimFunc를 단일 hart RVV C로 로워링(10-31)
      • 4.2 PTO-ISA/til

elangGym이 PTO-ISA/tilelang에서 독립 저장소로 분리되었으며, 메인 회귀 게이트는 Ascend SimtVF 경로입니다 (10-30)

  • 4.3 뉴스 측면: Google News 중·영문 쿼리, Hacker News, arXiv 모두 모니터링 기간 내 실질적 적중 없음 (11-02)
    1. 동향 관찰
  • 부록. 자료 및 검증 설명

오늘의 핵심

이번 모니터링 기간의 가장 큰 움직임은 덜어내기입니다. 메인라인은 3주 전에 병합된 정밀 스레드 카운트(#3291)를 전량 롤백했습니다(#3733). 이유는 매우 직설적으로 적혀 있습니다. 그 변경은 비용이 큰 카운트 쿼리를 풀 블록 배리어 검사 이전으로 옮겼고, 모든 스레드가 배리어에 도달하더라도 증명을 실행해야 하며, 현재 TVM의 솔버 예산이 소진되어 부분 카운트를 반환하므로, 본래 합법적인 remap 및 sequence 보조 kernel이 「Cannot prove a thread rectangle」로 컴파일 실패하게 됩니다. 롤백은 두 가지 ThreadSync 구현 모두에서 동시에 진행되었으며, #3291이 도입한 테스트를 삭제하고, 「루프가 대부분 스레드에 대해 비어 있으나 256 스레드가 여전히 전부 배리어에 도달하는」 리덕션 케이스를 추가했습니다. 상 가드, 정밀 참여 스레드 경계, 비직사각형 스레드 집합 등 정확성 개선도 함께 철회되었으며, 새로 열린 #3732는 마이그레이션 후 재최적화 추적 티켓으로 남겨두었습니다. 이는 명시적으로 임시 조치로 표기된 마이그레이션 양보입니다. TileLang은 현재 TVM fork 의존성에서 벗어나기 위한 TIRX 마이그레이션을 준비 중이며, 이에 수반되는 TVM fork 최적화 제안은 이미 종료되었습니다.

두 번째 축은 Ascend가 기반을 보강하는 것입니다. #1948은 호스트 측의 연산자 규모와 무관한 고정 오버헤드를 수정했습니다. Cython wrapper가 매 호출마다 무조건 TVM Analyzer를 생성하는데, 실측 결과 910B3에서 약 0.28밀리초가 소요되며, 이는 형상이 기호 표현식일 때만 사용됩니다. 지연 생성으로 변경한 후 vec add 단일 블록은 442~467마이크로초에서 105.1마이크로초로, 32 블록은 441~493마이크로초에서 108.0마이크로초로, kpool 압축은 447~450마이크로초에서 111.3마이크로초로 감소했습니다. 하한선인 순수 ctypes 호출은 약 12마이크로초로 변동이 없었습니다. 같은 시기 #1971의 RFC는 동기화 구현의 대체 경로를 제시했습니다. 메인라인 동기화 코어를 이식하고 독립적인 Z3 증명 서비스를 추가하며, 구 TIR 어댑터 계층으로 기존 컴파일 파이프라인에 연결하고, TVM, 스케줄링, 메모리 플래닝, AscendC/PTO codegen 및 런타임 ABI를 유지하며, 검증 통과 후 기본 구현으로 전환하고 구 알고리즘을 삭제하는 것입니다. 이는 #1975에서 전체 Ascend 저장소를 메인라인 멀티 백엔드 아키텍처로 이전하는 것과 동일한 사안의 두 가지 착지점입니다.

세 번째 축은 연산자 라이브러리입니다. TileOPs는 11-01 하루에 13건을 병합했으며, 거의 전부가 성능 패치이고 각 건마다 H200 실측 표와 manifest 회귀 임계값이 첨부되어 있습니다. 그중 #2808은 GQA dense forward의 프로그램 선택을 하나의 규칙으로 통합하여, 최고 속도 후보 프로그램 대비 선택 후회를 1.818에서 1.017로 낮추었으며(2~128개 query 위치, 586회 호출), 메인라인의 실제 버그 하나를 수정했습니다. score scale이 log2(e)를 곱한 후 FP32에서 소실되어 출력이 NaN 또는 잘못된 값이 되는 문제입니다.

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

1.1 메인라인: ThreadSync의 정밀 스레드 카운트 롤백, TIRX 마이그레이션을 위한 양보 (#3733)

11-01 14:50에 병합. 롤백된 #3291은 정밀 스레드 카운트를 풀 블록 배리어 검사 이전으로 옮기고 직사각형 인증을 추가했으며, 그 대가로 모든 스레드가 배리어에 도달하더라도 CountSatisfyingValues를 한 번 실행해야 했습니다. TVM 측 솔버 예산이 소진된 후 부분 카운트를 반환하면서, remap 및 sequence 보조 kernel이 즉시 「[ThreadSync] Cannot prove a thread rectangle」로 실패했습니다. 회귀는 메인라인 cc903dcb9에 변경되지 않은 TVM pin 4211874e9e를 더한 환경에서 재현되었으며, 새 케이스는 이를 「루프가 대부분 스레드에 대해 비어 있으나 256 스레드가 여전히 전부 배리어에 도달하는」 형태로 축소했습니다. 롤백은 공유 및 Ascend 두 구현을 모두 포괄하며, 풀 블록 빠른 경로와 기존의 부분 카운트 동작을 복원합니다. 원문은 이것이 TIRX 마이그레이션 기간의 임시 완화임을 명확히 밝히고 있으며, TVM fork 측 최적화 제안 tile-ai/tvm#89는 종료되었고 후속 최적화는 #3732로 이전되었습니다.

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

1.2 메인라인: CUDA 크로스 warp 전체 리덕션을 단일 공유 메모리 교환으로 통합, 256 스레드가 6개 배리어에서 2개로 감소 (#3715)

10-30 20:26에 병합. 기존 구현은 CUDA 리덕션 헤더 파일 reduce.h에서 threads/2부터 scale까지 재귀하며, offset이 32 이상인 모든 단계마다 전체 공유 메모리 쓰기-읽기 교환을 수행하고 배리어 2개를 배치했습니다. 256 스레드 블록의 세 단계(128, 64, 32)는 총 3회 공유 메모리 왕복, 6개 배리어였고, 512 스레드 블록은 4회, 8개였습니다. 변경 방법은 모든 크로스 warp 단계를 단일 교환으로 압축하는 것입니다(요소당 st.shared 1회, 총 2개 배리어). 먼저 warp 내에서 scale이 32 미만인 구간을 수행하고, 그 다음 각 warp 내 그룹에서 cross_radix 이하의 여러 lane이 협력하여 폴딩하며, 레지스터 압력은 상한이 있습니다. 스칼라 및 배치 두 경로 모두 변경되었습니다.

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

1.3 메인라인: PTO 백엔드의 T.Parallel 벡터화 첫 마일스톤 달성, 원소별 연산자가 VMI 명령 사용 (#3512)

10-30 01:29에 병합되었습니다. 목표는 사용자가 T.SimdVF(lanes=L) 안에서 스칼라 병렬의 원소별 T.Parallel을 작성하면, 컴파일러가 이를 PTO VMI 명령으로 벡터화하고, 다시 PTOAS를 거쳐 Ascend 950에서 실행되도록 하는 것입니다. 역할 분담은 명확하게 기술되어 있습니다. TileLang 측은 논리 lane 레이아웃 추론, 벡터화 pass 파이프라인, VMI 명령 생성 및 Fragment 논리 매핑을 담당하고, PTOAS 측은 물리 레지스터 분할, 마스크 구체화, Fragment 버퍼와 벡터 레지스터 할당, mem_bar 삽입을 담당합니다. 로드맵은 다섯 개 마일스톤으로 나뉘며, 이번 커밋은 M1 핵심 프레임워크(원소별 연산, extent가 lanes로 나누어떨어지는 경우와 그렇지 않은 경우 모두 포함)이고, M2 제어 흐름, M3 reduce, M4 UB 용량 계획과 레지스터 할당, M5 비연속 접근과 더 넓은 PACK 데이터 타입이 차례로 이어집니다.

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

1.4 메인라인: Ascend DeepGEMM의 scale 버퍼를 쓰기 카운트 다중 버퍼로 전환, 데이터 로드를 조건부 로드 앞으로 이동 가능(#3629)

10-30 05:48에 병합되었습니다. DeepGEMM은 여러 K 블록 사이에서 scale factor를 재사용하는데, 이번 커밋은 명시적 쓰기 카운트 다중 버퍼 모드를 새로 추가하여 scale 버퍼의 버전이 실제 쓰기를 기준으로 전진하도록 합니다. 각 owner 루프 반복에서 실제로 실행되는 첫 번째 쓰기 작업에서만 한 번 전진하고, 이후 쓰기는 해당 버전을 공유하며, 쓰기가 없는 반복은 이전 값을 유지하고, 여러 개의 서로소 owner와 공유 클록을 지원합니다. 이로써 데이터 로드를 조건화된 scale 로드 앞에 배치할 수 있게 되었습니다. 동시에 DeepGEMM의 K 루프를 평탄화하여 L1과 L0 모두에서 데이터를 먼저 가져오고 scale을 나중에 가져오도록 했으며, 가드 업데이트, 중첩 업데이트, 보존 읽기, flag와 버퍼 재사용, 그리고 FP4의 K 주순서와 MN 주순서에 꼬리 블록과 링 랩어라운드가 있는 사용 사례를 보강했습니다. 이 모드는 명시적 버전 수를 요구하며, 각 owner 반복에서 최대 한 번의 업데이트를 가정합니다.

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

1.5 메인라인: 두 건의 수정 병합, CPU 스칼라 GEMM은 대상 영역만 초기화, 음수 인덱스 합법화는 도달 불가 분기 건너뜀(#3709 / #3712)

두 건 모두 10-30에 병합되었으며, 지난 호에서 심사 중이던 항목의 반영입니다. #3709는 CPU 스칼라 GEMM이 대상 영역만 정리하도록 하여, 이전에는 이번 출력에 속하지 않는 범위까지 함께 지워지던 문제를 해결했습니다. #3712는 LegalizeNegativeIndex가 도달 불가 분기를 건너뛰도록 하여, 실행되지 않을 코드에 대해 음수 인덱스 재작성을 수행하지 않도록 했습니다. 이 두 건은 지난 호 리포트의 #3708, #3710과 같은 배치의 결함 정리이며, 그중 #3708의 비트 폭 int64 변경과 #3710의 캐시 비활성화 시 두 계층의 프로세스 내 캐시를 우회하는 건은 여전히 심사 중입니다.

출처: https://github.com/tile-ai/tilelang/pull/3709, https://github.com/tile-ai/tilelang/pull/3712

1.6 메인라인: 언어 계층에서 proxy 내 사용되지 않는 Ref 자리표시자 두 개 제거(#3722)

10-31 17:31에 병합되었습니다. tilelang/language/proxy.py에는 두 개의 Ref가 공존하고 있었습니다. 하나는 타입 검사용 tirx.Var 스텁이고, 다른 하나는 빈 런타임 클래스이며, 둘 다 트리 내에 소비자가 없습니다. 외부로 노출되는 T.Ref는 실제로 eager.builder에서 오며, 그곳의 구현이 매크로 참조 인자의 본체입니다. 이번 커밋은 두 자리표시자, _T 타입 변수, 그리고 그에 따라 무용해진 Generic, TypeVar 임포트를 삭제하고, T.Ref와 매크로 참조 동작은 유지하며, 사용되지 않는 직접 임포트 경로인 tilelang.language.proxy.Ref만 제거합니다. 프런트엔드 구성 테스트 179개 통과, 6개 건너뜀, 런타임 아이덴티티 검사로 T.Ref가 여전히 eager.builder.Ref임을 확인했습니다.

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

1.7 신규 결함 및 심사 중: TMEM 결과 조용히 유실, ROCm 심볼 미발견, TF32 최근접 짝수 반올림(#3721 / #3711 / #3714)

#3721은 이번 모니터링 기간에서 가장 경계해야 할 항목입니다. sm_100a에서 TMEM에 대해 T.gemm을 수행할 때 threads/32가 N을 나누어떨어뜨리지 못하면, 컴파일은 통과하지만 PTX에는 tcgen05.mma가 하나도 없고, T.copy(Ct, Cl)이 아무도 쓰지 않은 TMEM을 읽어 곱이 조용히 유실됩니다. 재현 사례에서 threads=256은 한 개가 나오고, threads=384는 0개가 나오며, 작성자는 src/cuda/op/gemm.cc가 TCGEN05에 {1, num_warps} 파티션을 반환할 때 커버리지 검사가 빠진 것으로 의심합니다. #3711은 ROCm 7.14와 10.0의 PyTorch wheel에서 kernel이 “named symbol not found”를 보고합니다. #3714는 심사 중이며, fp32에서 tf32로의 변환에 최근접 짝수 반올림(cvt.rn.tf32.f32 및 포화 형식)을 추가합니다. 비포화는 SM90 이상, 포화는 SM100 이상을 요구하며, 해당 기능 요청 #3713을 종료합니다.

출처: https://github.com/tile-ai/tilelang/issues/3721, https://github.com/tile-ai/tilelang/issues/3711, https://github.com/tile-ai/tilelang/pull/3714

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

2.1 Ascend: Cython wrapper가 TVM Analyzer를 지연 생성하여 910B3 단일 호출이 약 450마이크로초에서 약 106마이크로초로 단축 (#1948)

10-30 07:44에 병합됨. CythonKernelWrapper.forward()는 이전에 매 호출마다 무조건 TVM Analyzer()를 생성했으며, 910B3에서 약 0.28밀리초의 고정 호스트 오버헤드가 발생했는데, 이 객체는 result 또는 workspace 형상이 기호 표현식일 때만 사용되었고, 정적 shape의 kernel에서는 전혀 필요 없음에도 동일한 비용을 지불했다. 변경 방식은 생성을 tir.PrimExpr 분기 내로 지연하고 필요에 따라 재사용하는 것이다. 실측: vec add 단일 블록 442~467마이크로초에서 105.1마이크로초로, 32블록 441~493마이크로초에서 108.0마이크로초로, kpool 압축 447~450마이크로초에서 111.3마이크로초로, 순수 ctypes 호출 약 12.5마이크로초에서 약 12마이크로초로 단축됨. 정적 경로에서 생성하지 않음과 기호 경로에서 필요 시 생성함을 각각 검증하는 두 개의 회귀 테스트가 추가됨.

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

2.2 Ascend: RFC에서 동기화 구현을 메인라인 코어 + 독립 Z3 증명 서비스로 교체 제안 (#1971)

이 RFC는 속도 향상 결론을 미리 정하지 않고 비교 기준선만 고정한다: Ascend 저장소는 TileLang 43b0b69와 TVM c2921fd 기반, IR 기반은 구 tvm.tir; 메인라인은 TileLang 1c0f402와 TVM 4211874e9e 기반, IR 기반은 이미 tvm.tirx와 tvm.s_tir이며 객체와 등록 체계도 새로운 FFI와 리플렉션 시설로 교체되었다. 따라서 마이그레이션은 InsertSync 파일 하나를 복사하는 것으로 단순화할 수 없다. 제안 경로는 메인라인 동기화 코어와 독립 Z3 증명 서비스를 이식하고, 구 TIR 어댑터 계층을 통해 기존 컴파일 파이프라인에 연결하며, 기존 TVM, 스케줄링, 메모리 계획, AscendC/PTO codegen 및 런타임 ABI를 유지하고, 지원 범위, 정확성, 성능 및 설치 검수를 완료한 후 기본 구현을 전환하고 구 동기화 알고리즘을 삭제하는 것이다. 관련 PR은 #1965이다.

출처: https://github.com/tile-ai/tilelang-ascend/issues/1971

2.3 Ascend 심사 중: TileLang-Ascend 전체를 메인라인 멀티 백엔드 아키텍처로 이전 (#1975)

10-31에 제안되어 지속 업데이트 중이며, 목표는 제안 #1911에 따라 TileLang-Ascend를 TileLang 메인라인의 멀티 백엔드 아키텍처로 마이그레이션하는 것이다. 설명에는 두 가지 diff 대조가 제시되며, 하나는 ascendc_pto 브랜치 대비, 하나는 메인라인 대비이다. 이는 2.2의 RFC와 동일한 수렴 경로의 두 측면으로, 하나는 동기화 알고리즘을 바꾸고 하나는 저장소의 골격을 바꾼다.

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

2.4 Ascend 심사 중: GLM-5.3 kpool의 fp8_mqa_logits 예제, 듀얼 kernel + AIV/Cube 분업 (#1966)

메인 저장소 kpool 디렉터리의 GLM-5.3 fp8_mqa_logits 예제를 Ascend 버전으로 구현한 것이다. 연산자 의미는 GLM-5.3-Flash 희소 어텐션 indexer의 풀링 히스토리 스코어링이다: 각 query 행에 대해 유효 pool 구간 내에서 페이지 테이블에 따라 64-pool 페이지에서 E4M3 key를 읽고(uint8 바이트 + fp32 scale), 역양자화 후 32개 indexer 헤드의 MQA 스코어를 계산하고, per-pool scale을 곱한 뒤 ReLU gate와 헤드별 가중 리덕션을 거쳐 각 pool당 하나의 FP32 logit으로 만들며, 범위 밖이나 유효하지 않은 페이지는 음의 무한대를 출력한다. 구현은 듀얼 kernel이다: mask kernel은 64-pool tile 단위로 레인별 0 또는 음의 무한대 마스크를 생성하여 호스트 사전 채움을 대체하고, main kernel은 각 행의 각 페이지당 하나의 block으로, 두 AIV가 각각 32-pool 반쪽 페이지를 처리하며, Cube는 전체 페이지를 읽어 64×32×128 GEMM을 한 번 수행한다. 헤드 수 32, 차원 128, 페이지 크기 64.

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

2.5 Ascend: 4건의 문서 및 API 사실 검증 병합, clamp가 PTO에서 실제로 적용됨 (#1899 / #1893 / #1620 / #1600)

네 건 모두 10-30에 병합되었으며, 지속적으로 진행되는 API 사실 검증 배치에 속합니다: #1899는 T.tile.bitwise_xor 문서를 검증(#1681의 재제기), #1893은 T.tile.bitwise_rshift를 검증(#1676의 재제기), #1620은 T.tile.clamp가 PTO에서 실제로 작동하도록 하고 문서를 동기화하며 테스트와 API 문서를 보강, #1600은 axpy의 docstring을 갱신하고 테스트를 강화하며 API 문서를 보충합니다. 같은 배치에 #1968, #1967, #1974가 심사 중이며, 각각 sigmoid와 round 카운트 반감, 기호 형상 파싱 캐시, 전체 블록 행 분할 시 tile 크기가 잘못 반감되는 문제의 수정입니다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1620, https://github.com/tile-ai/tilelang-ascend/pull/1899

2.6 Ascend: 일일 테스트 2970개 케이스 전부 통과, 11-01에 workflow 실패 1회 (#1977 / #1976)

11-02 06:42의 일일 테스트 기록은 2970개 항목, 2970개 통과, 0개 실패로 통과율 100%입니다. 10-31 06:56 회차도 마찬가지로 2970개 전부 통과했습니다. 중간의 11-01 05:08 회차는 workflow test job failure로 표시되었으며, 이 모니터링 기간 내 #1975의 마이그레이션 변경, #1971의 RFC와 같은 시기에 나타나 주의가 필요한 변동이며 안정적인 회귀는 아닙니다.

출처: https://github.com/tile-ai/tilelang-ascend/issues/1977, https://github.com/tile-ai/tilelang-ascend/issues/1976

2.7 Hygon: MLS 경계 및 지속 WASP 동기화 수정 병합, gfx948 호환 심사 중 (#31 / #32)

#31은 10-30 10:56에 병합되었으며 세 가지 사항입니다: 동적 MLS 경계 추론과 MatrixStore의 stride 처리 개선, MMAC와 blockscaled GEMM의 clear_accum 하강 수정, 지속 WASP 경로에서 M 파티션 컨슈머와 독립적인 A/B 펜스로 전환하면서 slot과 phase 순환을 유지하고 테일 컨슈머에 동기화 추가. #32는 심사 중이며, gfx948을 gfx946 호환으로 처리하여 기존 명령이 새 모델에서 직접 작동하도록 합니다. 이전 모니터링 기간에 보도된 일반 복사 LDS swizzle(#30)은 같은 시기에 착지했습니다.

출처: https://github.com/tile-ai/tilelang-hygon/pull/31, https://github.com/tile-ai/tilelang-hygon/pull/32

2.8 mlir-ascend: CI에서 내부 캐시 서비스 참조 제거, Codex 워크플로 심사 중 (#206 / #208)

#206은 10-31 07:24에 병합되었으며, CI에서 내부 캐시 서비스를 가리키는 내용을 제거한 것으로 재현성 정리 작업입니다. #208은 심사 중이며, Operator Harness에 첫 Codex 워크플로를 추가합니다: 반복 질의로 프로젝트 초기화 완료, 기존 연산자 저장소 연결, 영속 컨텍스트로부터 개발·튜닝·분석 재개; 세 가지 skill이 .workspace/ 아래의 편집 가능한 Markdown 기록을 공유하며, 작업 인계 시 「현재 후보」와 「검증된 최적 후보」를 구분하고 실패 및 롤백 실험을 보존합니다. 작성자는 이번에 실제 NPU 실측이 없음을 명확히 하였으며, NPU 실행 적응과 npm 배포는 후속 작업으로 명시했습니다.

출처: https://github.com/tile-ai/tilelang-mlir-ascend/pull/206, https://github.com/tile-ai/tilelang-mlir-ascend/pull/208

2.9 Moore Threads와 MetaX는 모니터링 기간 내 푸시 없음 (11-02)

tilelang-metax의 마지막 푸시는 10-26에 멈춰 있고, tilelang-musa는 09-30에 멈춰 있으며, 두 저장소 모두 이 모니터링 기간의 72시간 내에 커밋이 없습니다. 이는 이번 호의 백엔드 동향이 Ascend, Hygon, mlir-ascend 세 측에 집중되었음을 의미합니다.

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

3. 생태계 및 채택 측 (TileOPs / TileFoundry / TileKernels / FlashQLA)

3.1 TileOPs: 72시간 내 15건 병합, 13건이 11-01에 집중, 거의 모든 건에 실측 표와 체크리스트 회귀 게이트 포함 (11-01)

v0.1.0 릴리스 후 TileOPs는 11-01 하루에 13건을 머지했으며, 10-30의 #2794와 10-31의 #2795를 더해 72시간 동안 총 15건을 기록했습니다. 변경 사항은 거의 전부 perf 계열로, 어텐션, GEMM, 원소별 연산, 리덕션, MoE를 포괄하며, 각 건마다 H200에서의 busy 및 back-to-back 두 가지 시간 또는 vLLM, FlashInfer와의 대조를 제시하고, manifest에 이미 측정된 행의 회귀가 없음을 선언했습니다. 이러한 ‘표 첨부 머지’ 방식은 연산자 라이브러리가 기능 보강 단계가 아니라 shape별 행 단위 교정 단계에 진입했음을 보여줍니다.

출처: https://github.com/tile-ai/TileOPs/commits/main

3.2 TileOPs: GQA 덴스 프리필이 하나의 실측 규칙으로 decode 또는 prefill 프로그램을 선택, 선택 리그렛 1.818에서 1.017로 감소 (#2808)

11-01 13:04에 머지. GQADenseFwdOp와 GQADenseWithLSEFwdOp가 동일한 규칙 AttentionCall.dense_decode_program에 의해 어떤 프로그램을 탈지 결정하도록 변경되었으며, 입력은 query 및 key 길이, batch, 헤드 그룹, head dim, RoPE, SM 수입니다. 2~64개의 query 위치는 warp 전용 decode를 타고, 하나의 KV 헤드 그룹 내 각 위치가 64행 블록에 패킹됩니다. 단일 query 위치일 때는 32개 세그먼트로 분할해 한 웨이브를 채우거나, 1 시퀀스 1~2개 KV 헤드의 batch-1 프로그램을 탑니다. 최소 사용 가능 프로그램 대비 선택 리그렛(기하 평균 및 최댓값): 단일 위치 1.068과 1.87에서 1.018과 1.69로, 2~128 위치 1.818과 25.2에서 1.017과 1.81로 감소했습니다. 동시에 메인라인의 두 가지 오류를 수정했습니다: score scale에 log2(e)를 곱한 후 FP32에서 소실되어 NaN 또는 오값이 발생하는 문제; D256 다중 위치 호출이 SM90에서 공유 메모리 부족으로 기동 실패하는 문제.

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

3.3 TileOPs: 페이지드 MHA decode를 공유 메모리에 들어가는 모든 head dim으로 확장, GQA 페이지드 프리필에서 구 decode 커널 제거 (#2807 / #2797)

#2807은 11-01 12:30에 머지되어, 페이지드 MHA decode를 ‘하나의 블록 공유 메모리에 들어간다’는 전제 하에 모든 head dim으로 확장했습니다. 이전에는 큰 head dim을 만나면 이 경로를 타지 않았습니다. #2797은 같은 날 03:54에 머지되어, GQAPagedFwdOp에서 구 paged decode kernel을 삭제하고 페이지드 MHA decode의 상한을 컴파일 가능한 head dim으로 낮췄으며, 이는 이전 버전이 남긴 범위 초과 능력을 회수한 것입니다. 두 건을 합치면 ‘먼저 능력 확장, 다음 중복 제거’의 마무리입니다.

출처: https://github.com/tile-ai/TileOPs/pull/2807, https://github.com/tile-ai/TileOPs/pull/2797

3.4 TileOPs: coop2 프로듀서를 단일 warp로 축소, 192폭 전치 블록 추가, W4A8 부분 웨이브를 1 CTA 1 tile로 변경 (#2806 / #2805)

#2806은 11-01 12:02에 머지되어, Hopper GEMM의 coop2 프로듀서를 다중 warp에서 단일 로드 warp로 축소하고, 공유 메모리가 3단 파이프라인을 허용할 때 192폭 균등 전치 블록을 제공했습니다. H200에서 wgrad-q-b가 384.9마이크로초에서 382.3으로, mlp-up fp16이 317.9마이크로초에서 311.9로 감소했으며(back-to-back 각각 458.5에서 449.6, 369.0에서 367.3), 작성자는 측정된 GEMM 행의 회귀가 없음을 선언했고, 이 건은 #2791 목표의 일부이며 듀얼 카드 확인은 아직 완료되지 않았습니다. #2805는 같은 날 10:36에 머지되어, SM의 최소 4분의 3을 채우는 부분 웨이브 W4A8 prefill 그리드를 stream-K 대신 1 CTA 1 tile로 변경했습니다. llama-8b-up-b33이 23.52마이크로초에서 17.98로(vLLM 20.42), b128이 31.33에서 23.39로(vLLM 24.05) 감소했으며, 목록 외 M 33~128의 44개 N/K 스캔에서 스케줄 형태 변경이 5%~36% 빨랐고 중앙값 18%였으며 느려진 경우는 없었습니다.

출처: https://github.com/tile-ai/TileOPs/pull/2806, https://github.com/tile-ai/TileOPs/pull/2805

3.5 TileOPs: 컴파일 캐시를 빌드 타깃 기준으로 키잉, 연산자 계층에서 출력과 입력 중첩을 통합 검사 (#2802 / #2801)

#2802는 11-01 11:31에 병합되었으며, 컴파일러 캐시의 크로스 디바이스 재사용 문제를 수정했다. 빌더와 Op 항목 모두 빌드 대상별로 키를 구성하고, 지연 컴파일과 자동 튜닝은 예상 디바이스에 유지되며, 검증은 99개 항목 대상 격리, 벤더 및 그룹 GEMM 테스트로 이루어졌다. #2801은 같은 날 10:33에 병합되었으며, “출력은 다른 어떤 호출 텐서와도 바이트를 공유해서는 안 된다”는 검사를 연산자 계층으로 끌어올리고, manifest의 overlaps 필드로 연산자별 예외를 선언하도록 했으며, MoE unpermute, MoE 그룹 GEMM, 지속 GEMM 템플릿 등 네 곳의 kernel 측 검사를 삭제했다. 아울러 메인라인의 경쟁 조건 하나를 수정했다. IndexedExpertMLPFwdOp가 out과 topk_weights의 중복을 허용했지만, 후자는 쓰기 시점에 여전히 리덕션 읽기로 사용되고 있었다. MoEPostPermuteFwdOp의 out을 포함한 호스트 소요 시간은 51.48마이크로초에서 39.33으로 감소했다.

출처: https://github.com/tile-ai/TileOPs/pull/2802, https://github.com/tile-ai/TileOPs/pull/2801

3.6 TileOPs 수정 중인 결함: Norm 4개 연산자가 0에서 먼 행에서 정밀도 손실, 약 1e36에서 오버플로로 NaN 발생 (#2812)

11-01 신규 오픈. LayerNorm, AdaLayerNorm(-Zero 포함), GroupNorm, BatchNorm은 현재 모두 원시 행의 fp32 직접 합산으로 평균과 분산을 구한다(BatchNorm은 E[x²]에서 평균 제곱을 뺀다). 이는 #2773에서 FusedAddLayerNorm을 위해 수정했고, #2811에서 InstanceNorm 훈련 커널을 위해 수정 중인 것과 동일한 유형의 결함이다. 증상은 다음과 같다. 상수 10000.1의 fp32 행은 정규화 후 출력이 bias가 아닌 대략 ±weight가 되고, 10000에 난수를 더한 행은 fp32 허용 오차에서 float64 참조와 일치하지 않으며, 1e36에 가까운 행은 합산 오버플로로 fp32와 bf16 모두 출력이 전부 NaN이 된다. 영향을 받는 경로는 fragment, 공유 메모리 순회, 행당 하나의 warp, 블로킹 및 스트리밍 등 여러 구현을 포함한다.

출처: https://github.com/tile-ai/TileOPs/issues/2812

3.7 TileFoundry: lowering이 issuing mesh와 윈도 분할을 한 곳의 결정으로 통합, 83개 스케줄링 픽스처 바이트 단위 일치 (#280 / #282)

#280은 10-31 00:47에 병합되어 “호출마다 무엇을 발행할지”의 결정을 한 곳으로 모았다. #282는 11-01 11:47에 병합되어 “어느 mesh가 발행할지”와 “윈도를 어떻게 분할할지” 역시 한 곳으로 모았다. analysis/emission.py에 copy_mesh, issue_owner, issue_mesh를 새로 추가하고, copy_mesh와 issue_owner는 하나의 헬퍼를 공유하며 두者的 유일한 차이(Partial을 읽는 방식)만 한 곳에 남겼다. 세 곳에 중복된 Slice 분할은 _cut 헬퍼로 통일했다. 검증 방식은 바이트 수준 대조다. 83개 스케줄링 픽스처 중 75개의 finalize 출력이 메인라인과 완전히 일치하고, 나머지 8개는 메인라인에서도 동일하게 실패한다. docs/spec/passes.md 7.1절에 이 mesh 결정 체계를 규범으로 기록했다.

출처: https://github.com/tile-ai/TileFoundry/pull/282, https://github.com/tile-ai/TileFoundry/pull/280

3.8 TileFoundry: 사이트에 0.3 버전 스케줄링 글 “Scheduling, Designed for Agents” 게시 (#4 / #281)

사이트 저장소는 10-30에 #4를 병합하고 0.3 버전 글 “Scheduling, Designed for Agents”를 게시했으며, 동시에 글과 이미지의 디렉터리 구조, 수식의 MathJax 지원, 캡션과 표 스타일, 각 차트 앵커를 보완했다. 홈페이지 Blog 카드는 블로그 인덱스를 열도록 변경되었고, 인덱스는 두 편의 글을 나열한다. upstream-ref는 변경되지 않았으며 사이트는 여전히 v0.0.3을 렌더링한다. 저장소 측 #281은 10-31에 README에서 이 글을 공지했다. TileFoundry 라인에서 이는 “에이전트를 위한 스케줄링”을 코드와 릴리스 노트에서 공개 서사로 옮기는 한 걸음이다.

출처: https://github.com/tile-ai/TileFoundry.github.io/pull/4, https://github.com/tile-ai/TileFoundry/pull/281

3.9 채택자: FlashQLA, SM120의 Dual-CP 전방/역방향 병합, RTX 5090 2장에서 32K~128K 시퀀스로 Inter-CP와 비교 (#51)

10-30 03:02 병합. 변경 사항은 SM120에 Dual-CP의 순방향과 역방향을 추가한 것이다. 순방향은 기존 state-scan kernel을 재사용해 카드 수준 상태 집계를 수행하고, 정렬된 M 곱은 32열 단위로 계산해 공유 메모리 내로 제어한다. 역방향은 전치된 M의 역순으로 상태를 전파하며, 순방향에서 캐시된 M 행렬을 재사용한다. 동시에 역방향 kernel의 공유 메모리 경합 한 건을 수정했으며, 방법은 7단계와 dV 저장 warp를 동기화하는 것이다. 성능은 2장의 RTX 5090에서 BF16, chunk 32, 시퀀스 길이 32K~128K로 Inter-CP와 비교했으며, 각 케이스당 3회 중앙값, 회당 25회 워밍업과 100회 타이밍으로 측정했다.

출처: https://github.com/QwenLM/FlashQLA/pull/51

3.10 채택 측: TileKernels 10-14 이후 푸시 없음 (11-02)

DeepSeek의 TileKernels 마지막 푸시는 10-14에 멈춰 있으며, 이번 모니터링 기간 동안 변화가 없다. 채택 측의 이번 기간 유효 동향은 전부 Tongyi FlashQLA에서 나왔다.

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

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

4.1 RISC-V 연구 체인 형성: 세 개 저장소가 Numba CUDA Python을 TileLang으로 변환한 뒤 제한된 PrimFunc을 단일 hart RVV C로 내림 (10-31)

커뮤니티 연구자 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 공식 등록 백엔드도, 범용 GPU 모델 시뮬레이터도 아니며, threadIdx에 의존하는 주소 지정, hart 간 동기화, 지원되지 않는 IR은 일괄 거부하고, block은 순차 실행되며, shared와 fragment는 단일 hart 로컬 배열로 매핑된다. 연구 환경은 TileLang 0.1.15, Clang 23.1.3, QEMU 11.1.0이며 다른 버전 호환성은 보장하지 않고, 업스트림 스레드 수준 산출물을 해당 백엔드가 수용 가능한 입력으로 자동 간주할 수 없음을 명시했다.

출처: https://github.com/zishu0131-png/tilelang-riscv

4.2 PTO-ISA/tilelangGym, PTO-ISA/tilelang에서 독립 저장소로 분리, 주 회귀 게이트는 Ascend SimtVF 경로 (10-30)

이 민감도 및 회귀 테스트 세트는 PTO-ISA/tilelang의 examples/ascend/t_parallel_st_suite에서 분리되었으며, 그 이유는 테스트 세트가 TileLang과 PTOAS의 버전을 독립적으로 고정할 수 있게 하기 위해서다. 현재 세 가지 주 경로는 Ascend SoftLib SimtVF(target=ascend, cython, 주 회귀 게이트, remat 버전에서 전부 통과), VectorizeParallelToPTO와 VL의 vload/vstore, 그리고 benches 아래 @pto.jit로 작성된 수동 SIMD DSL 대조 구현이다. 각 테스트 케이스 디렉터리에는 config.toml과 VERDICT.md가 있으며, 별도의 주간 및 CI 회귀 스크립트가 있다.

출처: https://github.com/PTO-ISA/tilelangGym

4.3 뉴스 측: Google News 중영문 쿼리, Hacker News, arXiv 모두 기간 내 실질적 적중 없음 (11-02)

사흘간의 Google News 중영문 쿼리는 기간 내에 해당하는 TileLang 관련 보도를 반환하지 않았고, Hacker News와 arXiv 역시 새 항목이 없었다. 후보 목록의 20건은 전부 GitHub 저장소 푸시에서 나온 것으로, 일반적인 코드 활동이지 뉴스가 아니다. 이번 기간의 정보 면은 따라서 전적으로 저장소 측과 커뮤니티 저장소로 구성되었다.

출처: https://news.google.com/

5. 동향 관찰

이번 기간에서 가장 기록할 만한 판단은 업스트림 스스로 나온 것이다. 정확성 개선을 되돌리더라도 ThreadSync의 복잡도를 핵심 경로에서 빼겠다는 것이다. 롤백 설명에는 그 대가가 명확히 적혀 있다. 몫 가드, 정확한 참여 스레드 경계, 비직사각형 스레드 집합이라는 세 가지 개선을 함께 철회하고 기존의 보수적 경로만 유지하며, #3732는 TIRX 마이그레이션 완료 후에나 최적화를 논하기로 남겨두었다. 한 컴파일러 프로젝트가 “우리 증명기 예산이 부족하니 우선 보수적 구현으로 되돌린다”고 공개적으로 쓰는 것은, 기능을 아무리 많이 쌓는 것보다도 현재 단계의 실제 제약을 더 잘 보여준다. TIRX 마이그레이션이 메인라인의 최우선 순위이며, 다른 최적화는 모두 그것에 길을 비켜줘야 한다.

이와 호응하는 것으로 Ascend 쪽의 두 가지 「메인라인으로 수렴」 텍스트가 있다: #1971의 RFC는 동기화 알고리즘을 메인라인 코어에 독립 Z3 증명 서비스를 더하는 방식으로 교체하려 하며, 먼저 조사하고 나중에 교체하며 중간에 구 TIR 적응 계층을 남긴다; #1975는 전체 Ascend 저장소를 메인라인 다중 백엔드 아키텍처로 직접 이전한다. 두 경로 모두 마이그레이션 비용을 치러야 하므로, 최근의 가시적인 작업은 기반类 작업 쪽으로 치우칠 것이다: #1948은 매 호출의 고정 호스트 오버헤드를 약 450마이크로초에서 약 106마이크로초로 줄였고, 네 건의 문서 사실 검증, CI 내부 캐시 참조 정리는 모두 이런 유형의 작업이다.

TileOPs의 리듬은 「릴리스」에서 「튜닝」으로 전환되었다. 72시간 동안 15건의 병합 중 13건이 같은 날에 집중되었고, 거의 모든 건에 H200 실측 표와 「manifest 무회귀」 선언이 첨부되었으며, 선택 후회와 같은 지표로 스케줄링 규칙의 좋고 나쁨을 설명하기 시작했다(#2808은 1.818에서 1.017로). 이는 연산자 라이브러리의 경쟁 지점이 이미 「이 연산자가 있느냐」에서 「각 shape마다 프로그램을 올바르게 선택하느냐」로 바뀌었음을 의미하며, 이것이 바로 TileLang 같은 DSL이 수작업 kernel 대비 갖는 장기적 가치다.

생태계 쪽에는 두 가지 새로운 분기가 나타났다. 하나는 하드웨어 방언이다: Hygon은 gfx948을 gfx946 호환으로 취급하고, Ascend는 PTO 문서와 심볼 형태를 계속 보완한다. 다른 하나는 언어 출구다: RISC-V 연구 체인은 제한된 PrimFunc를 단일 hart RVV C로 내리며, 스타 수는 매우 낮고 저자도 공식 백엔드가 아니라고 명확히 선언했지만, TileLang의 IR을 이식 가능한 자산으로 활용한다는 점에서 Ascend, Hygon의 저장소 수준 브랜치 적응과 같은 종류의 신호다. 여러 분기가 동시에 존재하는 상황에서 업스트림 TIRX 마이그레이션이 안정적인 IR 경계를 제시할 수 있는지가 이들 다운스트림 투자가 축적될지 재작업이 될지를 결정할 것이다.

부록, 자료와 검증 설명

출처 검증표

출처 검증 결과
tile-ai 조직(29 저장소) 모니터링 기간 내 10개 저장소 푸시
메인 저장소 tilelang 메인 브랜치 7건 병합; 신규 issue 4건; 문서 사이트 3건
TileOPs 메인 브랜치 15건 병합(13건은 11-01); 신규 issue 3건; 사이트 1건
TileFoundry 3건 병합; 사이트 1건
tilelang-ascend 5건 병합; 신규 issue 4건(일일 테스트 기록과 RFC 포함)
tilelang-hygon 2건 병합; 1건 심사 중 신규
tilelang-mlir-ascend 1건 병합; 1건 심사 중 신규
tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack 모니터링 기간 내 푸시 없음
릴리스 페이지(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / mlir-ascend) 모니터링 기간 내 새 버전 없음
FlashQLA 1건 병합(#51)
TileKernels 10-14 이후 푸시 없음
RISC-V 연구 체인 세 저장소 / PTO-ISA/tilelangGym 푸시 있음, 연구 체인 세 저장소는 같은 날 체인 생성
Google News / Hacker News / arXiv 모니터링 기간 내 실질적 적중 없음

전체 출처 목록

WASP 동기화 (#31) — https://github.com/tile-ai/tilelang-hygon/pull/31

** TileFoundry 릴리스 페이지 — https://github.com/tile-ai/TileFoundry/releases