모니터링 기간: 2026-11-04 07:00 ~ 2026-11-05 07:00 (베이징 시간, 24시간; 이전 기간 기준점은 11-04 07:00).

출처: GitHub tile-ai 조직 저장소별 점검(29개 저장소, 기간 내 푸시 8개) + 병합 및 심사 중 실제 값 검색(병합 12건, 심사 중 업데이트 42건) + 신규 issue 4건 + 메인 브랜치 커밋별 점검(tilelang 3, TileOPs 7, TileFoundry 1, tilelang-ascend 1, mlir-ascend 2) + 비기본 브랜치 점검(tilelang-hygon의 feat/hcu-scaled-cast 브랜치 4건) + 릴리스 페이지 및 tag 점검(5개 저장소, 기간 내 신규 릴리스 없음) + 채택 측 두 저장소 + Google News 중영문 4개 쿼리(프록시 경유) + Hacker News + arXiv.

이번 호 인덱스

  • 오늘의 중점: 10-14에 병합된 Ascend PTO 프레임워크 변경(#3512)이 집중 정리 기간에 들어서며 하루 만에 서로 독립적인 패치 다섯 개가 등장; TileOPs는 하루에 성능 변경 7건을 병합하며 연휴 이후 가장 밀도 높은 하루를 기록; 메인 저장소의 세 건 병합은 모두 “측정 기준”과 관련되며, L2 플러시가 0 쓰기에서 읽기 전용으로 변경
    1. 핵심 프로젝트 진행
      • 1.1 메인라인: CuTeDSL의 TMA 아키텍처 게이트가 sm_103a와 sm_107a를 수용, B300과 gr100에서 여섯 개 테스트 케이스 복구(#3772)(11-04)
      • 1.2 메인라인: L2 플러시가 0 쓰기에서 읽기 전용 리덕션으로 변경, do_bench에서 fast_flush 파라미터 제거(#3534)(11-04)
      • 1.3 메인라인: AscendC CI의 자동 튜닝 공급 테스트 케이스가 두 기기를 명시적으로 동기화, 비동기 복사 오류가 더 이상 arange 실패로 위장되지 않음(#3784)(11-04)
      • 1.4 심사 중: ThreadSync의 warp 일관성 증명에 warp 인덱스로 조합된 조합 술어 추가(#3787)(11-04)
      • 1.5 심사 중: T.copy 목적 영역 범위 초과가 컴파일 타임 오류로 변경, 전체 복사 fragment의 포함 증명에 스레드 소속 검증 추가(#3786 / #3788)(11-04)
      • 1.6 심사 중: SM120의 FP8 GEMM이 block-scaled MMA로 전환, RTX 5090에서 8192 큐브가 463에서 727 TFLOPS로(#3790)(11-04)
      • 1.7 심사 중: Profiler와 자동 튜닝의 기준 검증이 그룹으로 제출, 여섯 건이 같은 날 출시(#3797 / #3798 / #3799 / #3716 / #3717 / #3718)(11-04)
    1. 멀티 백엔드 적응
      • 2.1 Ascend PTO: 아키텍처 변경 한 건이 네 가지 유형의 회귀를 유발, 서로 독립적인 패치 다섯 개가 같은 날 제출(#3792 / #3793 / #3794 / #3795 / #3796)(11-04)
      • 2.2 Ascend: Cube와 Vector의 코어 비율을 구성 가능하게, 더 이상 1:2로 고정되지 않음(#3783)(11-04)
      • 2.3 Ascend: AutoSchedule의 주기 하한이 강화, Z3 해결 시간 10.04% 감소(#3776)(11-04)
      • 2.4 Ascend: AI Core 예외 덤프가 CANN 9.2.0과 호환, 버전 기준이 9.3에서 하향(#1956)(11-04)
      • 2.5 Hygon: feat/hcu-scaled-cast 브랜치 네 건, E8M0 스케일링 변환과 64 bank의 LDS 읽기 계획(11-04)
      • 2.6 mlir-ascend: BishengIR 툴체인 어댑터가 contrib로 이전(#209)(11-04)
      • 2.7 MetaX와 Moore Threads는 기간 내 푸시 없음(11-04)
    1. 생태계와 채택 측
      • 3.1 TileOPs: 하루 7건 병합, GQA varlen 패킹, ArgMax/ArgMin 스레드당 8벡터, 기본 캐시 전략(11-04)
      • 3.2 TileOPs 심사 중: 일원소별 단일 웨이브 시작, FP8 그룹 GEMM 연속 레이아웃, GEMM 선택적 bias(#2829 / #2824 / #2823)(11-04)
      • 3.3 TileFoundry: FSDP와 HSDP의 학습 시맨틱이 HIR 평가기에 진입(#284)(11-04)
      • 3.4 TileFoundry 심사 중: wgmma 동기화 계획(#286), 별도로 명명 통일 issue 개설(#287)(11-04)
      • 3.5 채택 측과 제3자: TileKernels와 FlashQLA는 푸시 없음, 제3자 저장소 세 곳이 24시간 내 활발(11-04)
    1. 커뮤니티, 튜토리얼 및 활동
      • 4.1 신규 issue 4건: warp 크기 64의 증명, PTO 부호 있는 원자 연산, CuTeDSL의 NVFP4 양자화기(11-04)
      • 4.2 뉴스 측: 4개 쿼리, HN 및 arXiv 모두 기간 내 적중 없음(11-04)
    1. 트렌드 관찰
  • 부록, 자료 및 검증 설명

오늘의 중점

이 날의 주된 흐름은 「무엇을 또 합쳤는가」가 아니라 「어제의 토대가 검증되기 시작했다」였다. 10-14에 머지된 #3512는 T.Parallel을 PTO의 VMI 파이프라인으로 벡터화한 것으로, Ascend 백엔드의 구조적 변경이었다: 정수 타입을 부호 없는 의미의 i*에서 명시적 부호 있는 si*로 바꾸고, SIMD lowering을 새로운 PTO 검증·합법화·벡터화 3단계 흐름이 담당하도록 했으며, 명시적 SIMD 병렬 루프의 판정도 다시 작성했다. 3주 후, 이 변경은 실제 연산자의 형상에서 네 가지 유형의 회귀를 하나씩 드러냈다: 정수 원자적 연산이 pto.si32를 받아들이지 않아 컴파일 실패(#3789 보고, #3792 수정), T.clear가 T.SimdVF 안에서 부동소수점만 지원하는 브로드캐스트 경로로 진입(#3795), 스칼라 cast의 경계 사실이 VMI 메모리 합법화에서 너무 일찍 바인딩됨(#3793), 명시적 SIMD의 병렬 루프가 스칼라 lane 접근으로 간주되어 거부됨(#3796), 그리고 정수 원자를 부호 없는 캐리어에 적응시키는 또 한 건(#3794). 다섯 개의 패치는 모두 메인라인 기반이며 서로 의존하지 않는다. 작성자는 설명에서 같은 문장을 거듭 썼다 — 「This PR is based directly on main and does not depend on the other PTO regression fixes」— 이는 의도된 것이다: 다섯 건이 각각 독립적으로 머지될 수 있게 하여, 과거처럼 거대한 롤백 묶음으로 쌓이지 않도록 한 것이다. 검증도 더 이상 「돌아갔다」가 아니라 구체적인 케이스 수를 동반한 하드웨어 실측이다: 47, 18, 16, 37 항목, 네 개의 독립 연산자(SwiGLU, Top-k, Cache, Engram 역방향)를 대표 케이스로 삼았다.

두 번째는 측정 기준의 자기 수정이며, 두 저장소가 같은 날 같은 일을 했다. 메인 저장소 #3534는 do_bench의 L2 플러시를 「버퍼에 0을 쓰기」에서 「버퍼를 한 번 읽기」로 바꿨다: 0을 쓰면 L2에 더티 행이 통째로 남아 다음 타이밍 구간에 이 행들을 다시 기록해야 하므로, 캐시 플러시 비용이 측정 대상 커널에 전가되는 셈이다. #3534는 cache.view(-1, 4096).sum(dim=1)로 블록 단위 읽기를 수행하고, 덤으로 fast_flush라는 공개 파라미터를 삭제했다 — 그 파라미터는 원래 int32와 int8 중 무엇으로 버퍼를 만들지 선택하는 용도였다. Ascend 측 msprof 백엔드는 더 근본적으로 바꿨다: 기존에는 SimtVF 루프로 UB에 0을 쓴 뒤 WTS_FV로 저장했지만, 이제는 NORMAL_FV에 따라 전역에서 UB로 읽어들이고 타입도 int32로 고정했다. TileOPs의 #2822는 같은 사상의 쌍둥이 변경이며, 추가로 플러시 방법을 기록하고 방법이 일치할 때만 과거 측정값과 비교한다. 세 변경을 함께 보면 하나의 신호다: 성능 수치가 이제 「버전 간 비교 가능」을 요구하기 시작했으며, 단지 「이번이 지난번보다 빠르다」가 아니다.

세 번째는 하드웨어 최전선의 뒤늦은 보충이다. B300이 보고하는 대상 아키텍처는 sm_103a인데, CuTeDSL의 TMA 헬퍼 함수는 화이트리스트로 아키텍처를 가로막고 있어 B300에서 여섯 개의 CuTeDSL 케이스가 DSL에 진입하기도 전에 ValueError를 던졌다. #3772는 sm_103a와 gr100(cc 10.7)이 보고하는 sm_107a를 함께 화이트리스트에 추가했고, 작성자는 B300 노드에서 8 통과 1 스킵, gr100에서 13 통과 1 스킵을 재현했다. 같은 기간에 #3790은 다른 길을 택했다: SM120에서 FP8에 fp32 누산을 더한 GEMM이 계속 SM89의 mma.sync 원자를 사용해 RTX 5090에서 cuBLASLt의 68%에 그쳤는데, ue8m0 단위 스케일링을 갖춘 block-scaled MMA로 교체한 후 8192 큐브가 463에서 727 TFLOPS로 올랐고 출력은 기존 원자와 비트 단위로 일치했다. 두 건의 성격은 다르다 — 하나는 「새 카드가 돌아가게」, 하나는 「새 카드를 최대한 활용하게」 — 하지만 모두 같은 것을 가리킨다: 아키텍처 지원의 경계가 Hopper에서 Blackwell과 국산 가속기 양쪽으로 동시에 확장되고 있다.

네 번째는 TileOPs의 커밋 리듬에 관한 것이다. 단일 일자에 머지된 7건 중 5건은 성능, 1건은 CI, 1건은 문서다: GQA varlen은 SM90에서 KV 그룹별로 요청을 패킹하여 짧은 요청의 key/value 읽기 횟수를 KV 헤드당 한 번으로 줄였다. ArgMax/ArgMin은 한 번에 상주할 수 있는 전제 아래 스레드당 16바이트 벡터 여덟 개를 갖도록 했고, (2048, 4096)의 bf16 ArgMin은 8.704마이크로초에서 7.393마이크로초로 줄어 FlagGems의 8.543보다 빨라졌다. 일원소별 연산과 정규화 계열 커널은 「한 번만 읽는」 피연산자를 기본 캐시 정책으로 되돌리고, 실제로 반복 읽기가 일어나는 곳에만 evict-last/evict-first를 유지했다. 이런 변경들의 공통점은 「더 빠르다」가 아니라 「왜 빠른지 안다」는 것이다 — 각 건마다 비교표를 첨부하고, 여전히 비교 대상에 뒤처지는 행을 표시했다(Softmax의 (32768, 4096)은 121.792에서 131.648마이크로초로 후퇴했다).

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

1.1 메인라인: CuTeDSL의 TMA 아키텍처 게이트가 sm_103a와 sm_107a를 받아들여 B300과 gr100에서 여섯 케이스 복구(#3772)

11-04 12:46(베이징 시간, 이하 동일)에 병합되었으며, 이번 기간 가장 이른 병합입니다. B300 리포트의 대상 아키텍처는 sm_103a인데, CuTeDSL의 TMA 보조 함수는 아키텍처 화이트리스트를 유지하며, 목록에 없는 값을 만나면 DSL 진입 전에 바로 오류를 발생시킵니다. 오류 메시지는 허용 값을 sm_90, sm_90a, sm_100a, sm_110, sm_120, sm_120a, sm_121, sm_121a로 나열합니다. fa69f7fa 메인라인에서 nvidia-cutlass-dsl 4.7.0 및 CUDA 13.3과 함께 사용하면, B300에서 6개 테스트 케이스가 실패합니다. test_tilelang_jit_cutedsl.py의 tma_load 경로 3개(GEMM JIT, 멀티스트림 커널, 동적 형상 각 1개)와 다른 tcgen05 CuTeDSL 테스트 파일의 tma_store 경로 3개 tmem_copy_roundtrip 변형입니다. DSL 자체는 sm_103a를 허용합니다. tcgen05 연산은 sm_100f 패밀리 전체를 허용하므로, 문제는 순전히 이 게이트에 있습니다. 이 항목을 추가한 후 두 파일은 B300(cc 10.3, CUDA 13.3, cutlass-dsl 4.8.1)에서 8개 통과 1개 스킵입니다. 같은 커밋에서 sm_107a도 추가했는데, gr100 브링업 보드(cc 10.7)가 이 값을 보고하며 같은 게이트에 부딪히기 때문입니다. 두 항목을 모두 추가한 후 같은 파일들은 gr100에서 13개 통과 1개 스킵이며, 여기에는 #3685에서 도입된 gather/scatter 파일도 포함됩니다.

기억할 만한 점은 작성자가 밝힌 검증 경계입니다. sm_103 하드웨어가 없어 변경은 또 다른 보고자(0z5a, #3685에서 최초 보고 및 bisect)가 제공한 B300 머신으로 검증했으며, CI는 sm_103a를 커버하지 않습니다. 이런 “코드는 병합, 검증은 머신 차용” 상황은 아키텍처 추격기에서 점점 많아질 것이며, 독자는 어떤 기능이 실측으로 뒷받침되는지, 어떤 것이 단지 화이트리스트에 추가된 것인지 알아야 합니다.

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

1.2 메인라인: L2 플러시가 제로 쓰기에서 읽기 전용 리덕션으로 변경, do_bench에서 fast_flush 매개변수 제거(#3534)

11-04 11:27에 병합되었으며, 7개 파일 변경, 36줄 추가, 49줄 삭제로, 그중 세 곳은 실제 구현이고 네 곳은 따라 변경된 테스트입니다. 기존에는 cache.zero_()였습니다. cache_size MB 버퍼에 0을 써서 L2를 밀어내는 방식인데, 문제는 0을 쓴 뒤 버퍼에 더티 라인이 가득 차서, 다음 라운드에서 피측정 커널이 실행될 때 이 라인들을 메모리에 다시 써야 하고, 이 쓰기 대역폭이 피측정 커널의 시간에 포함된다는 점입니다. 버퍼가 클수록 이 편차는 더 뚜렷해집니다. 이제 cache.view(-1, 4096).sum(dim=1, dtype=cache.dtype)으로 변경했습니다. 4096열 단위로 블록을 나눠 읽어, 행을 L2로 읽어들이고 리덕션 출력을 생성하므로 더티 라인 문제가 없습니다. 블록 분할은 전체 블록 리덕션으로 인한 전역 동기화를 피하기 위한 것이고, dtype 유지는 전체 버퍼의 int64 변환을 유발하지 않기 위한 것입니다. 이 변경은 bench.py의 estimate 루프, torch_bench.py의 event, cupti, cudagraph 세 경로, 그리고 msprof.py의 Ascend 전용 플러시 커널에도 적용되었습니다. 후자는 기존에 SimtVF로 0을 UB에 쓰고 WTS_FV로 다시 전역에 저장하는 전형적인 쓰기식 플러시였는데, 이제 NORMAL_FV로 전역에서 UB로 읽어들이고, 요소 타입도 int8/int32 중 선택에서 int32로 고정되었습니다.

공개 인터페이스의 변화에 주의해야 합니다. do_bench의 fast_flush 매개변수가 삭제되었습니다. 기존에는 int32와 int8 두 버퍼 요소 타입 간 전환(int32가 버퍼 생성이 더 빠름)에 불과했는데, 이제 버퍼가 int32로 고정되어 매개변수가 의미가 없습니다. do_bench(..., fast_flush=False)를 직접 호출하는 스크립트에는 파괴적 변경이므로, 저장소 내 테스트는 이미 동기화되어 수정되었고 외부 사용자에게는 릴리스 노트로 알려야 합니다.

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

1.3 메인라인: AscendC CI의 자동 튜닝 공급 테스트 케이스가 두 장치를 명시적으로 동기화, 비동기 복사 오류가 더 이상 arange 실패로 위장되지 않음(#3784)

11-04 17:01에 병합되었습니다. 현상은 AscendC의 CI 작업에서 자동 튜닝 set-inputs 테스트 케이스의 [other_npu] 파라미터화 항목이 다음 torch.arange() 호출 시 srclen=4, dstlen=4를 보고하는 것이었습니다. 4바이트 memcpy 크기 불일치인데, 이는 정확히 크로스 디바이스 복사 프로브의 형태입니다. 진짜 원인은 크로스 디바이스 복사 실패가 비동기로 보고되어, 오류가 발생 위치보다 한 단계 늦게 드러나면서 오류 지점이 전혀 무관한 arange처럼 보인 것입니다. 수정은 단순합니다. 프로브 앞뒤에 각각 장치 동기화(현재 장치와 npu:other)를 추가해 비동기 오류가 try ... except RuntimeError 범위 안으로 돌아오게 했습니다. 검증은 단일 머신 듀얼 카드로 해당 테스트 케이스 파일을 실행해 네 항목 모두 통과했으며, 크로스 디바이스 복사가 실제로 비동기 실패할 때 [other_npu] 파라미터화 항목은 스킵되고 전체 테스트 케이스가 실패하지 않습니다.

이 변경 자체는 작지만, 지난 호 #3546의 꼬리를 잇는 것입니다. 그 변경은 자동 튜닝의 입력 공급을 현재 디바이스 기준으로 해석하도록 바꾸었고, 검수 시 Ascend 듀얼 카드에서의 회귀가 언급되었는데, 여기서는 동일 시나리오의 CI 상 안정성을 보완합니다. 동일 기능을 두 모니터링 기간에 걸쳐 각각 한 번씩 수정했다는 것은 Ascend 측의 자동 튜닝이 여전히 「돌아가기는 하지만 경계가 불안정한」 상태임을 보여줍니다.

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

1.4 검토 중: ThreadSync의 warp 일관성 증명에 warp 인덱스로 구성된 조합 조건자 추가 (#3787)

이 변경은 지난 호 #3761, #3765가 남긴 동일 영역을 이어받습니다: ThreadSync는 「배리어 참여 스레드가 전체 warp 내에서 일관된다」를 증명해야 하고, 증명하지 못하면 lowering을 거부합니다. 트리거 형태는 192개 스레드에서의 조합 조건자 tx < 160 and tx // 32 % 2 == 0입니다. warp 0, 2, 4는 통째로 참여하고 warp 1, 3, 5는 통째로 불참하며, 어떤 warp도 쪼개지지 않으므로 증명 가능해야 하지만, 메인라인은 곧바로 내부 오류 「참여 스레드가 warp 일관되지 않음」을 던집니다. 작성자는 비교 표를 만들어 결함 범위를 정확히 좁혔습니다: 단일 조건자 tx // 32 % 2 == 0, tx % 64 < 32, tx < 160은 메인라인에서 모두 lowering되며, 「경계 항과 warp 인덱스 모듈로 항을 and로 결합한」 형태만 실패합니다. 반면 tx < 168(실제로 5번째 warp를 쪼갬)은 양쪽 모두에서 올바르게 거부됩니다. 근본 원인은 증명기의 인코딩에 있습니다: 스레드 번호를 두 개의 가상 스레드로 이름을 바꾸고, 산술 솔버로 linear1 / warp_size == linear2 / warp_size가 양측 조건자의 동일성을 도출함을 증명하게 하는데, FloorDiv와 FloorMod의 인코딩은 「상한 + 모듈로」라는 조합에서 두 나눗셈의 관계를 세우지 못합니다. 수정 방식은 threadIdx.x를 warp_index * warp_size + lane으로 재작성하여, 두 가상 스레드가 warp 인덱스를 공유하고 각자 자유 lane을 갖게 함으로써 정확히 「동일 warp」라는 관계를 표현하는 것입니다. 그리하여 warp 인덱스만으로 스레드 번호를 읽는 조건자는 양측에서 동일한 표현식으로 붕괴되어 일관성이 증명되고, 실제로 lane에 의존하는 조건자는 여전히 같지 않으므로 #3786에 나열된 tx < 168은 계속 거부됩니다—증명이 완화된 것이 아니라 누락된 한 부류를 보완한 것입니다.

같은 작성자는 같은 날 #3786도 제출했습니다: T.copy의 목적 영역이 버퍼를 초과할 때 범위 초과 접근을 생성하는 대신 컴파일 시점 오류로 바꾸는 것입니다. 두 변경을 합치면 같은 지향점입니다—「컴파일러가 조용히 잘못된 코드를 생성」하는 것을 「컴파일러가 아니라고 말함」으로 바꾸는 것입니다.

출처: https://github.com/tile-ai/tilelang/pull/3787, https://github.com/tile-ai/tilelang/pull/3786

1.5 검토 중: 전체 복제 fragment의 포함 증명에 스레드 귀속 검증 보완 (#3788)

ProveFragmentContains는 「큰 fragment가 작은 fragment를 커버하는가」를 판단할 때 빠른 경로가 하나 있습니다: 큰 fragment가 완전 복제라면 곧바로 참을 반환합니다. 이 경로는 스레드 귀속을 검사하기 전에 반환하므로 다음과 같은 오판이 발생합니다: 작은 fragment가 스레드 0~127을 소유하고, 큰 fragment는 완전 복제이지만 스레드 0~95만 소유하여 96~127이 전혀 커버되지 않는데도 함수는 포함이 성립한다고 말합니다. check_forward_index를 켜면 완전 복제이지만 물리 슬롯이 다른 두 fragment도 수용됩니다. 그 결과 레이아웃 추론이 호환되지 않는 매핑을 호환되는 것으로 간주하는, 「오류가 조용히 방치되는」 유형이며, 보고자는 #2399를 조사하다가 우연히 발견했습니다. 수정 방식은 먼저 물리 인덱스를 보고, 그다음 작은 fragment의 스레드 매핑을 큰 fragment의 좌표계로 환산하고(유계 스레드 범위 오프셋 고려), 산술 분석기로 복제 범위와 스레드 범위 내에 있음을 증명하며, 증명하지 못하면 거짓을 반환하는 것입니다. 검증은 A100에서 업스트림 메인라인의 새 빌드에 대해, 테스트 기간에 컴파일된 브리지를 통해 실제 C++ 구현을 직접 호출하여 스레드 수 128 대 96의 거부, 96 대 128과 96 대 96의 수용, 물리 슬롯, 비연속 귀속, 유계 범위, 기호 인덱스와 부분 복제를 커버합니다.

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

1.6 검토 중: SM120의 FP8 GEMM을 block-scaled MMA로 교체, RTX 5090에서 8192 큐브 463에서 727 TFLOPS로 (#3790)

이 항목은 이번 성능 측면에서 가장 비중 있는 항목이다. 문제는 명령어 선택에서 발생했다. SM120에서 FP8에 fp32 누산을 적용한 T.gemm은 RTX 5090, 8192 정방형 형상에서 약 465 TFLOPS에 불과한 반면 cuBLASLt는 681이었다. tile과 파이프라인 단계를 조정해도 소용없었는데, 병목이 명령어 자체에 있었기 때문이다. TileLang이 발행하는 것은 SM89 시대의 mma.sync ... f32.e4m3.e4m3.f32였고, 이러한 원자는 이 카드에서 절반 속도밖에 나오지 않는다. 작성자는 레지스터만 측정하는 프로브 세트를 만들었다. e4m3에서 f32로 가는 mma.sync는 512 TFLOPS인 반면, kind::mxf8f6f4.block_scale(단위 스케일링)을 거치면 1021, f16 누산을 쓰는 mma.sync는 1022였다. 차이는 전부 명령어에 있었다. ue8m0 스케일링이 0x7f(즉 2의 0제곱)를 취할 때 block-scaled와 일반 MMA 결과가 동일하므로, TL_SM120_BLOCK_SCALED_MMA_ENABLED가 켜져 있고 대상이 sm_120a/f, sm_121a/f일 때 네 개의 fp32 누산 FP8 디스패치 경로가 block-scaled 원자로 가도록 할 수 있으며, 나머지 대상, f16 누산, 명시적 비 a 접미사 sm_120은 그대로 유지된다. 비트 단위 일관성은 명령어 수준에서 검증했다. e4m3와 e5m2의 네 가지 조합, NaN과 Inf 인코딩을 포함한 무작위 레지스터 내용으로 확인했다.

수치: e4m3 GEMM, bf16 출력, 128×128×128 tile, 3단 파이프라인, 256 스레드 조건에서 4096 정방형은 408에서 673으로(cuBLASLt 530), 8192 정방형은 463에서 727로(681), 4096×14336×4096은 449에서 696으로(498), 256×8192×8192는 348에서 606으로(457) 상승했다. 유일하게 느려진 것은 M=16에 128행 tile을 조합한 형상으로, 31.5마이크로초에서 33.5마이크로초로 퇴보했으며, 작성자는 해당 형상이 대역폭 제한을 받으며 16 또는 64행 tile로 바꾸면 더 빨라진다고 설명했다. 새로 추가된 케이스는 PTX를 직접 검사한다. sm_120a는 반드시 block-scaled 원자를 받아야 하고, sm_89는 반드시 일반 원자를 받아야 한다.

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

1.7 검토 중: Profiler와 자동 튜닝의 기준 검증이 그룹으로 제출되어 여섯 건이 같은 날 나옴(#3797 / #3798 / #3799 / #3716 / #3717 / #3718)

여섯 건 모두 같은 작성자(Afloat16)로부터 나왔으며, 주제는 「검증기와 피검증 대상의 기준」을 정렬하는 것이다. assert_allclose는 이제 커널이 텐서를 줘야 할 위치에서 None을 반환하면 거부한다. 기존 구현은 한쪽이라도 None이면 비교를 건너뛰었기 때문에 빈 값을 반환하는 커널도 「통과」할 수 있었다(#3797). assert_consistent는 기존에 반환값에 바로 zip을 적용했는데, 스칼라 텐서는 TypeError를 던졌고, 출력이 하나 빠지거나 더 많으면 조용히 무시되었으며, 형상이 다르면 행을 자르거나 브로드캐스트로 얼버무릴 수 있었다. 이제는 먼저 출력 목록으로 정규화한 뒤 개수와 형상을 확인하고, 빈 텐서도 포함한다(#3798). CUPTI 벤치마크의 suppress_stdout_stderr는 기존에 Python 스트림에 fileno()를 호출했는데, 노트북과 StringIO는 파일 디스크립터가 없어서 크래시가 났다. 이제는 네이티브 디스크립터 1/2와 Python 스트림을 각각 리다이렉트하고, ExitStack으로 예외 경로에서도 복원을 보장한다(#3799). 나머지 세 건은 자동 튜닝 측에 있다. 대상 딕셔너리를 캐시 키로 쓸 때 문자열 표현을 사용했는데, 삽입 순서가 다른 동등한 딕셔너리가 서로 다른 키를 생성했으므로 정렬된 JSON 직렬화로 바꿨다(#3716). Tensor Core 감지는 우선 명시적 연산 능력을, 다음으로 명시적 또는 현재 대상 아키텍처를, 마지막으로 로컬 장치를 조회하도록 했으며, 그렇지 않으면 GPU가 없는 크로스 컴파일에서 지원하지 않는다고 잘못 보고하게 된다(#3718). Profiler의 동적 형상 치환은 m + 1 같은 복합 차원을 보완하고, 누락된 변수에 대해 TypeError 대신 ValueError를 낸다(#3717).

이 그룹의 기술적 함의는 변경 자체가 아니라 그것들이 드러낸 형상에 있다. 검증기는 오랫동안 「통과하면 정답」이라는 취약한 고리였고, 자동 튜닝의 캐시 키, 대상 해석 같은 기반 시설에 편차가 생기면 「결과를 재현할 수 없음」의 형태로 나타나 원인 귀속이 매우 어려워진다.

출처: https://github.com/tile-ai/tilelang/pull/3797, https://github.com/tile-ai/tilelang/pull/3798, https://github.com/tile-ai/tilelang/pull/3799

2. 멀티 백엔드 적응

2.1 Ascend PTO: 한 건의 아키텍처 변경이 네 가지 유형의 회귀를 불러일으키고, 서로 독립적인 다섯 개의 패치가 같은 날 제출됨(#3792 / #3793 / #3794 / #3795 / #3796)

먼저 배경을 명확히 하자. 10-14에 병합된 #3512([Ascend][PTO] Vectorize T.Parallel to PTO VMI)는 병렬 루프의 벡터화를 기존 SIMD 채널에서 PTO의 VMI 파이프라인으로 이전했다. 정수 타입은 부호 없는 i*에서 명시적 부호 있는 si*로 변경되었고, Lowering은 Verify, Legalize, Vectorize의 세 단계를 거치며, 명시적 SIMD 루프의 판정도 새로 구현되었다. 이 다섯 개의 패치들은 각각 실제 연산자에서 드러난 한 지점에 대응한다:

  • #3792 정수 원자 컴파일 실패: codegen_pto가 정수에 대해 pto.si32 / pto.si64를 발행하지만, PTO의 원자 인터페이스는 부호 없는 i32 / i64만 받아들이므로, T.atomic_add, T.atomic_max, T.atomic_min이 정수 버퍼에서 곧바로 컴파일 실패한다(issue #3789에 대응). 수정 방법은 tl.atomic_* 외부에 호환 래퍼를 두는 것이다. 포인터 요소 타입과 피연산자를 동일 폭의 부호 없는 타입으로 재작성하고, 정수 min/max에는 올바른 부호성을 전달하며, 반환값은 다시 원래 타입으로 변환한다. 테스트 케이스는 15개에서 18개로 확장되어 부호 있는/부호 없는 32/64비트의 덧셈, 뺄셈, 극값 계산, 음수 피연산자, 이전 값 반환, 32비트를 초과하는 카운터를 커버한다.
  • #3794 부호 있는 캐리어 적응: #3792와 같은 근원이지만 다른 경로를 탄다. tl.atomic_*의 의미를 바꾸지 않고, 생성된 PTODSL 호출 경계에서 포인터 뷰와 피연산자를 재작성하며, 주소 공간을 유지하고 부동 소수점 원자는 원래 경로를 유지한다. 대표 테스트 케이스는 SwiGLU의 clamp 카운트(bf16, hidden 2048, clamp 0.5, psum 레이아웃)이며, 18개의 원자 테스트 케이스가 모두 통과한다.
  • #3795 SIMD 채우기 손실: T.SimdVF의 int32 T.clear가 부동 소수점만 지원하는 브로드캐스트 경로로 빠졌다. 원인은 새 PTO 파이프라인이 기존 SIMD pass를 더 이상 실행하지 않는데, 그 pass가 채우기와 병렬 루프를 동시에 담당했기 때문이다. 수정 방법은 채우기만 수행하는 진입점을 추출하여 기존 LowerFill을 재사용하고, 버퍼 버전 정규화 이후이자 PTO 병렬 검증 이전에 배치하며, 여전히 원래의 vdup/vsts 시퀀스를 발행하고 연속 영역, 정렬, 전체 벡터, 다중 버퍼 버전 슬라이스 검사를 유지한다. 16개의 테스트 케이스가 통과하며, 대표 테스트 케이스는 Top-k(512 token, 1 전문가, top-1)이다.
  • #3793 스칼라 cast의 경계 사실이 너무 일찍 바인딩됨: u = uint32(i - 1), i가 [0, 63]에 있는 VMI 주소와 무관한 스칼라는 메모리 합법화에서 오류를 던진다. 서둘러 Bind를 수행하면 과소평가된 구간이 남아 IntSet의 cast 분석에서 다시 오류가 발생한다. 수정 방법은 정수 cast의 표현 가능성을 먼저 검사한 후 사실을 설치하여, 증명 엔진에 안전한 스칼라 정의를, IntSet에 보수적인 수치 구간을 제공하여 cast 끝점을 직접 계산하지 않도록 한다. 47개의 테스트 케이스가 통과하며, 대표 테스트 케이스는 Cache(block_size=64, n=65536)이다.
  • #3796 명시적 SIMD 병렬 루프가 스칼라 lane으로 취급됨: Engram 역방향의 초기화가 Parallel 중첩을 사용하는데, 여기서 명시적 SIMD 저장이 전체 벡터 블록 단위로 진행되며, 새로운 lane 레이아웃 추론이 이를 스칼라 접근으로 간주하여 거부한다(M=2, N=32, L=128). 수정 방법은 명시적 SIMD의 분류를 레이아웃 추론과 병렬 유닛 발견 사이에서 공유하도록 하여, 이러한 중첩은 #3512 이전의 범용 레이아웃 경로를 유지하고, 자동 요소별 유닛은 여전히 새로운 PTO lane 경로를 사용한다. 37개의 테스트 케이스가 통과하며, 대표 테스트 케이스는 Engram 역방향(512 token, hc 4, hidden 2048)이다. 작성자는 동시에 한 가지 제한을 명시했다. FP32/128 lane의 명시적+자동 혼합 테스트 케이스는 소스 코드 수준 커버리지만 생성되었으며, 현재 PTOAS/Bisheng 디바이스 컴파일러가 해당 조합에서 6 KiB를 초과하는 VF 스택 예산으로 인해 디바이스 실행 결과가 없다.

이 다섯 건을 나란히 보면 Ascend 측의 엔지니어링 리듬을 읽을 수 있다. #3512와 같은 구조적 변경이 먼저 착지하고, 이후 실제 연산자별로 경계가 하나씩 드러나며, 각 패치에는 하드웨어 실측과 테스트 케이스 수치가 포함되고 서로 독립적으로 유지된다. “프레임워크를 먼저 병합하고 회귀를 나중에 정리한다”는 이 전략은 이전 라운드의 아키텍처 마이그레이션(#1975)에서 이미 한 번 사용되었으며, 이제는 일상적인 프로세스가 되었다. 대가는 메인 브랜치가 이러한 모니터링 기간에 “알려진 몇 가지가 작동하지 않는” 상태에 놓여 외부 사용자가 마주칠 확률이 낮지 않다는 점이다.

출처: https://github.com/tile-ai/tilelang/pull/3792, https://github.com/tile-ai/tilelang/pull/3793, https://github.com/tile-ai/tilelang/pull/3794, https://github.com/tile-ai/tilelang/pull/3795, https://github.com/tile-ai/tilelang/pull/3796

2.2 Ascend: Cube와 Vector의 코어 비율을 설정 가능하게, 더 이상 1:2로 고정되지 않음(#3783)

Ascend의 AI Core는 Cube와 Vector 두 종류로 나뉘며, AscendC는 mix(1,2), mix(1,1), mix(1,0), mix(0,1) 등의 그룹화를 지원하지만 TileLang은 이전까지 일괄적으로 1:2로 기동해 왔으므로, 이는 하드웨어 자유도의 상한을 고정시켜 버리는 것이었고 향후 코어 비율이 다른 칩을 적응시키기 어렵게 만들기도 했다. 이번 변경은 코어 비율을 명시적 파라미터 또는 디바이스에 따라 추론되는 파라미터로 바꾼다: T.MixedKernel(..., sids=1) 또는 sids=2는 요청된 그룹 크기에 따라 동작하며 능력 검증을 수행하고, 생략하거나 None을 전달하면 현재 디바이스의 vector_core_num / cube_core_num을 사용하며 디바이스 정보를 얻을 수 없으면 2로 되돌아간다. T.Kernel은 기존 인터페이스와 코어 타입 판정을 유지하며, 혼합 커널로 인식될 때는 준비된 벡터 코어 수를 사용하고 단일 코어 경로는 변하지 않으며 명시적 T.Vector(1/2)의 그룹화 의미도 보존된다. 선택된 그룹 수는 sid 도메인, 태스크 할당, 동기화 및 코드 생성까지 일관되게 전달되어 디바이스 중복 조회와 중복 정규화를 피한다. 충돌하는 선언과 디바이스 능력을 초과하는 요청은 즉시 거부된다.

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

2.3 Ascend: AutoSchedule의 주기 하한 강화, Z3 솔빙 시간 10.04% 감소 (#3776)

AutoSchedule은 이분 탐색으로 파이프라인의 개시 간격(II)을 찾는데, 기존 하한은 모든 태스크 II의 최댓값을 취해 느슨했기 때문에 이분 탐색이 몇 라운드 더 필요했다. 이번 변경은 리소스 모델 아래 세 가지 필요 조건을 보강한다: 단일 태스크 경계(P >= max_i II_i 유지), 리소스 클리크 경계(양방향 리소스 충돌 클리크에 대해 잔차 순으로 정렬해 합산하여 P >= sum_i II_i 획득), 자기 의존 경계(S_u - S_u >= wait - d*P에서 P >= ceil(wait/d)를 도출하며 버전 관련 거리는 파이프라인 단 수에 따라 절단). 세 가지의 최댓값으로 이분 탐색을 초기화하고 나머지 로직은 변경하지 않는다. 테스트 환경은 Ascend 950PR, CANN 9.3.0 주간 업데이트 버전이며, 캐시를 끄고 1회 웜업 후 3회 측정을 교대로 수행했다: examples/ascend 전부 컴파일에 통과하고 스케줄링 결과는 변하지 않았으며, Z3 솔빙 단계의 누적 소요 시간이 10.04% 감소했다. 이처럼 「결과는 바꾸지 않고 검색 시작점만 바꾸는」 변경은 위험이 가장 낮으면서 이익은 전역적이다. 컴파일 타임의 솔빙 시간은 모든 사용자가 지불해야 하는 비용이기 때문이다.

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

2.4 Ascend: AI Core 예외 덤프가 CANN 9.2.0과 호환, 버전 기준 9.3에서 하향 (#1956)

이번 변경은 tilelang-ascend 저장소의 ascendc_pto 브랜치에 병합되었으며(11-04 10:18), 이번 기간 Ascend 저장소의 유일한 병합이다. AI Core 예외 덤프 기능은 원래 CANN 9.3.0 이상에 의존한다고 선언되었으나, 작성자는 9.2.0에서 실행 가능성을 검증하고 세 곳을 수정했다: acldumpTensorInfo의 필수 필드 보완(addrType=ACL_DUMP_ADDR_RAW, placement, shapeNum, originShapeNum). 9.2.0은 addrType을 강제 검증하는데 기존 구현은 calloc으로 0을 채워 얻은 ACL_DUMP_ADDR_PTR이 거부되기 때문이다. 빈 주소와 크기 0인 텐서는 건너뛰는데, 9.2.0은 tensorAddr이 NULL인 것을 허용하지 않는다. 등록 인터페이스는 폐기된 aclrtSetExceptionInfoCallback에서 aclrtExceptionInfoCallbackRegister로 교체했다. 도구 측 parse_exception_dump()는 9.2.0이 출력하는 .npy와 구버전의 .bin 두 형식을 모두 호환하며, 중·영문 튜토리얼의 버전 요구 사항도 CANN 9.2.0 이상으로 함께 변경했다. 아직 9.2를 사용 중인 사용자에게 이는 「기능 사용 불가」에서 「사용 가능」으로의 변화다.

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

2.5 Hygon: feat/hcu-scaled-cast 브랜치 4건, E8M0 스케일 변환과 64 bank의 LDS 읽기 계획 (11-04)

Hygon 저장소의 메인 브랜치는 11-02에 멈춰 있고, 이번 모니터링 기간의 움직임은 모두 기능 브랜치에 있으며, 4건의 커밋이 8시간 내에 집중되었다(베이징 시간 16:26부터 20:16까지). 작성자는 Hygon 소속이다(huangteng@hygon.cn 및 zhouweiwang@hygon.cn):

  • feat(hcu): add scaled_cast and unify GEMM LDS read planning: E8M0 스케일링의 FP8/BF8/FP4 변환을 새로 추가하고, 스칼라 및 패킹 두 가지 lowering을 지원하며, src/tl_templates/hcu/scaled_cast.h(223줄)와 lower_scaled_cast.cc를 새로 추가했다. 동시에 GEMM 소비 측 메타데이터를 행렬 읽기 선택과 LDS 전략 간에 공유하도록 하고, 단일 단계 LDS 슬라이싱을 지원하며, 주 변환 파일은 650줄 변경되었다.
  • fix(hcu): align GEMM LDS strategies with 64-bank read phases: bank-word 지오메트리에 따라 plain-copy의 swizzle를 유도하고, 행렬과 두 가지 네이티브 pack 크기가 공유하는 좁은 b16 tile에 대해 명시적 역변환을 제공한다. 혼합 또는 제한된 위상의 비동기 소비 측은 호환되는 공용 저장소를 우선 사용하며, 네이티브 wrap 선택 시 전체 b128 읽기 위상을 포함한다. 12개 파일, 1536줄 추가, 112줄 삭제, 테스트 파일 5개 추가.
  • fix(hcu): prefer mixed LDS wraps proven at most two-way: 혼합 LDS wrap에 “최대 2-way” 증명을 추가하고, 증명이 성립할 때만 우선 사용한다.
  • fix(hcu): use physical read width for mixed wrap checks: 혼합 wrap 검사에 물리적 읽기 폭을 사용하고, 기존에 kPack 기준으로 8바이트 per lane을 추론하던 방식을 제거했다.

주제는 공유 메모리의 bank 충돌과 위상 정렬이다. 이는 Hygon DCU에서 GEMM 성능의 핵심 경로이며, 해당 백엔드가 “컴파일 가능” 단계에서 “마이크로아키텍처 기준 튜닝” 단계로 진입했음을 보여준다. 네 건 모두 기능 브랜치에 있으며 아직 PR을 열지 않았고, 시기상 TileLang 메인라인과 한 차례 동기화한 후 커밋할 가능성이 있다.

출처: https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-scaled-cast

2.6 mlir-ascend: BishengIR 툴체인 어댑터를 contrib로 이전 (#209)

메인 브랜치에 병합(11-04 16:35)되어, BishengIR 툴체인의 어댑터 코드를 기존 위치에서 contrib 디렉터리로 옮겼다. 이는 순수한 구조 정리로, “업스트림 툴체인 진화에 따라 빈번히 수정해야 하는 글루 레이어”를 핵심 lowering과 분리하여, 향후 BishengIR 업그레이드 시 변경 범위가 contrib 내로 수렴되도록 했다. 같은 브랜치에는 refactor(contrib): move BishengIR toolchain adapters into contrib라는 본체 커밋도 있으며, 병합 커밋이 두 번째 커밋이다. 이런 정리는 단독으로는 기능적 가치가 없지만, “어댑터 레이어”를 독립적인 생애주기로 관리하기 시작한 출발점이다. Ascend 툴체인은 버전 반복이 빠르기 때문에 글루 코드의 교체 가능성이 우아함보다 더 중요하다.

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

2.7 MetaX와 Moore Threads, 모니터링 기간 내 푸시 없음 (11-04)

tile-ai/tilelang-metax의 마지막 푸시는 10-26 15:08, tile-ai/tilelang-musa는 09-30 17:50에 멈춰 있으며, 두 저장소 모두 이번 모니터링 기간에 커밋도, 리뷰 중인 PR 업데이트도 없었다. 국산 4개 백엔드 중 Ascend와 Hygon이 이번 모니터링 기간의 두 활성 라인이며, MetaX와 Moore Threads는 정적 기간에 있다. 지난 리포트에서도 MetaX는 푸시가 없었으며, 이로써 두 모니터링 기간 연속이다. Moore Threads는 09-30 이후 움직임이 없어, 향후 모니터링 기간에 내부 개발로 전환했는지 주의할 필요가 있다.

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

3. 생태계와 채택 측

3.1 TileOPs: 하루 7건 병합, GQA varlen 패킹, ArgMax/ArgMin 스레드당 8벡터, 기본 캐시 정책 (11-04)

7건 모두 모니터링 기간 내 병합되었으며, 이번 분기에서 가장 밀도 높은 하루였다. 건별로 살펴보면:

#2819는 SM90의 GQAVarlenFwdOp에 패킹된 warp 특화 프로그램을 추가했다. 하나의 요청에 대한 KV 헤드 그룹이 하나의 tile 내에서 위치별로 채워질 수 있을 때, 요청이 KV 그룹별로 tile을 채우도록 하여 각 KV 헤드가 key/value를 한 번만 읽도록 한다. 패킹 적용 여부는 비용 모델이 판단하며, “형태상 절감이 보장되는 양”이 실측된 9% 단일 헤드 오버헤드를 초과할 때만 사용한다. 넓고 좁은 요청은 동일한 실행 내에서 교대로 처리하며, 호스트 리드백 없이 동기화를 추가하지 않는다. 디바이스에서 집계된 분할 불가 항목이 SM 개수보다 적을 때, 단일 tile 패킹 요청은 key 스캔을 여러 세그먼트로 분할하고 fp32 작업 영역으로 병합한다. 또한 세 가지 메인라인 결함을 수정했다. sink에 점수 스케일링이 약 2.4를 초과할 때 보이는 key가 없는 행이 NaN을 반환하는 문제, 폴백 경로가 요청의 첫 key 이전에 RoPE 테이블을 읽는 문제, 그리고 하나의 네거티브 컨트롤이 호출자 스택 프레임을 참조 순환에 남겨 테스트 간 CUDA graph가 조기에 파괴되는 문제이다. 테스트 노드는 173에서 197로 증가했다.

#2828 행별 블록 분할 ArgMax/ArgMin에서 ‘단일 웨이브 상주’ 전제하에 스레드당 8개의 16바이트 벡터를 보유하도록 하고, 64스레드 블록을 튜닝 후보에 추가했다. 상주 판단은 SM의 스레드 수, 상주 CTA 상한과 레지스터가 공동으로 결정한다. (2048, 4096)의 bf16 ArgMin은 8.704마이크로초에서 7.393으로 감소했고(비교 항목 FlagGems 8.543, 1.155배 빠름), 백투백 시나리오는 7.808에서 6.000으로 감소했다. f16 ArgMax는 6.816에서 6.400으로 감소했다(비교 항목 9.344, 1.46배 빠름). 12개 벤치마크 행 중 4개 행이 빨라졌고 0개 행이 느려졌다.

#2820은 ‘각 원소를 한 번만 읽는’ 오퍼랜드들을 기본 캐시 전략으로 되돌렸다. 원소별, rope, norm, softmax, 리덕션 계열 커널은 더 이상 일률적으로 evict-last/evict-first를 사용하지 않고, 실제로 반복 읽기가 발생하는 곳에만 유지한다(이항 커널의 브로드캐스트 행, 다중 배치 다중 헤드가 공유하는 rope cos/sin 행, 2패스 RMSNorm의 두 번 읽기, 융합 덧셈의 처음 두 패스). 동시에 RMSNorm에 페어링 커널과 단계별 커널 두 가지 구현을 추가했고, warp 행 Softmax의 CTA를 SM을 정확히 채우는 warp 수로 축소했으며, 2D neox RoPE는 128스레드로 전환하고, 아키텍처별 상주 CTA 상한을 RESIDENT_CTAS에 넣었다. 대가는 Softmax의 (32768, 4096)이 121.792에서 131.648마이크로초로 퇴보한 것(비교 항목 quack 126.080)이며, 작성자는 표에 이를 사실대로 표기했다.

#2822는 L2 플러시의 읽기 전용 개정으로, 메인 저장소 #3534와 동일한 소스이며 168개 테스트가 통과했고, 추가로 플러시 방법을 기록하고 방법이 일치할 때만 과거 측정값과 비교한다. #2825는 단일 재적합 비용 모델로 W4A16의 법정 구성들을 정렬하고, split-K의 CTA는 자신의 양자화 메타데이터만 임시 저장하며, 공유 메모리 계정에 시작 여유분을 포함시킨다. #2826은配套 문서로, dense와 tiny-M의 캘리브레이션 계약(측정 방식, 적합 제약, 홀드아웃 검증)을 명확히 기술하며 런타임 변경은 없다. #2827은 off-CUDA 제품군을 야간 작업으로 옮겼다. 기존 preflight에서 이 두 단계는 ubuntu-latest에서 13~18분이 걸리고 다른 검사보다 10분 늦게 결과가 나왔으며, 그중 tests/ops의 스모크 실행은 연산자 참조 구현을 사용해 기기 무관성과 디스패치를 검사할 뿐 커널 결과를 검사하지 않는다. 옮긴 후 preflight에는 수집 검사, 폐기 테스트 보고서와 cuda_only 필수 실패 검사(약 3분)만 남기고, 1856개와 2789개 테스트 노드는 야간으로 전환했으며 릴리스 게이트는 여전히 야간 성공을 요구한다.

출처: https://github.com/tile-ai/TileOPs/pull/2819, https://github.com/tile-ai/TileOPs/pull/2828, https://github.com/tile-ai/TileOPs/pull/2820

3.2 TileOPs 검토 중: 일원 원소별 단일 웨이브 시작, FP8 그룹 GEMM 연속 레이아웃, GEMM 선택적 bias (#2829 / #2824 / #2823)

#2829는 ‘본문에 단일 명령 연산이 한두 개뿐인’ 일원 원소별 시작을 단일 웨이브로 실행되게 한다. 읽기/쓰기 바이트 수가 L2의 4분의 3을 넘지 않고, 결과와 오퍼랜드가 동일 폭이며, 본문에 각각 단일 명령으로 축소되는 연산이 최대 두 개인 세 조건을 모두 충족할 때만 256스레드 블록의 단일 웨이브로 축소한다. 본문 조건은 TIR 변수의 노드 카운트로 판단하며, cast, 나눗셈, 수학 호출, 가드가 있거나 순수 변수로 평가할 수 없는 본문은 기본 시작을 유지한다. 현재 해당하는 것은 relu, abs, neg, bitwise_not과 bool의 logical_not이다. 부수적으로 실제 문제 하나를 수정했다. ReLU가 NaN 입력에 0을 반환하던 것을 이제 NaN을 반환한다(16비트는 torch의 비트 패턴, float32는 표준 NaN 사용). 벤치마크에서 세 행의 ReLU와 한 행의 LogicalNotFwdOp가 빨라졌고 나머지 구성은 변동이 없으며, 작성자는 GPU 6에서 플러시가 있는 decode 행이 단지 동률임을 표기했다.

#2824는 SM90의 FP8 그룹 GEMM에 연속 레이아웃(컴팩트와 정렬 두 가지 패킹, 세그먼트 끝 또는 행별 두 가지 전문가 메타데이터)을 추가했고, SM89는 여전히 마스크 레이아웃만 지원하며, 전문가 크기 편향에 따른 실측 부하와 DeepGEMM 연속 레이아웃 비교를 보완했다. #2823은 GEMMFwdOp에 선택적 bias[N]을 추가하여 dense, GEMV, split-K 세 경로를 커버하고 cuBLASLt BIAS 비교를 첨부했다. 두 변경 모두 #2130에 나열된 격차에 대응하며, TileOPs가 무작위로 연산자를 추가하는 것이 아니라 ‘torch 생태계와 항목별 정렬’ 목록에 따라 추진되고 있음을 보여준다.

출처: https://github.com/tile-ai/TileOPs/pull/2829, https://github.com/tile-ai/TileOPs/pull/2824, https://github.com/tile-ai/TileOPs/pull/2823

3.3 TileFoundry: FSDP와 HSDP의 학습 의미론을 HIR 평가기에 도입 (#284)

11-04 16:36에 병합되었습니다. 기존 평가기는 reshard를 전역 텐서로 처리하여 서로 다른 참여자의 기여 차이를 무시했기 때문에, 훈련 의미론(특히 샤딩과 복제가 혼합된 HSDP)을 HIR 계층에서 올바르게 평가할 수 없었습니다. 이 구현 #255의 훈련 하위 집합은 참여자별 로컬 값, GPU 좌표, 호출과 루프, gather/scatter, 그리고 축 방향의 Partial(sum) 리덕션을 모두 공유된 프로젝션 규칙으로 처리합니다. Reshard의 앨리어스 증명은 소유권에 민감하도록 변경되어 #170의 구조화된 수명 주기와 연결되고, 튜플 실인수는 텐서 리프에서 검증됩니다. 이와 함께 modded-nanogpt에서 영감을 받은 2블록 GQA 연구를 추가했습니다: 명시적 순전파/역전파, 모멘텀 SGD, 독립적인 PyTorch 참조, FSDP, 2×2 및 4×4 HSDP, 두 가지 상주 전략을 포괄하며, 집합 통신 그룹, 페이로드 및 저장 점유를 보고하고 재사용 가능한 SVG를 삽입했으며, README의 모든 표를 코드에서 도출할 수 있습니다.

검증: 리팩토링된 코어 스위트 775개, 백엔드 스위트 778개(13개 스킵, sm90 카드 2장), 10개 HIR 구성이 CPU에서 통과했습니다. 두 가지 상주 전략은 4장의 sm90에서 실제 2×2 PyTorch HSDP와 정렬되었습니다(업데이트 2회 실행, 로컬 그래디언트, 파라미터 및 모멘텀 비교).

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

3.4 TileFoundry 검토 중: wgmma 동기화 계획(#286), 별도로 명명 통일 issue 개설(#287)

#286은 WGMMA의 비동기성을 다룹니다: 명령이 발행된 후에도 공유 메모리와 레지스터를 계속 읽을 수 있으므로, 후속 접근은 그 효과 뒤에 정렬되어야 합니다. 방식은 각 피연산자 효과에 완료 조건을 선언하고, WGMMA의 fence, commit, wait를 필요한 HIR 노드에 연결하며, lowering이 명령의 발행 지점에서 이러한 단계를 소비합니다. 만약 후속 스코프가 집합 대기를 수용할 수 없으면, 발행 스코프에서 연산을 완료합니다. 동시에 레지스터 Reshard의 앨리어스를 수정하여 각 요소가 복제본을 보유한 모든 스레드에서 동일한 레지스터에 유지되도록 하고, 이 저장소 식별로 의존성을 추적할 계획입니다. 작성자는 명확히 한 가지 위험을 명시했습니다: 행렬 명령의 발행기가 아직 연결되지 않았고, 이는 WGMMA의 실행 정확성을 증명하지 않으며 성능도 제공하지 않고, TMA의 완료 의미론은 별도로 처리됩니다.

#287은 문서 지향 issue로, 코드베이스에서 “하나의 개념에 다섯 개의 이름”(mesh의 한 멤버가 placement 133곳, program 110곳, instance 92곳, participant 44곳으로 불리며, wave_units 같은 unit이 붙은 표기도 있음)을 지적하며, unit이라는 단어가 빈 타입, 하드웨어 연산 유닛, mesh 멤버를 동시에 지칭한다는 점, 그리고 Python과 CUDA 런타임의 local_layout 시그니처가 일치하지 않는다는 점을 지적합니다. 이러한 issue는 보통 즉시 코드를 변경하지 않지만, 컴파일러 프로젝트가 “작성 가능”에서 “장기적으로 읽힐 수 있음”으로 가는 데 필수적인 단계입니다.

출처: https://github.com/tile-ai/TileFoundry/pull/286, https://github.com/tile-ai/TileFoundry/issues/287

3.5 채택자와 서드파티: TileKernels와 FlashQLA 푸시 없음, 세 개의 서드파티 저장소가 24시간 내 활성(11-04)

두 채택자 저장소는 이 모니터링 기간에 제출이 없습니다: 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에서 겪은 배리어 점유 회귀로, 컴파일러 변경이 “수동적 회귀” 방식으로 그들의 형태에 작용할 수 있음을 보여줍니다.

서드파티 측에는 저장소 발견으로 수집된 세 가지 활성 항목이 있으며, 모두 24시간 내에 생성되거나 푸시되었습니다: sablin39/tilelang-cuda-skills(149 스타, TileLang 작성과 CUDA 툴체인 디버깅을 위한 스킬 세트 제공), xwhzz/tilelang-hexagon(2 스타, TileLang을 Qualcomm Hexagon DSP로 이식), Index4NekoMeteor/Tilelang-AttnRes(1 스타, TileLang으로 Kimi K3의 어텐션 잔차 구조 구현). 세 가지 모두 작지만 방향이 흥미롭습니다: 하나는 툴체인 측 개발 경험을, 둘은 새로운 대상 아키텍처와 최전선 모델 구조의 이식을 다루며, 커뮤니티가 자발적으로 TileLang의 경계를 확장하는 유형의 작업입니다.

출처: https://github.com/deepseek-ai/TileKernels, https://github.com/QwenLM/FlashQLA, https://github.com/sablin39/tilelang-cuda-skills

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

4.1 신규 issue 4건: warp 크기 64의 증명, PTO 부호 있는 원자 연산, CuTeDSL의 NVFP4 양자화기(11-04)

주 저장소의 세 가지 항목. #3791은 「[BUG][ice-on-valid-code]」: ThreadSync가 warp 크기 32에서만 warp 일관 술어를 증명할 수 있고, 동일한 IR이 thread_warp_size가 64일 때 거부된다——이는 이번 기간 심사 중인 #3787과 매우 밀접하게 맞물려, 증명기가 warp 크기에 대해 가정하는 바가 하드코딩되어 있음을 보여주며, AMD와 일부 국내 가속기는 64폭 warp를 사용한다. #3789는 PTO 백엔드가 정수 T.atomic_add에 대해 부호 있는 pto.si32/si64를 발행하여 검증에 실패한다고 보고했고, 당일 #3792와 #3794 두 방향의 수정안이 나왔다. #3781은 CuTeDSL의 NVFP4 block-scaled 양자화기가 컴파일되지 않는다고 보고한다: FP4 패킹의 재해석과 누락된 ceil/floor. 여기에 TileFoundry의 #287(명명 통일)을 더해, 이번 모니터링 기간 총 4건의 신규 이슈가 있었으며, 전부 「내부 구현 세부사항이 사용자 가시면으로 누출된」 유형이고 성능 불만은 한 건도 없다.

출처: https://github.com/tile-ai/tilelang/issues/3791, https://github.com/tile-ai/tilelang/issues/3789, https://github.com/tile-ai/tilelang/issues/3781

4.2 뉴스 측: 네 그룹 쿼리, HN과 arXiv 모두 모니터링 기간 내 적중 없음 (11-04)

Google News의 네 그룹 쿼리(영문 TileLang GPU kernel, "TileLang" OR "tile-lang", 중문 TileLang 算子, TileLang 编译器)에서 최신 적중은 352시간 전(10-21 SemiAnalysis의 화웨이와 DeepSeek 보도)에 멈춰 있고, 나머지는 09-30의 「DeepSeek, 승등 버전 TileLang·DeepGEMM·DeepEP 오픈소스 공개」 보도 물결(856시간 전)에 집중되어 있으며, 모두 이번 모니터링 기간보다 이르다. Hacker News에서 tilelang으로 최근 24시간을 검색한 결과는 0건이고, arXiv의 all:tilelang 최신 적중은 여전히 07-24의 TileSight이며 모니터링 기간 내 신규 논문은 없다. GitHub 저장소 탐색으로 건진 9건의 후보 중 7건은 저장소 자체, 1건은 Swarm World 같은 주제와 무관한 Show HN, 1건은 무관한 2048류 프로젝트다. 이번 모니터링 기간 뉴스 측은 영(零)임을 확인할 수 있으며, 보고서 본문은 커밋과 PR에서만 나올 수 있다——이는 이미 여러 기간 연속으로 안정적인 형태다.

출처: https://news.google.com/, https://news.ycombinator.com/, https://arxiv.org/

5. 트렌드 관찰

첫째, 승등 PTO의 정산 패턴이 이미 정형화되었으며, 그 이익과 대가 모두 기록할 가치가 있다. #3512는 10-14에 머지되었고, 3주 내내 실제 연산자들이 계속해서 그 경계를 밟아 나갔으며, 오늘의 다섯 건 보완은 최신 배치일 뿐이다. 이러한 「먼저 아키텍처를 내려놓고, 이후 연산자별로 회귀를 정산하는」 리듬은 #1975 전 저장소 마이그레이션 때도 나타났다. 장점은 아키텍처 반복이 빠르고 단일 변경을 통제할 수 있다는 것이고, 단점은 메인라인이 상당히 긴 기간 동안 「알려진 몇 곳이 작동하지 않는」 상태에 놓인다는 것이다. 다섯 개의 패치가 의도적으로 서로 의존하지 않게 유지된 것은, 메인테이너가 롤백 단위를 의식적으로 줄이고 있음을 보여준다——이는 하나의 큰 롤백 묶음으로 모으는 것보다 성숙한 방식이지만, 동시에 더 강력한 릴리스 전 회귀 규율을 요구한다는 의미이기도 하다. 그렇지 않으면 「독립적으로 머지 가능」이 「독립적으로 누락 가능」으로 변할 수 있다.

둘째, 성능 측정이 제품 특성으로 취급되기 시작했다. 주 저장소 #3534와 TileOPs #2822는 같은 날 L2 플러시를 0 쓰기에서 읽기 전용으로 바꾸었으며, 이유는 모두 「더티 라인 쓰기 되돌림이 다음 라운드 타이밍을 오염시킨다」는 것이었다. TileOPs는 더 나아가 플러시 방법을 기록하고, 방법이 일치할 때만 과거 판독값을 비교한다. 이는 작은 수정이 아니다: 기준 수치가 버전 간에 비교 불가능하다면, 모든 「얼마나 빨라졌는가」의 결론을 다시 계산해야 한다. 측정 방법 자체를 버전 관리에 포함시키는 것은 성능 엔지니어링이 신뢰성으로 나아가는 전제다.

셋째, 연산자 라이브러리의 경쟁점이 「연산자 수」에서 「선택기 품질」로 이동하고 있다. TileOPs가 오늘 머지한 W4A16 단일 재적합 비용 모델과 조밀 GEMM의 캘리브레이션 계약 문서, 그리고 심사 중인 FP8 그룹 GEMM 연속 레이아웃과 GEMM bias는 하나의 명확한 라인을 구성한다: 먼저 각 형상의 유효 구성을 열거하고, 통일된 비용 모델로 정렬한 뒤, 캘리브레이션 과정을 재현 가능한 계약으로 작성한다. 이 라인이 제대로 되면 연산자 라이브러리의 성능은 더 이상 개별 커널의 수동 튜닝에 의존하지 않게 되고, 제대로 되지 않으면 「특정 형상은 유난히 빠르고 나머지 형상은 운에 맡기는」 상태로 퇴화한다.

넷째, 아키텍처 지원의 경계가 양쪽으로 확장되고 있다. 한쪽은 신형 카드다: B300의 sm_103a, gr100의 sm_107a가 화이트리스트에 들어가고, SM120의 FP8이 block-scaled MMA로 교체된다. 다른 한쪽은 국산 가속기다: 하이광은 64 bank의 LDS 페이즈에서 마이크로아키텍처 수준 튜닝을 하고, 승등은 코어 비율과 예외 덤프 같은 엔지니어링 세부사항을 보완한다. 이 두 확장 방식의 작업량 구조는 완전히 다르다——신형 카드는 주로 「경로가 통과되게 만드는 것」이고, 국산 카드는 더 나아가 「동일한 경로가 하드웨어에 근접한 효율을 내게 만드는 것」이며, 후자는 훨씬 높은 실측 밀도를 요구하며, 바로 그것이 하이광의 네 건 커밋에서 테스트 파일 비중이 그렇게 큰 이유다.

마지막 리스크 포인트: 한 모니터링 기간 내에 42건의 리뷰 중인 PR 업데이트가 있었고, 그중 12건은 Ascend/PTO 태그, 7건은 CUDA/SM/CuTeDSL 태그, 4건은 Profiler에 집중되어 있습니다. 이 구조는 저장소의 변경 압력이 「기능 추가」에서 「경계 수정」으로 이동하고 있음을 보여줍니다. 사용자 입장에서 좋은 소식은 결함이 하나씩 지목되고 있다는 것이고, 나쁜 소식은 이러한 결함 대부분이 검증기, 캐시 키, 증명기처럼 「평소에는 보이지 않는」 곳에 위치해 있어, 그것들이 드러날 때는 이미 결과의 신뢰성에 영향을 미친 뒤인 경우가 많다는 점입니다.

부록, 자료 및 검증 설명

출처 검증표

출처 검증 결과
tile-ai 조직(29개 저장소) 모니터링 기간 내 8개 저장소에 푸시 발생
메인 저장소 tilelang 메인 브랜치 3건 머지, 신규 이슈 3건, 리뷰 중 업데이트 38건(Ascend/PTO 12, CUDA 및 CuTeDSL 7, Profiler 4)
TileOPs 메인 브랜치 7건 머지, 리뷰 중 업데이트 3건, 신규 이슈 없음
TileFoundry 메인 브랜치 1건 머지, 리뷰 중 업데이트 1건, 신규 이슈 1건
tilelang-ascend 기본 브랜치 ascendc_pto 1건 머지, 리뷰 중 목록은 익명 인터페이스 한도로 가져오지 못해 메인 브랜치와 커밋 원자 출처로 검증 대체
tilelang-mlir-ascend 메인 브랜치 2건(#209 머지 포함)
tilelang-hygon 메인 브랜치 신규 커밋 없음(11-02에 정지), feat/hcu-scaled-cast 브랜치 4건
tilelang-metax / tilelang-musa 모니터링 기간 내 푸시 없음
TileOPs-nightly 브랜치 푸시만 있으며 메인 브랜치 신규 커밋 없음
릴리스 페이지(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) 모니터링 기간 내 신규 버전 없음, tilelang 최신은 여전히 09-30의 v0.1.15
FlashQLA / TileKernels 모니터링 기간 내 푸시 없음
Google News(중영문 4개 쿼리) / Hacker News / arXiv 모니터링 기간 내 실질적 적중 없음, 최신 적중은 352시간 전

전체 출처 목록

*[24]** tilelang-ascend 예외 덤프가 CANN 9.2.0과 호환됩니다(#1956) — https://github.com/tile-ai/tilelang-ascend/pull/1956

검증 설명

이번 호의 모든 병합·심사 중·이슈·커밋 수량은 GitHub 검색 및 원자 출처의 실측 출력에서 나온 것으로, 추정치를 사용하지 않았다. Ascend 저장소와 Hygon 저장소의 심사 중 PR 목록은 익명 인터페이스 할당량 소진으로 확보하지 못했으며, 관련 결론은 메인 브랜치 커밋과 브랜치 수준 원자 출처에만 근거하였고 출처 검증표에 명시해 두었다. tilelang-hygon의 네 건 커밋은 PR이 열리지 않은 기능 브랜치에 위치하여 브랜치 수준 동적에 속하며, 메인라인에 이미 병합되었음을 의미하지 않는다. Google News, Hacker News, arXiv 세 뉴스 채널은 본 모니터링 기간에 모두 히트가 없었으므로, 리포트 본문은 전부 코드 측에서 나왔다.