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

출처: GitHub(tile-ai 조직 29개 저장소 푸시 전수 점검, 모니터링 기간 내 9개 저장소에 푸시 있음: TileOPs, TileOPs-nightly, TileOPs.github.io, tilelang, tilelang.github.io, TileFoundry, tilelang-ascend, tilelang-hygon, tilelang-mlir-ascend; 조직 모니터링 기간 내 병합 36건〔TileOPs 19, tilelang 11, TileFoundry 3, TileOPs.github.io 1, tilelang-ascend 1, tilelang-mlir-ascend 1〕, 심사 중 52건〔tilelang 24, tilelang-ascend 15, TileOPs 11, tilelang-hygon 2〕; 비기본 브랜치 개별 점검〔ascendc_pto, snapshots, stats〕) + 릴리스 페이지 및 버전 파일 점검 + 스펙 통계 브랜치 점검 + 채택 측 저장소 점검(TileKernels / FlashQLA) + 커뮤니티 저장소 점검 + 야간 파이프라인 점검 + Google News 중영문 다중 쿼리(프록시 경유) + Hacker News + arXiv.

이번 호 인덱스

  • 오늘의 핵심: SM89에서 하루에 두 건의 GEMM이 정식 반영되고, TileOPs 사양 및 벤치마크 커버리지가 205/205로 동시 보고되며, 메인 저장소가 Windows ARM64를 CI에 포함함 (10-19)
    1. 핵심 프로젝트 진행 상황
      • 1.1 메인 저장소: Windows ARM64 CI 진입, CUDA 13.4 arm64 패키지 설치 및 clang 호환 레이어 (#2962) (10-19)
      • 1.2 메인 저장소: 공유 버퍼上的 T.cumsum, warp 특수화下 스레드 오프셋 보완 (#3536) (10-19)
      • 1.3 메인 저장소: call_extern의 베어 data 변수를 보수적으로 쓰기로 표시 (#2618) (10-19)
      • 1.4 메인 저장소: lower trace를 호스트와 디바이스 소스로 분리, 편집 가능 베이스라인으로 변경 (#3496) (10-19)
      • 1.5 메인 저장소: Autotuner profiler 파라미터 전달, ROCm FP4 상수 출력 및 두 건의 CI 유지보수 (#3337 / #3548 / #3565 / #3567) (10-19)
    1. 멀티 백엔드 적응 (Ascend / MetaX / Hygon / Moore Threads)
      • 2.1 Ascend: 호스트 연산자와 디바이스 커널이 msprof下 연관 관계 구축 (#3321) (10-19)
      • 2.2 Ascend: VF 도메인 외 스칼라 sqrt를 표준 수학 라이브러리로 처리, PTO Ramp를 타입화된 벡터로 발사 (#3466 / #3535) (10-19)
      • 2.3 Ascend: 네 개의 원소별 연산자 중국어 API 문서 및 사용 예시 보완 (#1657) (10-19)
      • 2.4 Ascend 심사 대기열 15건: GLM-5.3 압축 연산자 및 세 건의 역방향 연산자 (10-19)
      • 2.5 Hygon: 심사 중 두 건으로 MLS 윈도우 상수 K와 레지스터 파이프라인 보완 (#25 / #18) (10-19)
      • 2.6 MetaX / Moore Threads / MLIR-Ascend: 양측 신규 푸시 없음, MLIR 라인 TVM 격리 수정 (#204) (10-19)
    1. 생태계 및 채택자 (TileOPs / TileFoundry / 연산자 라이브러리)
      • 3.1 TileOPs: SM89 블록 스케일링 FP8 그룹 GEMM과 INT8 W8A8 동일자 병합 (#2659 / #2655) (10-19)
      • 3.2 TileOPs: 다섯 건의 foundry 연산자로 MLA 페이징, KV 쓰기 및 융합 정규화 보완 (#2642 / #2635 / #2639 / #2634 / #2651) (10-19)
      • 3.3 TileOPs: GQA 페이징 FP8 포워드 implemented로 전환, 36개 사례 FA3 대비 벤치마크 (#2648) (10-19)
      • 3.4 TileOPs: 선형 어텐션 및 GEMM의 지점별 성능 마무리 시리즈 (#2650 / #2645 / #2637 / #2636 / #2641 / #2640 / #2646 / #2652) (10-19)
      • 3.5 TileOPs: Mamba-1 선택적 스캔 구현, MoE 라우팅 열거에 수렴 상한 추가 (#2638 / #2653) (10-19)
      • 3.6 TileOPs: 사양 및 벤치마크 커버리지 205/205 동시 보고, 야간 리포트에 연산자 칼럼 추가 (#2656) (10-19)
      • 3.7 TileFoundry: Apple AMX 지원 및 Qwen 프리필 예제 M2 Pro에서 구동 성공 (#244 / #240 / #245) (10-19)
      • 3.8 채택자 및 야간 파이프라인: TileKernels / FlashQLA 푸시 없음, 스냅샷 10-20으로 진행 (10-20)
    1. 커뮤니티, 튜토리얼 및 활동
      • 4.1 문서: 여섯 개 신규 연산자 API 페이지에 작성, 메인 저장소 문서 사이트 자동 업데이트 (#77) (10-19)
      • 4.2 커뮤니티 저장소: W4A16 kernel 실습 저장소와 CUDA 디버깅 스킬 저장소 부상 (10-19)
      • 4.3 버전 라인: 메인 저장소 0.1.15 여전히 미교체, TileOPs 0.1.12 계속 고정 (10-19)
      • 4.4 미디어 및 학술: 뉴스 검색 3건 적중 모두 직접 관련 없음, arXiv 신규 논문 없음 (10-19)
    1. 트렌드 관찰
      • 5.1 SM89가 신규 연산자의 기본 착지점이 됨
      • 5.2 선형 어텐션이 「청크당 한 번의 행렬 곱 + SM 웨이브별 전략 선택」으로 고정
      • 5.3 TileFoundry CUDA를 벗어남: Apple AMX 및 Qwen 프리필
      • 5.4 관측 가능성이 컴파일러 하드닝의 주축이 됨
      • 5.5 Ascend 연산자 라이브러리가 대화 모델 전용 연산자로 전환
  • 부록: 자료 및 검증 설명
  • 전체 출처 목록

오늘의 핵심: SM89에서 하루에 두 건의 GEMM이 정식 반영되고, TileOPs 사양 및 벤치마크 커버리지가 205/205로 동시 보고되며, 메인 저장소가 Windows ARM64를 CI에 포함함 (10-19)

이 모니터링 기간의 호스트 측은 여전히 중국 시간 야간에 활동이 집중되었다. 조직 내 병합 36건 중 TileOPs가 19건을 차지했으며, 그중 SM89(Ada급)의 블록 스케일링 FP8 그룹 GEMM과 INT8 W8A8 GEMM이 같은 날 구현 완료되었고, 5개의 신규 연산자와 함께 MLA 페이징, KV 캐시 쓰기, QK-Norm-RoPE 융합, 토큰별 양자화 RMSNorm을 한 번에 보완했다. 조직 규격 통계 브랜치는 곧바로 「구현 완료 205/205, 벤치마크 완료 205/205」로 갱신되었으며, 이는 연산자 라이브러리가 자체 기술한 규격 커버리지와 벤치마크 커버리지가 동시에 최대치에 도달했음을 의미한다. 메인 저장소 tilelang은 11건을 병합했으며, 가장 비중 있는 것은 Windows ARM64를 릴리스 파이프라인에 포함시킨 것(12개 파일, CUDA 13.4 arm64 설치 스크립트와 clang 호환 헤더)이고, 그 외에 정확성 수정 2건(공유 버퍼상의 T.cumsum, call_extern 쓰기 표기)과 디버깅 인프라 개편 1건(lower trace의 호스트와 디바이스 소스 분리)이 있다. 백엔드 측에서는 Ascend가 호스트 연산자와 디바이스 커널을 msprof 하에서 대응시켰고, Hygon의 2건은 레지스터 파이프라인과 상수 K 주소 지정을 보완하는 중이며, TileFoundry는 처음으로 Qwen 프리필 예제를 Apple M2 Pro의 AMX 유닛에서 128 token까지 실행하는 데 성공했다.

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

1.1 메인 저장소: Windows ARM64 CI 진입, CUDA 13.4 arm64 패키지 설치와 clang 호환 레이어(#2962)

이 커밋은 Windows on ARM을 정식 빌드 타깃으로 릴리스 파이프라인에 연결한다. PowerShell 설치 스크립트를 신규 추가하여 NVIDIA 프리릴리스 채널에서 CUDA 13.4의 windows_arm64 로컬 설치 패키지를 받아 고정 경로에 배치하고, CMake의 FindPipCUDAToolkit 모듈과 windows_arm64_clang_compat.h 호환 헤더를 신규 추가하여 ARM64에서 MSVC 툴체인의 헤더 파일 및 내장 기능 차이를 해결하며, pyproject와 두 개의 requirements를 동기적으로 조정하고 레이아웃 모듈과 MSVC 보조 모듈을 변경하여 새 플랫폼에 적응시킨다. 총 12개 파일, 356줄 추가, 45줄 삭제. 그동안 조직 측은 ROCm과 Ascend 양쪽을 연속으로 보완해 왔으며, 이 커밋은 Windows ARM64를 세 번째 비 x86 전달 면으로 추가한 것으로, Snapdragon X 시리즈 노트북으로 로컬 커널 개발을 하는 사용자에게 직접적인 사용성 개선이다.

출처: tile-ai/tilelang#2962 (10-19)

1.2 메인 저장소: 공유 버퍼상의 T.cumsum이 warp 특수화 하에서 스레드 오프셋 보완(#3536)

warp 특수화가 생산자와 소비자를 서로 다른 warp에 배분한 후, 스캔 계열 프리미티브가 공유 버퍼상에서 계산한 스레드 번호가 여전히 전체 CUDA 블록 기준으로 계산되어 cumsum의 행 귀속이 어긋났다. 이 변경은 CUDA와 HIP 두 스캔 템플릿에 thread_offset 파라미터를 도입하여 하위 파라미터의 스레드 하한(thread_bounds의 최솟값)에서 전달받고, 템플릿이 이를 기반으로 lane과 행 인덱스를 다시 정렬한다. 백엔드 범용 구현과 Python 측 테스트 케이스도 동기적으로 갱신되었으며, 총 4개 파일, 101줄 추가, 38줄 삭제. 이러한 문제는 실제로 warp 특수화를 활성화하고 버퍼가 공유 메모리에 위치할 때만 노출되며, 규모화된 사용 이후 역류된 경계 조건에 속한다.

출처: tile-ai/tilelang#3536 (10-19)

1.3 메인 저장소: call_extern의 로우 data 변수를 보수적으로 쓰기로 표기(#2618)

읽기 전용 파라미터 표기는 원래 별칭 분석으로 버퍼가 쓰였는지 판단했으나, 외부 호출(call_extern)이 로우 포인터를 통해 쓸 때 분석이 쓰기 동작을 보지 못해 읽기 전용으로 오판되어 상수 폴딩 또는 조기 재사용이 발생했다. 이 변경은 해당 pass에서 call_extern을 특별히 식별하여 그 로우 데이터 변수를 일괄적으로 보수적으로 쓰기 완료로 표기하고, 「call_extern이 출력 버퍼에 쓰기」하는 표현을 커버하는 69줄의 테스트 케이스를 추가했다. 방향상으로는 최적화를 덜 하더라도 잘못된 최적화는 하지 않는다는 것으로, 최근 메인 저장소의 정확성 수정 배치와 궤를 같이한다.

출처: tile-ai/tilelang#2618 (10-19)

1.4 메인 저장소: lower trace가 호스트와 디바이스 소스를 분리, 편집 가능한 베이스라인으로 변경(#3496)

디버깅 도구 lower trace는 원래 생성 소스 한 부만 출력하여 호스트 측 래핑과 디바이스 측 커널이 섞여 있어 한 곳을 수정해도 되돌리기 어려웠다. 이 변경은 양측을 독립 산출물로 분리한다. 디바이스 측은 codegen_device.cpp 및 그 .original 베이스라인과 .latest 최신 버전을 출력하고, 호스트 측은 호스트 코드 생성 실행 시 codegen_host.cpp와 동일 구조의 베이스라인 파일을 별도로 출력하며, report.html은 최근 실행을 가리키는 소프트 링크로 연결된다. 문서도 편집 및 재컴파일 흐름을 동기적으로 개편했으며, 총 6개 파일, 121줄 추가, 63줄 삭제.

출처: tile-ai/tilelang#3496 (10-19)

1.5 메인 저장소: Autotuner profiler 파라미터 전달, ROCm FP4 상수 출력 및 CI 유지보수 2건 (#3337 / #3548 / #3565 / #3567)

#3337은 단 3줄이지만 자동 튜닝의 측정 기준을 수정한 것이다. 데코레이터의 warmup, rep, timeout이 이전에는 튜너에 전달되지 않아 사용자가 데코레이터에서 이 세 파라미터를 바꿔도 실제로 적용되지 않았는데, 이제 파라미터를 그대로 전달한다. #3548은 ROCm 측 FP4(e2m1) 부동소수점 상수 출력을 보완한 것으로, 해당 타입은 원시 바이트로만 생성할 수 있기 때문에 디바이스 측에서 헬퍼 함수를 통해 float32에서 변환하여 상수를 생성하도록 변경해 HIP 코드 생성 오류를 방지한다. 나머지 2건은 유지보수성 변경으로, cibuildwheel 버전을 4.2에서 4.3으로 올리고 ROCm 라인에 CodeRabbit 리뷰어 1명을 추가한 것이다.

출처: tile-ai/tilelang#3337, #3548, #3565, #3567 (10-19)

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

2.1 Ascend: 호스트 연산자와 디바이스 커널을 msprof에서 연관시키기 (#3321)

Ascend 측 profiler는 이전에 디바이스 커널 소요 시간만 보고할 수 있었고 호스트 연산자와 커널 간에 대응 관계가 없어 자동 튜닝이 얻은 샘플을 특정 호출로 귀속시킬 수 없었다. 이번 변경은 네이티브 profiling 채널을 추가한다. 빌드 시 CANN의 profiling 헤더를 탐지하며(msprof와 profiling 두 설치 레이아웃 호환), 설치되어 있으면 활성화하고 없으면 건너뛰며 사용자 대상 컴파일 스위치를 도입하지 않는다. 런타임 측에는 ascend_profiling.h와 호스트 측 profile 모듈을 보완하여 호스트 연산자의 발행과 디바이스 커널 실행을 호출 순서대로 대응시키고, 그룹 컴파일 경로가 동일한 보고 체계를 재사용하도록 했다. 총 13개 파일, 676줄 추가, 13줄 삭제로 Ascend 라인의 측정 능력에 대한 구조적 보완이다.

출처: tile-ai/tilelang#3321 (10-19)

2.2 Ascend: VF 도메인 밖 스칼라 sqrt는 표준 수학 라이브러리 사용, PTO Ramp는 타입화된 벡터로 발행 (#3466 / #3535)

#3466은 스칼라 sqrt가 벡터화 도메인(VF) 밖에 위치하는 상황을 처리한다. 하위 계층은 원래 이를 일괄적으로 sqrtf로 내렸지만 스칼라 코어에서는 해당 구현을 사용할 수 없어, 도메인 밖이고 타입이 float32일 때 AscendC::Std::sqrt를 사용하도록 변경하고 필요에 따라 표준 수학 헤더를 도입했으며 테스트 케이스 포함 126줄을 추가했다. #3535는 PTO 코드 생성에서 Ramp 표현식의 발행 방식을 수정하여 각 lane 표현식으로 타입화된 벡터를 조립하도록 변경했고, 테스트 케이스는 병렬 인덱스와 명시적 스트라이드 두 가지 Ramp의 벡터화 저장 생성 및 실행을 커버한다.

출처: tile-ai/tilelang#3466, #3535 (10-19)

2.3 Ascend: 4개 원소별 연산자의 중국어 API 문서 및 테스트 케이스 보완 (#1657)

Ascend 적응 저장소는 T.tile 아래의 relu, leaky_relu, sigmoid, silu 4개 원소별 연산자에 대해 중국어 API 문서를 보완하여 기능, 파라미터, 반환값 및 전형적인 사용법(예: leaky_relu의 음수 기울기 의미)을 명시하고, 테스트 케이스와 문서 진입점을 동시에 보완했으며 총 6개 파일, 451줄 추가, 19줄 삭제이다. 이는 Ascend 라인에서 드물게 문서를 주체로 한 병합으로, 해당 저장소의 연산자 면이 이미 외부에 의미를 약속해야 하는 단계에 진입했음을 보여준다.

출처: tile-ai/tilelang-ascend#1657 (10-19)

2.4 Ascend 심사 대기 큐 15건: GLM-5.3 압축 연산자와 역방향 연산자 3건 (10-19)

Ascend 저장소는 이번 모니터링 기간에 심사가 가장 집중된 곳이다(15건). 연산자 측에는 뚜렷한 대화 모델 지향성이 나타난다. GLM-5.3의 kpool 압축 연산자(#1910), KDA의 chunk 역방향(#1895), indexer 역방향(#1887), situ_glu 및 그 그래디언트(#1891 / #1877)가 있다. 예제 측에는 cann-bench의 lstm과 gru 2건(#1892 / #1902)을 보완했다. 컴파일러 측 2건은 안정성에 속한다. #1907은 UB에서 L1으로의 임시 저장에서 블록 주석과 ND 오프셋을 보존하고, #1903은 동시 커널 캐시의 발행과 캐시 삭제 후 복구를 수정한다. 또한 5건은 비트 연산 연산자의 API 사실 확인(#1651 / #1655 / #1656 / #1659 / #1893)으로, 그중 #1893은 지난주 롤백 후 재제출이고 나머지 4건은 신규 확인이다.

출처: tile-ai/tilelang-ascend#1910, #1895, #1887, #1891, #1907, #1903 (10-19)

2.5 Hygon: 검토 중인 두 건이 MLS 윈도 상수 K와 레지스터 파이프라인을 보완 (#25 / #18)

Hygon 저장소의 기본 브랜치 모니터링 기간 내에는 머지된 커밋이 없으나, 검토 중인 두 건 모두 실질적인 진전이 있었다. #25는 루프 언롤링 이후 MLS 리소스 승격에서 상수 윈도 오프셋이 유실되는 문제를 수정한다. 먼저 갱신되지 않은 축을 검증한 뒤 단일 축 갱신 또는 전방 차분 주소 지정을 채택할지 결정하고, 필요 시 보수적으로 윈도를 재구성하며, 기존 MLS 테스트에 회귀 케이스를 추가한다. #18은 HCU 백엔드에 레지스터 파이프라인을 추가하여 기존의 공유 메모리 파이프라인과 중첩될 수 있게 하고, 파이프라인 단계에 따라 해당 커밋 및 배리어 명령을 삽입하며, warp 수가 4를 초과할 때 warp_diverge를 활성화한다. PR 설명에 따르면 example_gemm의 성능이 이미 rocblas에 근접했다고 한다.

출처: tile-ai/tilelang-hygon#25, #18 (10-19)

2.6 MetaX / Moore Threads / MLIR-Ascend: 양측 신규 푸시 없음, MLIR 라인은 TVM 격리 수정 (#204)

tilelang-metax의 마지막 푸시는 10-16에, tilelang-musa는 09-30에 멈춰 있으며, 두 저장소 모두 모니터링 기간 내 커밋이 없어 정상적인 정적 상태이다. tilelang-mlir-ascend는 TVM 격리 수정(#204)을 머지했는데, 해당 저장소와 TVM 네임스페이스 및 심볼 가시성 간의 상호 오염을 처리한다. 규모는 크지 않으나 멀티 백엔드 공존 시 링크 안정성에 직접 영향을 미친다.

출처: tile-ai/tilelang-mlir-ascend#204 (10-19)

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

3.1 TileOPs: SM89 블록 스케일링 FP8 그룹 GEMM과 INT8 W8A8 동일자 머지 (#2659 / #2655)

#2659는 블록 스케일링 FP8 그룹 GEMM(MoE 시나리오)을 SM89로 확장한다. 네 가지 명시적 구성, 복사가 포함된 여섯 개의 벤치마크 케이스, 테스트는 알 수 없는 키, 잘못된 tile 크기, 마스크 카운트 및 slab 테일 블록, 튜닝, 출력 별칭, 비정렬 입력, 유휴 전문가 유지 및 입력 불변성을 커버하며, 두 가지 TileLang 베이스라인(0.1.12 및 0.1.15 개발 버전) 각각 35개 항목이 통과했다. 케이스는 기존 42개에서 77개로 늘어났으며, 12개의 SM89 케이스가 새로 추가되었다. #2655는 동일 아키텍처에서 INT8 W8A8 GEMM을 활성화하며, 13개의 SM89 인터페이스 케이스와 기존 29개의 INT8 벤치마크 워크로드 모두 vLLM 대상 정확성 검증을 통과했고, SM90 실행 결과는 주장하지 않는다고 명시했다. 두 건을 합치면 TileOPs의 저정밀도 GEMM이 더 이상 Hopper만을 유일한 대상으로 삼지 않음을 의미한다.

출처: tile-ai/TileOPs#2659, #2655 (10-19)

3.2 TileOPs: foundry 연산자 다섯 건이 MLA 페이징, KV 쓰기 및 융합 정규화 보완 (#2642 / #2635 / #2639 / #2634 / #2651)

이 배치의 항목은 모두 TileFoundry의 스펙 파이프라인에서 산출되었으며, 각 연산자는 먼저 HIR 스펙과 참조 구현을 가진 뒤 배치, 스케줄링, 런타임 트윈을 갖는다. #2642는 MLA 페이지드 포워드를 새로 추가했으며, 요청별로 페이지 테이블을 순회하고 fp32 누산으로 스코어링하며 온라인 softmax를 수행하고 출력과 lse를 반환한다. 6개 행의 테스트 케이스는 4e-3 절대 오차 내에서 참조와 정렬되며, 행 내에서 576과 512 열로 두 번의 곱-누산을 판정하여 연산 능력 제한으로 판정된다. #2635는 MLA의 KV 캐시 쓰기를 새로 추가했으며, 토큰별로 kv_c와 k_pe를 연결하고 fp32로 회전한 뒤 디스크에 기록하며, scale에 따라 스케일링하고 클램핑하여 E4M3로 변환한다. 6개 행의 테스트 케이스는 vLLM의 동명 인터페이스를 통과한다. #2639의 QK-Norm-RoPE 융합은 7가지 행/정밀도 조합을 커버하며 NeoX와 인접 쌍 두 가지 회전 레이아웃을 지원한다. #2634의 토큰별 양자화 RMSNorm은 잔차 덧셈을 융합하며, 12개 워크로드가 vLLM 동명 인터페이스와 정렬되고, 한 행에 하나의 CTA, 4096 열 이내에서는 128 스레드, 더 긴 경우에는 512 스레드로 구현된다. #2651은 INT8 W8A8의 연산자 래퍼를 보완한다.

출처: tile-ai/TileOPs#2642, #2635, #2639, #2634, #2651 (10-19)

3.3 TileOPs: GQA 페이지드 FP8 포워드가 implemented로 전환, 36개 사례가 FA3와 비교 (#2648)

이 커밋은 GQA 페이지드 포워드가 FP8의 Q와 KV를 직접 받도록 하며, 연산자 상태를 스펙에서 구현 완료로 끌어올린다. 참조 구현은 요청과 KV 헤드 차원에 따라 HIR을 작성하고, 쿼리 헤드를 그룹으로 쌓아 스코어링 행을 만들며, 각 가시 페이지를 하나의 키 tile로 삼고, FP8 피연산자를 각자의 scale로 역양자화한 뒤 곱한다. 36개 테스트 케이스가 모두 통과하며 NaN과 Inf 검사를 포함한다. 비교 대상은 FA3의 동종 페이지드 구현을 선택했으며, “요청별, KV 헤드별” descala에 따라 해당 형태로 전개된다. 부수 변경은 17개 파일, 2601줄 추가를 포함하며, 페이지드 디코딩, 배치 1 디코딩, FP8 디코딩, warp 특화, RoPE 전처리 등 여러 커널 경로를 포함한다.

출처: tile-ai/TileOPs#2648 (10-19)

3.4 TileOPs: 선형 어텐션과 GEMM의 지점별 성능 마무리 시리즈 (#2650 / #2645 / #2637 / #2636 / #2641 / #2640 / #2646 / #2652)

8개의 [Perf] 커밋은 지난주 마무리 흐름을 이어가며, 패턴이 매우 일관된다. 각 chunk의 행렬 곱을 한 번으로 압축하고, 동기화 지점을 앞당기거나 병합하며, 하드웨어 웨이브에 따라 전략을 결정한다. #2650은 GDN prefill의 chunk 역행렬을 warp 그룹 단위 솔브로 변경하고, “SM당 약 7개 warp 그룹 솔브를 수용”을 기준으로 하며, 한 웨이브 미만일 때는 여전히 단일 warp가 더 유리하다. #2645는 DeltaNet 포워드 상태가 chunk마다 행렬 곱을 한 번만 수행하도록 한다. #2637은 GLA 역방향 상태 그래디언트가 그 체인에서 더 이상 행렬 곱을 삽입하지 않도록 한다. #2636은 배치 GEMM의 각 tile 저장을 다음 tile의 첫 WGMMA 아래로 옮겨 쓰기 반환을 중첩한다. #2641은 긴 행 리덕션에 512 스레드 폴딩을 사용한다. #2640은 최대 풀링 2D 인덱스 커널의 행을 evict-last로 L2에 읽어 들인다. #2646은 스케일 팩터가 정확히 1일 때 bf16 나눗셈의 전 구간 스케일링을 우회한다. #2652는 근사 역수를 사용하여 float32 나머지와 내림 나눗셈의 내림을 구현한다.

출처: tile-ai/TileOPs#2650, #2645, #2637, #2636, #2641, #2640, #2646, #2652 (10-19)

3.5 TileOPs: Mamba-1 선택적 스캔 구현, MoE 라우팅 열거에 수렴 상한 추가 (#2638 / #2653)

#2638 Mamba-1의 선택적 스캔 전방향을 추가했으며, 9개 워크로드가 Mamba-2.8B의 4096 토큰 프리필과 단일 토큰 디코딩, 2.8B 배치 64 디코딩, 1.4B 프리필, 130M 무게이팅, 시퀀스 인덱스가 포함된 패킹 프리필, 그룹 상태와 베어 스캔을 포괄한다. 참조 구현과 레이아웃, 스케줄 생성기가 함께 저장소에 반영되어 상태 공간 모델이 공식적으로 연산자 목록에 진입했음을 보여준다. #2653은 MoE 그룹 GEMM의 라우팅 열거를 수정했다. 이전에는 전체 그룹 라우팅을 열거하여 조합 수가 한도를 초과할 때 불필요한 시간을 소모했으나, 합집합이 상한을 초과하면 열거를 중단하도록 변경했다.

출처: tile-ai/TileOPs#2638, #2653 (10-19)

3.6 TileOPs: 사양 및 벤치마크 커버리지 모두 205/205 보고, 야간 리포트에 연산자 칼럼 추가 (#2656)

조직 사양 통계 브랜치는 마지막 병합 직후 새로고침되었으며, 배지 데이터는 「구현 완료 205/205, 벤치마크 완료 205/205」로 두 항목 모두 만점을 기록하여 이번 기간에 연속으로 보충된 연산자와 상호 검증된다. #2656은 야간 리포트에 두 가지를 추가했다. 해당 기간에 측정되지 않은 연산자를 「새 연산자」 칼럼으로 별도 나열하고, 이전 데이터 항목의 읽기를 구식 문자열 표현에서 JSON 파싱으로 변경하여 매핑되지 않는 내용은 모두 폐기하도록 했다. 테스트 케이스는 43개에서 47개로 증가했다. #2656이 자술한 과거 데이터 한 조각도 기록할 만하다. 10-18의 야간 리포트(통계 기간 10-05~10-17)는 109개 개선, 6개 새 연산자, 총 121행을 나열했으며, 그중 W4A8 GEMM이 98행을 차지했다.

출처: tile-ai/TileOPs#2656, TileOPs 사양 통계 브랜치 (10-19)

3.7 TileFoundry: Apple AMX 지원 및 Qwen 프리필 예제 M2 Pro에서 실행 성공 (#244 / #240 / #245)

#244는 네 가지 AMX 연산(X/Y 로드, FP16 외적에 FP32 누적, Z 저장)을 TileFoundry에 통합하고, Copy와 Fill이 Apple 코어 및 AMX 실행기 스코프에 배치될 수 있도록 했다. 예제 디렉터리는 전체 모델 HIR, 스케줄링 후 HIR, Python 런타임 트윈, C++ 프로젝션 커널을 제공하며, 최초 사용 시 로컬 커널을 빌드하고 캐시하고, 가중치는 미리 준비된 GGUF Q8_0를 재사용한다. 실측 결과 M2 Pro에서 Qwen3-1.7B와 Qwen3.5-4B 모두 128토큰 프리필을 완료했고, 출력 logits는 유한하며 다음 토큰 번호는 각각 1813과 306이었다. 이는 TileLang 생태계가 처음으로 전체 모델 추론 예제를 비CUDA 가속기에서 실행한 사례다. #240은 CUDA 측 네 가지 저하 결함(효과 형식 Cast에 생성 핸들러 없음, Fill이 커널 파라미터 원시 포인터에 기록, 슬라이스 좌표가 Dim 산술로 해석 불가)을 수정했으며, 코드 생성 디렉터리만 변경하여 45행 추가 10행 삭제했다. #245는 분석 계층의 추정 컨텍스트와 스코프 스냅샷을 통일했다.

출처: tile-ai/TileFoundry#244, #240, #245 (10-19)

3.8 채택자 및 야간 파이프라인: TileKernels / FlashQLA 푸시 없음, 스냅샷 10-20으로 전진 (10-20)

deepseek-ai/TileKernels는 여전히 10-14의 v2.1.0 릴리스 병합에 머물러 있고, QwenLM/FlashQLA는 10-14의 수정 커밋에 머물러 있으며, 둘 다 기간 내 푸시가 없어 채택자 측은 이날 조용했다. TileOPs-nightly의 snapshots 브랜치는 기간 내 스냅샷을 한 번 추가했으며(10-20 05:39), 이전 기록은 10-19 05:41로 야간 벤치마크는 하루 한 번의 리듬을 유지했다.

출처: TileOPs-nightly snapshots (10-20)

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

4.1 문서: 여섯 개 새 연산자 API 페이지에 작성, 메인 저장소 문서 사이트 자동 업데이트 (#77)

TileOPs 문서 사이트는 연산자 병합에 따라 여섯 개 새 구현을 API 페이지에 보충했으며(#77), 호출 규약과 현재 상태를 항목별로 명시하여 지난주에 established된 「구현 완료 즉시 페이지 반영」 규칙을 따랐다. tilelang 메인 저장소 문서 사이트는 기간 내 자동 빌드 푸시가 한 번 있었으며 수동 콘텐츠 변경은 없었다.

출처: tile-ai/TileOPs.github.io#77 (10-19)

4.2 커뮤니티 저장소: W4A16 kernel 실습 저장소와 CUDA 디버깅 스킬 저장소가 시야에 진입 (10-19)

커뮤니티 측에서 TileLang을 둘러싼 서드파티 저장소 몇 가지가 등장했다: fanyuxiao-ui/w4a16-tilelang은 10-10에 생성되었으며 TileLang과 수기 CUDA를 사용해 A800(SM80)에서 W4A16 양자화 추론의 GEMM과 GEMV를 구현한다; sablin39/tilelang-cuda-skills는 올해 5월에 생성되어 이미 145개의 스타를 보유하고 있으며, TileLang 작성과 CUDA 툴체인으로 디버깅하는 스킬 패키지를 제공하고, 저자는 동시에 TileFoundry의 활발한 기여자이기도 하다; localizethedocs/tilelang-docs-l10n은 공식 문서 현지화를 진행한다; chaigua/tilelang-agent는 10-19에 새로 생성되었다. 네 가지 모두 커뮤니티 자발적이며 조직 공식 저장소가 아니다.

출처: fanyuxiao-ui/w4a16-tilelang, sablin39/tilelang-cuda-skills, localizethedocs/tilelang-docs-l10n(10-19)

4.3 버전 라인: 메인 저장소 0.1.15 여전히 세대 교체 없음, TileOPs는 계속 0.1.12 고정(10-19)

메인 저장소 버전 파일은 여전히 0.1.15이며, 09-30 릴리스에 해당하고 모니터링 기간 내 새로운 태그나 릴리스가 없다. TileOPs의 의존성 선언은 여전히 0.1.12로 고정되어 있으며, 지난주 0.1.15로 올리려던 PR은 이미 닫혔고, 이번 기간 새로 병합된 SM89 사용 사례도 0.1.12와 0.1.15 개발 버전을 동시에 실행했으며, 이는 업그레이드 게이트가 여전히 이중 베이스라인 검증을 전제로 한다는 것을 보여준다.

출처: tile-ai/tilelang VERSION, tile-ai/TileOPs pyproject(10-19)

4.4 미디어 및 학술: 뉴스 검색 3건 적중 모두 직접 관련 없음, arXiv 새 논문 없음(10-19)

Google News 중영문 여러 쿼리(TileLang, tile-ai, TileOPs, TileFoundry, tilelang 연산자 등)는 모니터링 기간 내 단 3건 적중했다: DeepSeek와 화웨이 칩에 관한 산업 보도 2건(The Information, Business Upturn), 그리고 Prime Intellect의 prime-rl 0.10.0 릴리스 기사; 각 기사를 개별 확인한 결과 읽을 수 있는 내용에서 TileLang 구체적 사실을 찾지 못해 직접 관련 없음으로 처리하고 항목에 포함하지 않았다. Hacker News의 tilelang 검색은 관련 적중이 없었고, arXiv 전체 필드 검색에서 최근 관련 논문은 여전히 07-24의 TileSight(tile 중심 GPU 성능 분석 모델)이며 모니터링 기간 내 신규는 없다.

출처: Google News TileLang 검색, arXiv TileLang 검색(10-19)

5. 트렌드 관찰

5.1 SM89가 새 연산자의 기본 착지점이 되다

어제 TileOPs의 중점은 여전히 Hopper의 W4A8과 SM90 클러스터 협업이었으나, 이번 모니터링 기간 두 건의 새 GEMM 모두 SM89를 최초 목표로 삼았고 PR에서 SM90 실행 결과를 주장하지 않음을 명확히 선언했다. Ada급 그래픽카드(L40S, 4090류)는 국내 개발자들이 보유한 주력 하드웨어로, 저정밀 GEMM을 먼저 SM89에 착지시키는 것은 연산자 라이브러리가 ‘획득 가능성’을 기준으로 우선순위를 정하기 시작했음을 의미하며 단지 ‘피크 성능’만이 아니다.

5.2 선형 어텐션이 ‘chunk당 한 번의 행렬 곱 + SM 웨이브별 전략 선택’으로 고정되다

GLA, GDN, KDA, DeltaNet 네 라인은 이번 기간 총 8건의 성능 변경을 이루었으며, 수법은 두 가지로 수렴된다: chunk 체인의 행렬 곱을 블록당 한 번으로 압축하는 것, 그리고 SM이 수용할 수 있는 warp 그룹 수에 따라 그룹 솔버를 사용할지 단일 warp를 사용할지 결정하는 것. 이러한 ‘먼저 웨이브를 계산한 뒤 커널을 선택’하는 방식은 이전에는 수기 GEMM 라이브러리에서만 보였으나, 이제 연산자 라이브러리의 기본 경로로 진입했다.

5.3 TileFoundry, CUDA를 벗어나다: Apple AMX와 Qwen 프리필

TileFoundry는 줄곧 ‘agent용 연산자 공장’으로 자리매김했으나, 이전의 모든 예시는 NVIDIA에서 이루어졌다. 이번 기간 AMX 지원과 Qwen 전체 모델 프리필이 M2 Pro에서 실행됨으로써, 그 스펙·배치·스케줄 3계층 추상이 CUDA에 묶여 있지 않고 Apple의 행렬 보조 프로세서로 내려갈 수 있음을 보여준다. 이 라인이 계속된다면 TileFoundry의 자리매김은 ‘TileLang의 프론트엔드’에서 ‘크로스 가속기 스펙-투-커널 컴파일러’로 바뀔 것이다.

5.4 관측 가능성이 컴파일러 하드닝의 주축이 되다

이 모니터링 기간의 세 가지 변경은 같은 목표를 가리킨다. 컴파일 산출물을 측정 가능하고 원인 추적이 가능하게 만드는 것이다. lower trace는 호스트와 디바이스 소스 코드를 분리하여 저장하고, Autotuner는 데코레이터 파라미터를 실제로 측정까지 전달하며, Ascend는 호스트 연산자와 디바이스 커널을 msprof 아래에서 대응시킨다. 연산자 측면이 거의 채워진 상황(규격 통계가 이미 205/205로 보고됨)에서, 남은 격차는 연산자를 추가하는 것이 아니라 측정을 통해 찾아야 한다.

5.5 Ascend 연산자 라이브러리가 대화 모델 전용 연산자로 전환

Ascend의 리뷰 대기 큐에 GLM-5.3의 kpool 압축, KDA와 indexer의 역방향, situ_glu 및 그 그래디언트가 등장했는데, 이들은 범용 연산자가 아니라 특정 대화 모델 구조의 구성 요소다. 지난주 PTO 자동 벡터화와 함께, Ascend 라인의 작업 중심은 「TileLang을 Ascend에서 실행하기」에서 「국내 주류 모델의 핵심 연산자를 먼저 Ascend에서 구현하기」로 옮겨갔다.

부록: 자료 및 검증 설명

  • 이 모니터링 기간은 통상적인 24시간 기간으로, 10-19 07:00부터 10-20 07:00까지를 포함하며, 이전 호(10-16 07:00부터 10-19 07:00까지의 3일 기간)와 겹치지 않는다.
  • 조직 저장소 목록과 푸시 상태는 GitHub API로 전량 수집했으며, 29개 저장소 중 9개가 기간 내에 푸시가 있었다. 각 저장소의 메인 브랜치 커밋은 Atom feed로 항목별로 확인했고, 비기본 브랜치(tilelang-ascend의 ascendc_pto, TileOPs-nightly의 snapshots, TileOPs의 stats)는 별도로 검증했다.
  • 병합 및 리뷰 대기 수량은 GitHub 검색 API의 실제 검색 결과(is:merged와 is:open 각 1회, 기간 시작 타임스탬프 포함)에서 가져왔으며, 저장소별 목록을 가져와 누적하는 방식은 사용하지 않아 중복 집계를 피했다.
  • 조직 API의 익명 할당량이 수집 후반부에 소진되어, PR 본문과 변경 통계는 웹 페이지와 패치 파일 읽기로 대체했으며 항목 사실 검증에는 영향을 주지 않았다. 병합 완료 상태는 일괄적으로 is:merged 검색 결과를 기준으로 한다.
  • 규격 및 벤치마크 커버리지 수치는 조직 통계 브랜치의 배지 데이터에서 가져왔으며, 해당 브랜치는 기간 내 마지막 병합 이후 갱신되었고 프로젝트 자체 기술 기준에 속한다.
  • 야간 스냅샷 수량은 snapshots 브랜치 Atom feed의 UTC 시간을 베이징 시간으로 환산한 뒤 기간 내 포함 여부를 판정했다.
  • 뉴스 측 세 건의 적중은 각각 확인했으며, 그중 두 산업 보도는 읽을 수 있는 내용에서 TileLang과의 구체적 연관성을 확인할 수 없어 미검증으로 처리하고 이번 호 항목에 포함하지 않았다. arXiv와 Hacker News는 모두 신규 건수가 없다.
  • 보고서에는 어떠한 로컬 환경 정보도 포함되지 않는다.

전체 출처 목록