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

출처: GitHub tile-ai 조직 저장소별 점검(tilelang / TileOPs / TileOPs-nightly / TileFoundry / tilelang-hygon / tilelang-metax / tilelang-ascend / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 문서 사이트 및 야간 벤치마크 저장소) + 머지 및 심사 중 실값 검색(머지 20건, 심사 중 업데이트 45건) + 신규 issue 전수 확인(7건) + 커밋 수준 diff 점검 + 릴리스 페이지 및 PyPI 아티팩트 점검 + 외부 백엔드 저장소 점검(Hexagon / CANN 연산자 저장소 / PPU 래퍼 / 모델 스니펫 저장소) + Google News 중영문 다중 쿼리(프록시 경유) + Hacker News + arXiv.

이번 호 인덱스

  • 오늘의 핵심: TileOPs 0.1.0 정식 릴리스, 206개 연산자, 2265행 벤치마크 중 2262행이 최고 속도 구현보다 느리지 않음; 24시간 내 20건 병합, 심사 중 업데이트 45건, 신규 issue 7건; Hygon scale 버퍼 듀얼 bank 1389행 병합; 서드파티 Hexagon 백엔드 하루 15건
    1. 핵심 프로젝트 진행 상황
      • 1.1 메인 저장소: ROCm의 exp/sin/cos가 더 이상 기본적으로 빠른 내장 함수로 낮춰지지 않으며, sin/cos 오차가 7.8e-04에서 6.6e-08로 감소 (#3621) (10-27)
      • 1.2 메인 저장소: 자동 타깃이 현재 HIP 디바이스를 읽도록 변경, 멀티 GPU 머신에서 더 이상 0번 카드 아키텍처로 고정되지 않음 (#3594) (10-27)
      • 1.3 메인 저장소: TMA gather/scatter의 상수 행 인덱스를 int32로 변환, 리터럴 표기법이 더 이상 컴파일 크래시를 일으키지 않음 (#3365) (10-27)
      • 1.4 메인 저장소: NVRTC와 CuTeDSL이 첫 번째 입력 텐서가 위치한 디바이스 기준으로 현재 스트림을 가져옴 (#3359) (10-27)
      • 1.5 메인 저장소: 자동 튜닝 캐시 키에 compile_flags 포함, 동일 커널에서 컴파일 옵션 변경 시 더 이상 이전 바이너리를 재사용하지 않음 (#3338) (10-27)
      • 1.6 메인 저장소: ThreadSync가 address_of를 쓰기로 기록, 불투명 호출 이후 펜스 보완 (#3346) (10-27)
      • 1.7 메인 저장소: 런타임 행 수의 파이프라인 epilogue 직접 전개, 페이지드 디코딩 컴파일 복구 (#3356) (10-27)
      • 1.8 메인 저장소: 슬라이스 패딩과 음수 인덱스의 정적 경계 두 건 수정, 범위 초과 쓰기와 무음 음수 주소 비활성화 (#3352 / #3347) (10-27)
      • 1.9 메인 저장소: T.ws(-1) 영역이 모든 스레드에서 실행, 단독 표기법에서 더 이상 저장 유실 없음 (#3354) (10-27)
      • 1.10 메인 저장소: profiler에서 디바이스 미지정 시 현재 디바이스 사용, 캐시 플러시가 더 이상 0번 카드에 적용되지 않음 (#3357) (10-27)
      • 1.11 신규 결함: 커널 캐시 히트가 콜드 컴파일을 가정한 여섯 개 테스트를 깨뜨림; CuTeDSL의 CUDA graph 캡처가 빈 그래프가 됨 (#3694 / #3693) (10-27)
    1. 멀티 백엔드 적응 (Ascend / Hygon / ROCm / Hexagon)
      • 2.1 Hygon: HCU scale 버퍼를 듀얼 bank로 분할, blockscaled GEMM이 수동 warp 특수화 하에서 동일한 행 매핑 세트를 공유 (#29) (10-27)
      • 2.2 Ascend 메인 저장소: PerCoreTask 후보 본문에서 내부 if/else 허용, 조건부 L1 패딩이 더 이상 전체 거부되지 않음 (#3686) (10-27)
      • 2.3 Ascend 메인 저장소: CI가 AscendC와 PTO를 두 개의 매트릭스 작업으로 분리, 각각 재실행 가능 (#3687) (10-27)
      • 2.4 Ascend 저장소 심사 중: Cython 래퍼가 분석기를 지연 생성, 어제 기록된 호출별 재생성 문제를 직접 겨냥 (#1948) (10-27)
      • 2.5 Ascend 저장소 심사 중: GLM-5.3의 kpool 압축 연산자, PTO 상의 텐서 시프트량 (#1910 / #1949) (10-27)
      • 2.6 Ascend 메인 저장소 심사 중: PTO의 T.Parallel 벡터화, AutoSimVF 채널, 배치 3D ND2NZ, 쓰기 카운트 멀티 버퍼, Bisheng 빌드 아이덴티티를 캐시 키에 포함 (#3512 / #3320 / #3689 / #3629 / #3691) (10-27)
      • 2.7 ROCm 심사 중: pip 버전 ROCm 10의 아키텍처 탐지와 이미지, FP4 변환과 E8M0 타입 (#3549 / #3665) (10-27)
      • 2.8 Ascend 저장소 일일 테스트: 2903항목 전체 통과, 이틀 연속 실패 없음 (#1950) (10-28)
      • 2.9 MetaX: 조직 뷰에 푸시 타임스탬프 존재, 여섯 개 브랜치 모두 모니터링 기간 내 신규 커밋 없음 (10-27)
      • 2.10 서드파티: Hexagon DSP 백엔드 하루 15건, HVX 벡터화, VTCM 아레나, Crouton 레이아웃 및 DMA 내장 함수 완비 (10-27)
    1. 생태계와 도입 측 (TileOPs / TileFoundry)
      • 3.1 TileOPs 0.1.0 릴리스: 206개 연산자의 첫 정식 버전, 공개 API 확정 (10-27)
      • 3.2 TileOPs 0.1.0 벤치마크: H200에서 2265행 중 2262행이 최고 속도 구현보다 느리지 않음, 기하평균 1.37배 (10-27)
      • 3.3 TileOPs 0.1.0의 알려진 격차와 다음 단계: 융합 연산자, 전체 모델, 연산-통신 융합 (10-27)
      • 3.4 TileOPs 4건 병합: where와 masked_fill 튜닝 가능, D128 인과 대각선을 64 키 기준으로 스코어링, 희소 장문 GDN 프리필 분할, 변경 로그 릴리스 날짜 기준 날짜 지정 (#2782 / #2783 / #2779 / #2781) (10-27)
      • 3.5 TileFoundry 2건 병합: 전체 윈도우를 오퍼랜드 자체 기준으로 전송, 동적 공유 메모리를 커널 실제 사용량 기준으로 크기 결정 (#273 / #271) (10-27)
      • 3.6 외부 생태계: CANN 커뮤니티의 TileLang 연산자 저장소, PPU 버전 무수정 래퍼, DeepSeek V4.1 Flash 스니펫 예제 (10-27)
    1. 커뮤니티, 튜토리얼 및 활동
      • 4.1 문서 사이트 하루 6건 자동 업데이트 (10-27)
      • 4.2 Ascend 저장소의 문서 팩트 체크 4건, 단일 코멘트 최다

46 라운드(#1628 / #1655 / #1656 / #1659)(10-27)

  • 4.3 심사 중: lower trace를 독립 개발 툴킷으로 분리(#3583)(10-27)
  • 4.4 뉴스 측면: Google News와 HN에서 24시간 내 이번 모니터링 기간의 신규 보도 없음(10-27)
    1. 트렌드 관찰
  • 부록. 자료 및 검증 설명

오늘의 핵심

TileOPs 0.1.0이 오늘 모니터링 기간 내에 릴리스되면서, 여러 날 이어진 “마지막 한 칸에 걸린” 상태가 끝났다. PyPI의 0.1.0 아티팩트는 10-27 07:45(베이징 시간)에 업로드되었으며, 이전 데일리 리포트의 마감 시점보다 45분 늦다. GitHub 릴리스 페이지는 13:15에 게시되었고, 저장소에 처음으로 정식 tag가 등장했다. 릴리스 노트의 핵심 지표는 세 가지다. 206개 연산자 전부에 manifest 항목이 있으며(시그니처, 테스트 및 벤치마크 워크로드, roofline 공식, CI에서 항목별 검증), H200에서의 2265줄 야간 벤치마크 중 2262줄이 가장 빠른 비교 구현보다 느리지 않았고 기하평균 1.37배를 기록했다(비교 대상은 torch.compile, FlashAttention 3, FlashInfer, vLLM, DeepGEMM, FlagGems). 206개 연산자 중 195건의 foundry 태그가 붙은 병합은 TileFoundry가 커널을 생성하고 동일한 CI를 통해 저장소에 반영된 것이다. 릴리스와 함께 공개 API가 동결되어 이후 0.x에서는 추가만 가능하고 변경은 불가하며, API 호환성 검사 스크립트가 새로 추가되어 이전 버전 표기법으로 호출하는 PR은 즉시 실패한다.

24시간 내 tile-ai 조직에서 20건이 병합되었고(tilelang 13, TileOPs 4, TileFoundry 2, tilelang-hygon 1), 심사 중 업데이트 45건, 신규 issue 7건(그중 2건은 Ascend 저장소의 일일 테스트 봇)이 열렸다. 병합의 중심은 두 갈래다. 하나는 부동소수점 의미론과 실행 컨텍스트로, ROCm은 일곱 개 기본 수학 함수를 빠른 내장에서 정밀 구현으로 되돌렸고(MI300X에서 sin/cos 오차가 7.8e-04에서 6.6e-08로 감소), 자동 타깃이 고정된 0번 카드 대신 현재 디바이스를 읽도록 변경했으며, NVRTC와 CuTeDSL이 입력 텐서가 위치한 디바이스 기준으로 스트림을 가져오도록 하고, 자동 튜닝 캐시 키에 compile_flags를 추가했다. 다른 하나는 정적 경계와 스레드 도메인으로, 슬라이스 채우기를 길이가 아닌 영역 끝 기준으로 검증하고, 범위를 벗어난 음수 인덱스는 컴파일 시점에 거부하며, T.ws(-1) 영역이 모든 스레드에서 실행되도록 변경했다. Hygon은 12개 파일 1389줄을 병합하여 scale 버퍼를 독립적인 이중 bank로 분리하고 LDS bank와 분리했다. 서드파티 측에서는 Hexagon DSP 백엔드가 하루 15건으로, HVX 벡터화, VTCM 아레나, Crouton 레이아웃, DMA 내장을 Crouton 상의 매핑과 리덕션까지 이어서 구축했다.

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

1.1 메인 저장소: ROCm의 exp/sin/cos가 더 이상 기본으로 빠른 내장에 저하되지 않음, sin/cos 오차 7.8e-04에서 6.6e-08로 감소(#3621)

10-27 22:39에 병합, 5개 파일 +107/-17, #3618 수정. HIP에서 T.exp, T.exp10, T.log, T.log2, T.log10, T.sin, T.cos 일곱 개 기본 함수는 tl.enable_fast_math 설정과 무관하게 __expf 같은 빠른 내장으로 저하되었으나, CUDA 측은 #875의 규약에 따라 정밀 구현을 사용했다. 반대로 명시적 T.__exp 등 빠른 연산자는 HIP에서 전혀 내려가지 않고 tilelang.rocm.language에서 곧바로 AttributeError를 던졌다. 이번 커밋은 HIP과 CUDA가 동일한 전략을 사용하도록 한다. 일곱 개 기본 연산자는 HIPMath로 전환하고, 밑줄이 붙은 여덟 개 빠른 연산자는 HIPFastMath를 가리키는 최저 저하 규칙을 새로 추가하며, 모두 FAST_MATH_INTRINSICS 목록으로 통합해 ROCm 언어 계층에서 내보낸다. 정밀도 비교는 #3618의 재현을 사용한다. MI300X에서 exp와 exp10은 상대 오차, sin과 cos는 절대 오차(x는 1e4까지)를 보며, 네 열은 순서대로 T.exp, T.exp10, T.sin, T.cos이고, 메인 브랜치는 3.7e-06 / 4.1e-06 / 7.8e-04 / 7.8e-04, 이번 커밋은 8.3e-08 / 8.1e-08 / 6.6e-08 / 6.4e-08이다.

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

1.2 메인 저장소: 자동 타깃이 현재 HIP 디바이스를 읽도록 변경, 멀티 GPU 머신에서 더 이상 0번 카드 아키텍처로 고정되지 않음(#3594)

10-27 22:37에 병합되었으며, 2개 파일 +24/-1. determine_target(“auto”)는 ROCm에서 등록된 HIP 프로브로 귀결되는데, 해당 프로브는 torch.cuda.get_device_properties(0)를 무조건 읽으므로 현재 HIP 디바이스가 0번 카드가 아닐 때 자동 프로브는 0번 카드의 아키텍처를 반환한다. CUDA 측은 이미 torch.cuda.current_device()를 사용하고 있었으나 ROCm은 따라가지 못했다. 수정 방식은 먼저 현재 디바이스 인덱스를 가져온 뒤 속성을 조회하고, normalize_rocm_arch로 기능 접미사를 제거하는 것이다. 새 테스트는 current_device()를 스텁으로 1을 반환하게 하고, 프로브된 모든 인덱스를 기록하며, gfx942가 프로브되었고 1번 카드만 건드렸음을 단언한다. 수정 전에는 이 단언이 probed == [0]을 얻었다. 작성자는 본인 기기에 ROCm 디바이스가 없으며, 테스트는 스텁 torch로만 프로브 헬퍼를 검증했고 실제 HIP 런타임은 실행하지 않았다고 명시했다. 같은 날의 다른 변경 #3549와는 다른 함수를 다룬다(#3549는 tilelang/contrib/rocm.py, 이번 변경은 tilelang/rocm/target.py).

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

1.3 메인 저장소: TMA gather/scatter의 상수 행 인덱스를 int32로 변환, 리터럴 표기가 더 이상 컴파일 크래시를 일으키지 않음(#3365)

10-27 09:15에 병합되었으며, 2개 파일 +44/-2, #2601 수정. tma_gather4와 tma_scatter4 두 프런트엔드 헬퍼는 Python 정수 형태의 행 인덱스를 받으면 그대로 전달하는데, TIR에 도달했을 때 정수형 폭이 확정되지 않아 Blackwell에서 컴파일 크래시가 발생했다. 다른 행 표현식은 그대로 유지된다. 이번 변경은 Python 정수를 int32 TIR 상수로 변환하고, Blackwell 게이트가 걸린 컴파일 회귀 테스트를 추가하여 리터럴 행 인덱스로 gather와 scatter를 각각 컴파일한 뒤 생성된 CUDA 소스에 해당 명령이 실제로 등장하는지 확인한다. 작성자는 이 테스트가 컴파일 가능 여부만 검증하며 런타임 수치 정확성은 검증하지 않는다고 밝혔다. 또한 제출자의 로컬 Windows Python 3.13 환경에서는 Home Assistant의 pytest 플러그인이 fcntl에 의존하여 pytest 수집 단계에서 실패했으므로, 이번 변경의 TileLang 테스트는 로컬에서 실제로 실행되지 않았다.

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

1.4 메인 저장소: NVRTC와 CuTeDSL이 첫 입력 텐서의 디바이스에 따라 현재 스트림을 가져옴(#3359)

10-27 08:11에 병합되었으며, 2개 파일 +5/-4, #3301의 F149와 #3408에 이어짐. 두 어댑터는 이전에 문자열 접두사로 대상 유형을 판단했고, 기본적으로 호출자의 현재 디바이스의 현재 스트림을 가져왔다. 따라서 호출자의 현재 GPU와 텐서가 위치한 GPU가 일치하지 않으면 스트림을 잘못 선택했다. 이번 변경은 대상 종류 판단으로 바꾸고, 첫 입력 텐서의 디바이스에 따라 현재 스트림을 조회하며, 명시적으로 전달된 스트림 핸들은 여전히 최우선이고, 텐서가 없는 NVRTC 호출은 기존 의미를 유지한다. 이는 어제의 out_idx 출력 할당 스트림(#3649)에 이은 스트림 의미 정리의 두 번째 변경이다. 작성자는 경계를 명확히 했다. 이번 변경은 스트림 선택만 수정하며, 크로스 디바이스의 기본 스트림 시작 컨텍스트는 여전히 미해결 항목으로, 이러한 호출자는 시작 전에 텐서가 위치한 디바이스를 현재 디바이스로 설정해야 하며, 패치에는 호출별 디바이스 가드가 없다. 검증 시 실제 Torch 스트림과 GPU 시작으로 동일 디바이스 및 크로스 디바이스의 비기본 스트림을 커버했으나, 크로스 디바이스 기본 스트림 컨텍스트의 정확성은 확립하지 않았다.

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

1.5 메인 저장소: 자동 튜닝 캐시 키에 compile_flags 포함, 동일 커널에서 컴파일 옵션을 바꿔도 이전 바이너리를 재사용하지 않음(#3338)

10-27 16:57에 병합되었으며, 3개 파일 +28/-1, #3301의 F131에 이어짐. 자동 튜너는 컴파일 결과를 영속 캐시했으나, 캐시 식별자에 compile_flags가 없어 동일 커널에서 디바이스 컴파일러 옵션을 바꿔도 이전 바이너리가 여전히 적중되어 측정된 성능과 실제 컴파일 결과가 맞지 않았다. 이번 변경은 CompileArgs가 실제로 저장하고 기존 순서대로 해시에 참여하게 하며, flag를 JIT 구현에서 AutoTuner.set_compile_args()를 거쳐 후보 컴파일, 구성별 pass 병합, 그룹 디바이스 컴파일, 캐시 커널 로드까지 전달한다. 검증은 세 가지 기존 캐시 키 테스트와 네이티브 런타임을 사용한 CUDA 정밀 비교를 커버한다. 리스트와 문자열 두 가지 flag 형식 모두 구성별 오버라이드에서 살아남는다. 작성자는 기존 병합 로직이 컴파일러 매크로를 잃어버렸으며, 바로 이것이 이번에 막으려는 구멍이라고 지적했다.

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

1.6 메인 저장소: ThreadSync가 address_of를 쓰기로 기록, 불투명 호출 이후의 펜스를 보충(#3346)

10-27 05:49에 병합되었으며, 2개 파일 +62/-1, #3301의 F305에 이어 어제의 펜스 의미론 주제를 계속한다. address_of 자체는 접근 정보를 지니지 않으며 포인터를 하위로 넘길 뿐이지만, 접근 분류기가 그 타입을 읽기로 하드코딩해 두었기 때문에 공유 버퍼가 불투명 호출을 통해 쓰여진 후 writes 집합에 들어가지 않고, 이후의 전체 블록 읽기가 충돌하는 쓰기를 찾지 못해 펜스를 계획하지 않으며, 스레드는 이렇게 쓰기를 앞질러 쓰기 전 내용을 읽게 된다. 실제 측정된 CUDA 출력은 매우 명확하다: T.address_of(B[0]) 형태는 extern 이후에 펜스가 없고, T.access_ptr(B, “w”) 형태는 있다. 후자는 접근 타입을 명시적으로 선언하므로 이 구멍이 지금까지 발견되지 않았다. loop_unswitching.cc는 이미 address_of를 쓰기로 보수적으로 처리하고 있으며, ThreadSync가 유일한 예외다. 이번 커밋은 address_of를 쓰기로 기록하며, 방향은 보수적이다: 포인터 용도가 알려지지 않았다면 쓰일 수 있고, 펜스 계획은 안전한 쪽에서 한 번 틀리면 된다. 두 개의 테스트 케이스를 추가하여 해당 커널이 하강된 후 extern 이후에 __syncthreads()가 따라오는지 단언한다.

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

1.7 메인 저장소: 런타임 트립 카운트의 파이프라인 epilogue 직접 전개, 페이지드 디코딩 컴파일 복구 (#3356)

10-27 07:02에 병합되었으며, 1개 파일 +5/-1. T.Pipelined 루프의 트립 카운트가 런타임 값일 때, epilogue 전개가 두 단계 이상이고 비동기 복사가 있으면 CUDA codegen이 TypeError: Downcast from tirx.Sub to ir.IntImm failed를 보고한다. 근본 원인은 EmitImpl에 있다: prologue와 epilogue의 extent가 상수일 때만 전개하여 각 반복의 대기 카운트가 상수가 되도록 보장하는데, epilogue의 extent는 end - start, 즉 (n + max_stage) - n이며, n이 상수일 때는 생성 시점에 폴딩되고, n이 런타임일 때는 Sub로 유지되어 as_const_int가 빈 값을 반환하므로 epilogue가 T.unroll 루프로 발행되고, 그 대기가 T.ptx_wait_group(1 - k)처럼 루프 변수를 읽는 형태로 작성되지만 ptx_wait_group은 즉시값만 받는다. 수정법은 판단 전에 먼저 단순화된 extent를 보는 것으로, 상수 경계의 경우는 자기 자신으로 단순화되어 기존과 동일하다. 이 결함은 하위 TileOPs#1852에서 보고되어 afcebed1 이후로 이분되었으며, 피해 대상은 페이지드 디코딩으로, 그 트립 카운트는 KV 길이에서 온다. 작성자는 sm_89, CUDA 13.0에서 두 버전 컴파일을 비교했으며, 추가된 테스트는 메인 브랜치에서 실패하고 이번 커밋에서 통과한다.

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

1.8 메인 저장소: 슬라이스 채우기와 음수 인덱스의 정적 경계 두 건 수정, 범위 초과 쓰기와 무음 음수 주소 차단 (#3352 / #3347)

10-27 05:37과 05:14에 병합되었으며, 합계 4개 파일 +100/-3. 전자는 Fill 생성자의 정적 경계 검사에 대한 것으로: 원래 차원별로 min >= 0과 extent <= shape만 검증했는데, extent <= shape는 버퍼 시작점에서 시작하는 영역에만 성립한다. 시작점이 안쪽에 있는 슬라이스는 두 항목 모두 통과하지만 범위 초과 쓰기로 하강된다 — T.fill(As[6:12], 7)은 길이 8인 버퍼에서 컴파일을 통과하여 As[threadIdx.x + 6] = 7을 발행하고, 6부터 11까지 쓴다. 수정법은 영역 끝도 대상 형상에 대조하여 검증하는 것이며, 새 조건은 시작점과 길이가 모두 정적일 것을 요구하므로 0에 고정된 영역과 심볼릭 시작점 영역은 영향을 받지 않는다. T.clear는 동일 노드를 공유하므로 같은 결함도 함께 차단된다. 후자는 LegalizeNegativeIndex에 대한 것으로: 원래 음수 인덱스에 버퍼 길이를 한 번 더하면 그만이었고 범위 검사가 없어, 길이 8인 버퍼에서 A[-9]는 -1을, A[-10]은 -2를 얻어 할당 시작점을 벗어나면서도 아무런 진단이 없었으나, NumPy와 PyTorch는 같은 인덱스에 대해 IndexError를 던진다. 수정법은 분석기가 랩어라운드 후에도 여전히 음수임을 증명할 수 있을 때만 거부하고 인덱스와 길이를 보고하는 것이며, 범위 내의 모든 상수와 모든 동적 인덱스는 기존 동작을 유지한다.

출처: https://github.com/tile-ai/tilelang/pull/3352, https://github.com/tile-ai/tilelang/pull/3347

1.9 메인 저장소: T.ws(-1) 영역이 모든 스레드에서 실행되어, 단독 작성이 더 이상 저장을 잃지 않음 (#3354)

10-27 04:52에 병합됨, 2개 파일 +40/-3, #2978에 이어짐. WarpSpecialize 문서는 -1을 「threadIdx.x 바인딩 건너뛰기」로 정의하는데, 즉 해당 영역이 threadIdx 제약을 받지 않는다는 뜻이다. 그러나 구현에서는 -1을 일반 warp 그룹 인덱스로 취급하여, 어떤 스레드에도 성립하지 않는 [-warp_group_size, 0) 구간을 생성한다. 그 결과 해당 영역은 0개의 스레드에서 실행되고, 구간 내의 모든 저장이 진단 없이 폐기된다. 최소 재현은 명확하다. 256개 스레드에서 먼저 모두 0을 쓴 뒤 T.ws(-1) 안에서 5를 쓰면, 결과적으로 256개 스레드가 쓴 값은 여전히 0이다. 동일한 저장을 ws 래핑 없이 수행하면 256개 스레드 모두 5를 쓴다. 수정 방법은 id에 -1이 포함된 경우 어떤 범위 제약도 발생시키지 않고, 실제 그룹 인덱스는 자체 스레드 제약을 따르며, 문서 문자열에 -1 예시를 추가하는 것이다. 원래 문서의 예시에는 정확히 이것이 빠져 있었다. 작성자는 전체 저장소를 검토한 결과 ws(-1) 사용례가 나타나지 않았으므로, 기존 코드가 그 빈 영역에 의존하지 않는다고 밝혔다.

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

1.10 메인 저장소: profiler가 장치를 지정하지 않으면 현재 장치 사용, 캐시 플러시가 더 이상 0번 카드에 떨어지지 않음 (#3357)

10-27 07:26에 병합됨, 1개 파일 +1/-1. do_bench에서 장치를 명시적으로 지정하지 않았고 현재 CUDA 장치가 0이 아닐 때, _cache_device(None)이 모듈 수준 기본값인 cuda:0으로 귀결되어, 캐시 플러시를 위한 할당이 실제 측정 작업과 다른 GPU에 떨어질 수 있고, 측정된 수치가 실제 부하와 같은 카드에 있지 않게 된다. 수정 방법은 기본 문자열을 인덱스 없는 “cuda”로 변경하여 PyTorch가 버퍼 할당 시 현재 장치를 해석하도록 하는 것이다. 명시적 장치 지정과 NPU, MPS의 기본값은 변경되지 않는다. 상위 CI(run 37915443170)는 Quick Lint, CUDA, ROCm, Metal, Ascend 및 CuTeDSL 예시에서 통과했다. 작성자는 후속 PR 트리거 실행은 유지관리자 승인이 필요하며 아직 테스트가 실행되지 않았고, 로컬에 CUDA 하드웨어나 타이밍 벤치마크도 없다고 명시했다.

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

1.11 신규 결함: 커널 캐시 적중이 콜드 컴파일을 가정한 6개 테스트를 깨뜨림; CuTeDSL의 CUDA graph 캡처가 빈 그래프 (#3694 / #3693)

두 이슈 모두 lucifer1004가 10-27에 개설했다. 전자는 6개 테스트가 콜드 캐시에서는 통과하지만 동일한 캐시 디렉터리로 다시 실행하면 실패한다는 점을 지적한다. 원인은 이 테스트들이 컴파일이 반드시 발생한다고 가정하기 때문이다. 예를 들어 test_tilelang_kernel_gemm_sm75.py의 sm75 경계 케이스는 네 개의 파라미터에서 AttributeError: ‘NoneType’ object …를 보고한다. 후자는 더 심각하다. CuTeDSL 백엔드에서 발행된 커널 런치가 CUDA graph에 기록되지 않아 torch.cuda.CUDAGraph 캡처가 빈 그래프를 얻고, replay()가 아무것도 실행하지 않는다. 반면 동일한 커널을 cuda 백엔드로 컴파일하면 캡처 재생이 정상이다. 보고서에는 최소 재현이 첨부되어 있다. 두 가지 모두 「캐시와 실행 컨텍스트」 결함군의 연장선으로, 1.4와 1.5 두 건의 병합과 같은 영역을 가리킨다.

출처: https://github.com/tile-ai/tilelang/issues/3694, https://github.com/tile-ai/tilelang/issues/3693

2. 멀티 백엔드 적응 (Ascend / Hygon / ROCm / Hexagon)

2.1 Hygon: HCU scale 버퍼를 듀얼 bank 파티션으로 분할, blockscaled GEMM이 수동 warp 특화 하에서 하나의 행 매핑을 공유 (#29)

10-27 08:26에 병합됨, 12개 파일 +1389/-314, 이번 모니터링 기간 소프트웨어 측 최대 규모다. 변경은 scale 버퍼의 행을 독립적인 듀얼 bank 파티션으로 계획하여 LDS bank와 분리하고, 동일한 행 매핑을 할당, scale 복사, MMAC 하강 세 곳에 동시에 적용하여 세 곳이 각자 레이아웃을 추론하는 것을 방지한다. 세부 사항으로는 물리적 구간화 수명 주기와 m0 정렬 페이징 유지, scale 복사 발행 횟수 압축, 유휴 wave의 작업을 두 scale bank에 분산, 이 레이아웃을 수동 warp 특화 하의 blockscaled GEMM에 사용하는 것이 포함된다. 이번 변경은 어제 병합된 네이티브 B16 저장과 동적 WASP GEMM(#28)에 이어지며, Hygon은 이틀 동안 세 건의 메인라인 병합을 완료했다.

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

2.2 Ascend 메인 저장소: PerCoreTask 후보체가 내부 if/else 허용, 조건부 L1 채우기가 더 이상 전체 거부되지 않음 (#3686)

10-27 07:24에 병합, 2개 파일 +85/-24. T.PerCoreTask 내의 T.Task는 본문에 if/else가 포함되면 거부되었습니다. 비록 분기가 내부 작업만 제어하더라도—예를 들어 한 번 복사한 후 조건에 따라 L1 채우기를 보충하는 경우—마찬가지였습니다. d85a70a9에서 재현되었으며, 구체화 단계에서 「선택 후보의 제어 흐름은 task 스코프 밖에 있어야 합니다」라는 오류가 발생했습니다. 이번 커밋은 이 일괄 거부를 제거하고 task 본문에 대한 통합 검증으로 변경했습니다. 동기 배리어가 전체 후보를 감싸고, 각 분기는 여전히 파이프라인과 코어의 호환성, 중첩 task, 그리고 지원되지 않는 크로스 코어 동기화를 항목별로 검사합니다. 회귀 테스트는 내부 if와 if/else를 수용하고, 후보 본문을 유지하며, 기존의 거부 케이스도 보존합니다. 또한 완료 플래그가 외부 가드 아래에서 전체 후보를 따라가는지 확인합니다. task, 스케줄링 유닛, 지연 세 가지 세트 총 113개 케이스가 모두 통과했으며, 조건부 L1 채우기가 있는 커널도 완전한 Ascend 소스 하강을 완료할 수 있습니다. 디바이스 실행은 테스트되지 않았습니다.

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

2.3 Ascend 메인 저장소: CI가 AscendC와 PTO를 두 개의 매트릭스 작업으로 분리, 개별 재실행 가능 (#3687)

10-27 07:05에 병합, 7개 파일 +56/-27. 이전에는 AscendC와 PTO가 동일한 Test for Ascend 작업의 두 단계였기 때문에, 어느 한 백엔드에 특화된 실패가 발생하면 전체 검사가 빨간색으로 표시되었고, 재실행 시 반드시 다른 한쪽 예제도 함께 실행되었습니다. 이번 커밋은 예제와 단위 테스트를 두 개의 매트릭스 작업으로 분리했습니다. Test for AscendC는 examples/ascend(target은 ascend, 동시성 4)와 testing/ascend -m “not pto”(동시성 32)를 실행하고, Test for PTO는 examples/ascend(target은 pto)와 testing/ascend -m pto(동시성 32)를 실행합니다. 구분은 기존의 pto 마커를 활용하며, PTO만 실행하는 테스트와 이중 백엔드 테스트 내의 PTO 파라미터에 마커를 추가했습니다. 두 마커 표현식은 상호 배타적이며 합치면 전체 세트를 커버하므로 중복 실행이 없습니다. fail-fast는 false로 설정하여 한 작업이 실패해도 다른 작업이 취소되지 않습니다. 양쪽 모두 「예제 실패 시에도 빌드가 성공하면 단위 테스트를 계속 실행」하는 동작을 유지하며, 각자 빌드 및 테스트 요약을 게시합니다. 공용 테스트(동일한 프론트엔드 IR을 두 코드 생성기에 각각 통과시키는 계약 검사 포함)는 AscendC 측에서만 한 번 실행하며, 이 합의는 testing/ascend/README.md에 기록되었습니다.

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

2.4 Ascend 저장소 리뷰 중: Cython 래퍼의 지연 생성 분석기, 어제 기록된 호출마다 재생성 문제를 직접 겨냥 (#1948)

10-27 07:42 업데이트, 작성자 zwh1025. 어제 데일리 리포트에 기록된 새로운 결함 #1945는 Ascend 저장소의 Cython 래퍼가 호출마다 새로운 TVM 분석기를 생성하여 실측 약 0.28밀리초의 고정 오버헤드가 발생한다고 지적했습니다. 이번 커밋은 「TVM 분석기 지연 생성」을 제목으로 이 오버헤드를 정면으로 다룹니다. 이러한 소규모 오버헤드는 대규모 모델 추론의 레이어별 호출에서 관측 가능한 지연으로 누적될 수 있으며, Ascend 측 성능 청소의 손쉬운 성과에 해당합니다. 현재 리뷰 중입니다.

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

2.5 Ascend 저장소 리뷰 중: GLM-5.3의 kpool 압축 연산자, PTO상의 텐서 시프트량 (#1910 / #1949)

두 건 모두 10-27에 업데이트되었습니다. #1910은 Ascend NPU에 GLM-5.3의 kpool 압축 연산자 예제를 추가합니다—이는 특정 모델의 구조 요소를 예제 저장소로 가져오는 방식으로, TileOPs 측의 「연산자는 모델 요구에 따라 성장한다」는 사고와 일치합니다. #1949는 T.tile.bitwise_lshift가 PTO에서 텐서 형태의 시프트량을 지원하도록 하며, 이전에는 상수만 허용했습니다. PTO는 Ascend 측에서 상대적으로 독립적인 코드 생성 경로이며, 시프트량의 텐서화는 해당 경로가 마스크 및 비트 패킹류 연산자에서 범용 표현에 한 걸음 더 다가갔음을 의미합니다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1910, https://github.com/tile-ai/tilelang-ascend/pull/1949

2.6 Ascend 메인 저장소 리뷰 중: PTO의 T.Parallel 벡터화, AutoSimVF 채널, 배치 3D ND2NZ, 쓰기 카운트 다중 버퍼링, Bisheng 빌드 아이덴티티를 캐시 키에 포함 (#3512 / #3320 / #3689 / #3629 / #3691)

다섯 건의 리뷰 중 항목이 이번 모니터링 기간 Ascend 측의 대기 큐를 구성하며, 모두 10-27에 업데이트되었습니다. #3512는 PTO에서 T.Parallel을 벡터화하는 파이프라인 프레임워크와 요소별 지원으로, 구조적 변경에 해당합니다. #3320은 AutoSimtVF 채널을 신규 추가하고 Ascend 파이프라인에 등록하며, 방향은 SIMT와 벡터 사이에서 자동으로 경로를 선택하는 것입니다. #3689는 배치 3D GM에서 L1으로의 ND2NZ 복사를 지원하며, 배치 어텐션과 블록 GEMM의 전송 요구에 직접 대응합니다. #3629는 쓰기 카운트 방식 다중 버퍼링을 추가하고 DeepGEMM의 로딩을 재배치합니다. #3691은 Bisheng 컴파일러 빌드 아이덴티티를 바이너리 캐시 키에 포함하며, 1.5의 compile_flags를 캐시 키에 포함하는 것과 같은 종류의 사고입니다—Ascend 측에도 컴파일러 교체 후 이전 바이너리에 잘못 히트하는 문제가 동일하게 존재합니다.

출처: https://github.com/tile-ai/tilelang/pull/3512, https://github.com/tile-ai/tilelang/pull/3320, https://github.com/tile-ai/tilelang/pull/3689, https://github.com/tile-ai/tilelang/pull/3629, https://github.com/tile-ai/tilelang/pull/3691

2.7 ROCm 심사 중: pip 버전 ROCm 10의 아키텍처 탐지와 이미지, FP4 변환 및 E8M0 타입 (#3549 / #3665)

#3549는 pip로 설치하는 ROCm 10 SDK를 지원하며, 변경 사항에는 아키텍처 탐지와 Dockerfile이 포함된다. 이는 1.2의 auto 타깃 수정과 관련이 있지만 서로 다른 함수에 속한다. 이번 변경은 /opt/rocm 기반의 get_rocm_arch 탐지를 수정하고, 1.2는 tilelang/rocm/target.py의 탐지 헬퍼를 수정한다. #3665는 HIP 템플릿의 부동소수점에서 FP4로의 변환을 수정하고 E8M0 타입 지원을 보강하며, 이는 블록 스케일링과 저정밀도 양자화의 ROCm 가용성을 겨냥한다. 두 항목 모두 심사 중이다.

출처: https://github.com/tile-ai/tilelang/pull/3549, https://github.com/tile-ai/tilelang/pull/3665

2.8 Ascend 저장소 일일 테스트: 2903개 항목 전부 통과, 이틀 연속 실패 없음 (#1950)

Ascend 저장소의 일일 테스트 봇이 10-28 06:49에 보고한 결과에 따르면 총 2903개 항목, 통과 2903개, 통과율 100%이며 워크플로 실행 번호는 40590261193이다. 전날 보고서(#1947)도 마찬가지로 2903개 항목 전부 통과였다. 여전히 빠르게 진화하는 NPU 백엔드에게 이틀 연속 실패 없는 일일 테스트는 메인 브랜치의 가용성이 안정적임을 보여주며, 이는 2.3의 CI 분할에 기준선을 제공한다. 분할 후에도 두 작업 각각이 이 숫자를 유지해야 한다.

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

2.9 MetaX: 조직 뷰에 푸시 타임스탬프 존재, 여섯 개 브랜치 모두 모니터링 기간 내 신규 커밋 없음 (10-27)

tilelang-metax 저장소의 pushed_at은 10-26 07:08로, 모니터링 기간 시작점 이전에 해당한다. 여섯 개 브랜치를 브랜치별로 점검한 결과 모니터링 기간 내 신규 커밋이 없어 실질적 변경 없음으로 처리하며, 항목 수준의 내용은 나열하지 않는다.

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

2.10 서드파티: Hexagon DSP 백엔드 하루 15건, HVX 벡터화, VTCM 아레나, Crouton 레이아웃 및 DMA 내장 기능 완비 (10-27)

xwhzz/tilelang-hexagon의 hexagon-backend 브랜치는 10-27 하루에 15건의 커밋을 기록했으며, 이는 모니터링 기간 내 백엔드 측에서 가장 밀도 높은 외부 활동이다. 당일 주된 흐름은 세 단계로 나눌 수 있다. 저층에서는 CodeGenTileLangCBase로부터 코드 생성기를 재구축하고, DMA를 커밋 그룹을 갖춘 IR 내장 기능으로 만들고, 시작 그리드와 워커 스레드 풀을 TIR로 내리고, TIR에서 VTCM 아레나를 계획했다. 중간층에서는 HVX의 단일 벡터화 경로(산술 매핑, 벡터 수학 및 원소별 경로 하나)를 구축하고, HVX 전치 트리로 다중 행 리덕션을 수행하고, 소프트웨어 파이프라이닝된 HVX 매핑을 적용하고, 프래그먼트와 T.reduce_*를 지원했다. 상위 계층에서는 Crouton 레이아웃과 영역 복사를 도입하고, HVX 매핑과 리덕션이 Crouton 타일 위에서 직접 실행되게 하고, 세션마다 DSP 주파수에 투표하고, 마지막으로 빌드 캐시와 병렬 디바이스 테스트를 보강했다. 커밋에는 파괴적 변경 표시(feat(hexagon)!)가 포함되어 있으며, 이는 Crouton 레이아웃과 영역 복사가 기존 인터페이스를 변경했음을 의미한다. Hexagon은 휴대폰과 임베디드 측의 DSP로, 이 라인이 성공한다면 TileLang의 하드웨어 커버리지는 GPU, NPU에서 DSP로 확장된다.

출처: https://github.com/xwhzz/tilelang-hexagon

3. 생태계와 채택측 (TileOPs / TileFoundry)

3.1 TileOPs 0.1.0 릴리스: 206개 연산자의 첫 정식 버전, 공개 API 확정 (10-27)

10-27 13:15에 v0.1.0을 공개했으며, 이는 저장소의 첫 GitHub release이다. PyPI의 0.1.0 아티팩트는 더 이른 10-27 07:45(베이징 시간, 즉 10-26 23:45 UTC)에 업로드되었는데, 이는 직전 데일리 리포트의 마감 시점보다 45분 늦은 시점이다——이는 어제 데일리 리포트에서 「첫 릴리스가 여전히 마지막 칸에 있다」고 쓴 직후에 바로 일어난 변화이다. 릴리스 노트의 포지셔닝은 매우 명확하다. agent 시대를 위해 다시 만든 연산자 라이브러리로, 각 연산자는 먼저 YAML 명세(시그니처 선언, 테스트 및 벤치마크용 워크로드, roofline 공식)를 작성하고, 코딩 agent가 커널을 개발하고 튜닝하며, 명세에서 파생된 테스트와 벤치마크가 어떤 것이 라이브러리에 들어갈 수 있는지를 결정한다. foundry 태그가 붙은 195건의 병합이 이런 방식으로 들어왔다. 이번 버전은 공개 API를 고정한다. 연산자는 tileops..로 임포트하며, docs/stability.md가 공개 모듈을 나열하고, scripts/ci/check_api_compat.py가 이전 버전 작성 방식을 깨뜨리는 모든 PR을 즉시 실패하게 만든다. 이후 0.x에서는 추가만 하고 변경하지 않는다. 몇 가지 기능도 함께 버전에 들어간다. torch.compile(fullgraph=True)이 모든 manifest 항목 중 호출 시점 텐서 입력을 포함하고 조합이 없는 연산자에 대해 사용 가능하며, H200에서 그래프 테스트 그룹 내 연산자의 CUDA Graph 캡처를 검증했고, op.request_tune() 후 자동 튜닝되어 컴파일 산출물과 튜닝 결과가 TILELANG_CACHE_DIR 아래에 캐시되어 프로세스 간 재사용되며, tileops.backend는 백엔드 배포판이 register_kernel_builder를 통해 자신의 디바이스상의 전체 연산자를 연결할 수 있게 하고, tileops.trace는 커널 내 타임라인을 기록한다. 릴리스 감사에서는 @FeathBow 한 사람이 98건의 병합을 기여했다고 밝혔다.

출처: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0, https://pypi.org/project/tileops/

3.2 TileOPs 0.1.0 벤치마크: H200에서 2265행 중 2262행이 최속 구현보다 느리지 않고 기하평균 1.37배 (10-27)

야간 벤치마크는 각 TileOPs 커널을 동일 연산자의 다른 모든 구현과 같은 GPU에서 비교하며, 단일 행의 비율은 가장 빠른 비교 대상의 소요 시간을 TileOPs 소요 시간으로 나눈 값으로, 1보다 크면 더 빠름을 의미한다. 전체 2265행, 16개 연산자 패밀리에서 기하평균 1.37배이며, 2262행이 0.99 이상을 달성했다. 패밀리별로 차이가 뚜렷하다. sequence_modeling 4.00배, sampling 2.95배, convolution 1.86배, mamba 1.61배, linear_attention 1.54배, pool 1.51배, quantization 1.45배, attention 1.42배이며, 규모가 가장 큰 elementwise는 527행에 1.09배, gemm은 235행에 1.28배이다. 비교 대상에는 torch, torch.compile, FlashAttention 2와 3, FlashInfer, FlashMLA, vLLM, DeepGEMM, flash-linear-attention, FlagGems, QuACK, Mamba, Marlin, cuBLASLt가 포함된다. 0.99에 뒤진 세 행은 각각 16M float32 요소에서의 Where와 ClampScalar, 그리고 DeepSeek-V3 프리필 출력 프로젝션에서의 FP8 GEMM으로, 격차는 1.1%에서 1.7%이다. 이 분포는 이득이 주로 롱테일 연산자 패밀리에서 나오며 주류 GEMM과 원소별 연산은 아니라는 것을 보여준다——후자는 원래 각 구현이 가장 치열하게 경쟁하는 영역이다.

출처: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.3 TileOPs 0.1.0의 알려진 격차와 다음 단계: 융합 연산자, 전체 모델, 연산-통신 융합 (10-27)

릴리스 노트는 네 가지 알려진 문제를 나열했다. FusedAddLayerNormFwdOp가 0에서 멀리 떨어진 행에서 정밀도를 잃고 1e36에 가까운 행에서 오버플로한다(#2773, 어제 데일리 리포트에서 이미 이 결함을 기록했다). 동일 호스트에 서로 다른 아키텍처의 GPU가 있으면 한 GPU용으로 컴파일된 커널이 다른 GPU에서 실행될 수 있다(#2302). Bwd 계열 연산자는 직접 호출해야 하며 torch.autograd에 연결되지 않는다. 단일 연산자 인스턴스는 동시에 하나의 스레드와 하나의 CUDA 스트림만 지원한다. 검증 범위는 H200, Python 3.12와 3.13만을 포함하며, 일부 커널은 별도로 SM80, SM86, SM89 경로가 있어 해당 디바이스에서 실행되지만 릴리스 검증 범위에는 속하지 않는다. 로드맵은 세 단계이다. 먼저 연산자 경계를 넘는 복잡한 융합(어텐션과 전후 프로젝션 및 정규화, 전체 레이어 MoE, GEMM epilogue 내 양자화)을 하고, 다음으로 전체 모델의 프리필과 디코딩을 서빙 엔진과 엔드투엔드로 비교하며, 마지막으로 연산과 통신의 융합(텐서 병렬 GEMM과 all-reduce, 전문가 병렬 MoE와 all-to-all)을 한다. 이 세 단계가 실현되면 TileOPs의 경쟁 상대는 단일 연산자 라이브러리에서 추론 프레임워크로 바뀐다.

출처: https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.4 TileOPs 4건 병합: where와 masked_fill 튜닝 가능, D128 인과 대각 64 키 기준 채점, 희소 장 GDN 프리필 분할, 변경 로그 릴리스 날짜 기준 확정 (#2782 / #2783 / #2779 / #2781)

4건이 10-27에 병합되었으며, 이는 릴리스 직전의 마무리 작업이다. #2782는 request_tune()이 WhereFwdOp, MaskedFillTensorFwdOp, MaskedFillScalarFwdOp에서 ValueError를 발생시키는 문제를 수정했다. 자동 튜닝이 무작위로 생성된 정수 입력을 거부하는데, 이 세 커널이 bool 조건이나 마스크를 uint8로 전달하기 때문이다. 무작위 선택기는 각 요소가 어떤 피연산자를 취할지만 변경하고 후보의 작업량은 바꾸지 않으므로 무작위 정수 입력을 안전하다고 선언했다. 테스트 노드 3개를 새로 추가했으며, 메인 브랜치에서 세 사례 모두 실패했다. #2783은 D128 밀집 프로그램의 인과 대각을 최적화했다. 인과 호출에 key 구간이 하나뿐이고 두 길이가 모두 128행 경계에 있을 때, 첫 번째 소비 warpgroup의 대각 key 타일이 모든 행에 대해 상위 64 키를 마스킹하므로, 이 단계를 m64n64의 QK에 64 키 softmax를 조합하여 실행하도록 변경했다. 이 선택은 컴파일 타임 플래그(diag_aligned, 커널 키에 포함)인데, 런타임에 64와 128 너비 QK 사이를 분기하면 ptxas가 일부 특수화에서 모든 WGMMA를 직렬화하기 때문이다(C7511). 플래그로 전환한 후 manifest의 모든 D128 프리필 특수화에서 더 이상 C75xx가 발생하지 않는다. #2779는 SM89에서 희소 장 시퀀스의 GDN 프리필을 상태 스캔과 독립 출력 프로그램으로 분할했으며, 밀집 key-major D128 입력이면서 최소 16384 토큰, 상태 스캔 CTA가 SM 수의 4분의 1을 넘지 않는 경우에 적용된다. 벤치마크에서 (1, 16384, 4, 128) bf16의 소요 시간이 약 0.612밀리초에서 약 0.470밀리초로 감소했으며, flash-linear-attention의 0.441밀리초 대비 0.939배이다. #2781은 CHANGELOG에서 0.1.0의 날짜를 PyPI에 등록된 날(2026-10-27)로 확정하고, 기여자 문서에서 릴리스 1회 설정과 Actions 클릭 경로를 삭제했다.

출처: https://github.com/tile-ai/TileOPs/pull/2782, https://github.com/tile-ai/TileOPs/pull/2783, https://github.com/tile-ai/TileOPs/pull/2779, https://github.com/tile-ai/TileOPs/pull/2781

3.5 TileFoundry 2건 병합: 전체 윈도를 피연산자 자체로 하향 전달, 동적 공유 메모리를 커널 실제 사용량 기준으로 크기 결정 (#273 / #271)

두 건 모두 10-27에 병합되었다. #273은 sm_89 GEMM의 하향 전달을 수정했다. schedule finalize가 작성한 내용이 컴파일되지 않는데, 하향 전달이 전체 누산기를 0에서 시작하고 범위가 자기 자신과 같은 윈도 T.ptr_of(acc[0:0+16, 0:0+8])로 감싸기 때문이며, CUDA codegen이 분할 레지스터 타일의 포인터 슬라이스를 거부하여 “일반 소스 레이아웃 요구” 오류를 낸다. 또한 MMA가 레지스터 프래그먼트를 lhs_frag_dst[0:0+16, 0:0+32] 같은 원시 Slice로 읽는데 codegen 역시 처리기가 없다. 수정은 하향 전달 측에서 이루어졌다. StoragePlan.operand는 원래 “0에서 시작하고 피연산자 전체를 걸치는 윈도”를 피연산자 자체로 취급하지만, 구체화와 명령 하향 전달이 여전히 Slice를 구성한다. 이를 한 번만 판정하도록 변경했다. MLIR의 TrivialSubViewOpFolder에 대해 동일한 치환을 수행한다. #271은 동적 공유 메모리를 수정했다. 합성된 호스트 진입점이 각 TIR 커널에 대해 dynamic_smem=0으로 시작하므로, 정수 TensorView 오프셋으로 공유 메모리에 접근하는 커널은 공간을 할당받지 못하며, 48 KiB를 초과하면 cudaFuncAttributeMaxDynamicSharedMemorySize를 설정해야 하는데 런처는 이를 전혀 설정하지 않는다. 작성자는 #267의 “목표 용량만큼 최대한 할당” 제안을 채택하지 않고 커널 실제 사용량 기준으로 크기를 결정하는 방식으로 변경했으며, 그 이유는 정적 공유 메모리와 동적 공유 메모리가 동일한 한도를 공유하고 각 블록의 요청량이 하나의 SM에 몇 개를 배치할 수 있는지를 결정하기 때문이다. sm_89에서 99 KiB를 채우면 512바이트를 사용하는 커널도 하나만 배치할 수 있다. CUTLASS와 TileLang 모두 같은 방식으로 크기를 결정한다.

출처: https://github.com/tile-ai/TileFoundry/pull/273, https://github.com/tile-ai/TileFoundry/pull/271

3.6 외부 생태계: CANN 커뮤니티의 TileLang 연산자 저장소, PPU 버전 무수정 래퍼, DeepSeek V4.1 Flash 프래그먼트 예제 (10-27)

세 개의 외부 저장소 모두 모니터링 기간 내에 움직임이 있었으며, 종합해 보면 ‘Ascend 측 연산자 자산’이 형성되는 과정에 해당한다. Atlas-Ascend-Fork/ops-tilelang은 CANN 커뮤니티가 Ascend를 대상으로 만든 TileLang 연산자 저장소로, 10-20에 저장소를 생성하고 10-27에 “자동 발견 증분 테스트 및 연산자 매핑 테이블 제거”를 커밋했다. README는 Ascend 950 이상 아키텍처만 지원하고 CANN 9.3.0 이상 및 커뮤니티 weekly 버전 사용을 요구한다고 명시하고 있으며, 설계 원칙은 인터페이스와 스케줄링 분리, 비즈니스 도메인별 집계, 성능보다 정확성 우선, 벤치마크와 테스트 분리 및 베이스라인을 테스트 파일과 칩 클러스터 단위로 분할 저장하는 것이다. kevin3113/models-tilelang은 이 저장소의 일곱 개 연산자로 DeepSeek V4.1 Flash의 “Attention 출력 + FFN 전처리” 구간을 엮었으며(mhc_post_forward, 제곱합 GEMM, 리덕션 + RMSNorm, mix 분할, sinkhorn, apply_mix, ffn_norm), HIDDEN=5120, MHC=4, 기본 decode 형태 SEQ_LEN=1이고 SuperKernel과 npugraph_ex의 ACL graph를 지원하며, B를 1부터 256까지 실측했을 때 SuperKernel 이득은 1.17%에서 2.75%였다. flytiger-eco/tilelang-wrapper-for-sail은 다른 방향의 마이그레이션을 해결한다. 즉 CUDA 스타일 TileLang 코드를 무수정으로 tilelang-sail(PPU 백엔드 배포판) 위에서 실행하는 것으로, 방식은 순수 런타임에서 import 프로토콜과 site 메커니즘을 가로채는 것이며 업스트림 소스 코드를 수정하지도 재패키징하지도 않는다. PPU에 대응물이 없는 CUDA 기능은 조용히 저하되는 대신 명시적으로 오류를 낸다. 이 셋은 각각 연산자 공급, 모델 적용, 백엔드 마이그레이션이라는 세 단계에 대응한다.

출처: https://github.com/Atlas-Ascend-Fork/ops-tilelang, https://github.com/kevin3113/models-tilelang, https://github.com/flytiger-eco/tilelang-wrapper-for-sail

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

4.1 문서 사이트 하루 여섯 건 자동 업데이트 (10-27)

tilelang.github.io는 모니터링 기간 내에 여섯 차례 Update docs 커밋(05:20, 06:22, 07:19, 09:28, 22:42, 22:43)이 있었으며, 모두 봇이 푸시한 것으로 내용은 메인 저장소 병합 후 문서 재생성에서 비롯되었다. 배포 당일 문서 사이트의 업데이트 빈도는 메인 브랜치 병합 리듬과 동기화되어 있으며, 이는 정상적인 연동으로 별도 내용으로 다루지 않는다.

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

4.2 Ascend 저장소의 문서 사실 검증 4건, 단일 코멘트 최대 46회 (#1628 / #1655 / #1656 / #1659)

Ascend 저장소는 10-27에 “사실 검증”이라는 제목의 문서 PR 업데이트가 네 건 있었다. T.tile.clear(#1628), T.tile.bitwise_or(#1655), T.tile.bitwise_not(#1656), T.tile.bitwise_lshift(#1659)이며, 코멘트 수는 각각 10, 46, 18, 20이고 그중 세 건에 봇 표시가 있다. 한 차례 검증이 46개 코멘트까지 굴러간다는 것은 논의가 단순한 표현 문제가 아니라는 뜻이다. 이러한 API별 동작 대조는 실질적으로 Ascend 측의 의미 규약을 기록으로 남기는 작업이다. 새로운 백엔드 라인에게 문서와 구현의 정합성 비용은 흔히 과소평가되는데, 이 PR 묶음은 가시적인 투입이다.

출처: https://github.com/tile-ai/tilelang-ascend/pull/1655, https://github.com/tile-ai/tilelang-ascend/pull/1656

4.3 심사 중: lower trace를 독립 개발 툴킷으로 분리 (#3583)

10-27에 업데이트된 이 심사 중 PR은 lower trace를 메인 저장소에서 tilelang-devkit으로 분리한다. 하향 과정의 추적 정보는 백엔드 적응 디버깅에 매우 유용하지만, 메인 저장소에 두면 컴파일 경로와 의존성 부담이 늘어난다. 독립 툴킷으로 분리하면 메인 저장소는 경량을 유지하고 필요할 때 다시 붙이면 된다. 이러한 “디버깅 능력의 외부화” 변경은 보통 프로젝트에 여러 다운스트림 백엔드가 생기기 시작할 때 나타나는데, 이번 모니터링 기간 내에 Hexagon, PPU, Ascend 세 라인이 동시에 움직이고 있는 것이 바로 그런 시점이다.

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

4.4 뉴스 측: Google News와 HN에서 24시간 내 이번 모니터링 기간의 신규 보도 없음 (10-27)

네 가지 쿼리(영문·중문, TileLang, TileOPs, TileLang 연산자, tile-ai 즈싼 kernel)에서 반환된 항목은 모두 9월 말에서 10월 중순에 걸쳐 있습니다. 한 무리는 DeepSeek과 Huawei가 9월 29일~30일에 Ascend 버전 TileLang을 오픈소스로 공개했다는 보도(SCMP, Reuters, Tom’s Hardware, Leiphone, InfoQ 등)이며, 가장 최신 항목은 10월 22일 SemiAnalysis의 Huawei와 DeepSeek 공생 관계 분석으로, 모두 모니터링 기간 밖입니다. TileOPs 정확 쿼리는 0건입니다. Hacker News 모니터링 기간 내에는 GameMaker 재구현 관련 항목 하나뿐으로, 노이즈로 제외했습니다. arXiv에서 TileLang 키워드 최근 논문은 2026-07-24입니다. 오늘 콘텐츠의 중심은 전부 GitHub 쪽에 있으며, 이는 릴리스일 시점과 일치합니다 — 정식 릴리스가 아직 미디어 보도 주기에 진입하지 않았습니다.

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

5. 트렌드 관찰

첫째, TileOPs 0.1.0은 「연산자 라이브러리」라는 일의 비용 구조를 다른 알고리즘으로 바꿨습니다. 전통적 방식은 연산자 수가 하드웨어와 모델 요구에 따라 선형으로 늘어나고 인력이 병목입니다. 여기서의 방식은 연산자마다 먼저 YAML 명세를 한 부 작성해 시그니처, 워크로드, roofline을 명확히 선언하고, 그다음 agent가 커널을 생성·튜닝하며, 마지막으로 명세에서 파생된 테스트와 벤치마크로 진입 심사를 합니다. 206개 연산자, 195건의 foundry 태그 머지는 이 프로세스의 첫 전량 납품입니다. 이것이 지속될 수 있는지는 두 가지에 달려 있습니다: 명세가 복잡한 융합 연산자를 커버할 수 있는지(로드맵 첫 단계가 바로 이것), 그리고 벤치마크가 「벤치마크에 맞춰 튜닝됐지만 실제로는 견고하지 않은」 커널을 막을 수 있는지 — FusedAddLayerNormFwdOp가 0에서 먼 행에서 정밀도를 잃는 이 알려진 문제가 바로 이런 리스크의 한 표본입니다.

둘째, 이 모니터링 기간의 머지는 거의 전부 「조용한 오류」입니다. ROCm의 수학 함수가 기본으로 빠른 내장 함수를 타면서 sin/cos가 세 자릿수만큼 차이 나게 하고, address_of가 읽기로 취급되면서 불투명 호출 이후의 공유 메모리에서 쓰기 전 내용을 읽게 하며, T.ws(-1)의 영역이 스레드 0개에서 실행되어 저장이 전부 버려지고 오류도 나지 않으며, 슬라이스 채우기가 길이가 아닌 끝점으로 검증되어 그대로 범위를 넘어 쓰고, 음수 인덱스가 랩어라운드 후에도 음수로 남아 할당 시작점 이전으로 떨어집니다. 이 다섯 가지의 공통점은 컴파일은 통과하고 실행 시 오류는 없으며 결과가 틀린다는 것이고, 모두 최근 2주 안에 집중적으로 발견됐습니다. 대형 모델용 커널을 작성하는 언어에게 이런 결함의 대가는 컴파일 시점 오류보다 훨씬 크며, 이들은 같은 엔지니어링 요구를 가리킵니다: 「오류는 없지만 동작이 틀린」 표기법을 컴파일 시점에 거부하는 것. 이 모니터링 기간의 경계 검증 두 건 수정과 ws(-1) 수정이 하는 일이 바로 이것입니다.

셋째, 백엔드 커버리지가 「벤더 주도」에서 「커뮤니티 자발」로 옮겨가고 있습니다. 공식 유지보수되는 Ascend, Hygon, MetaX, Moore Threads 네 라인 외에, Hexagon DSP 백엔드가 하루에 15건을 머지하며 선형대수 원시 연산을 Crouton 타일에 깔았고, PPU 버전에는 무변경 래퍼가 등장했으며, CANN 커뮤니티가 자체 연산자 저장소를 세우고 그것으로 DeepSeek V4.1 Flash의 모델 조각을 엮어냈습니다. 이 작업들은 모두 tile-ai 조직의 저장소에 있지 않지만, 모두 TileLang 언어 계층의 안정성에 의존합니다. 이것이야말로 0.1.0이 공개 API를 동결하고 API 호환성 검사를 CI 게이트로 만든 의미이기도 합니다 — 다운스트림이 많을수록 하위 호환의 가치가 높아집니다.

넷째, 캐시 아이덴티티 문제가 두 백엔드에 동시에 포착됐습니다. 메인 저장소는 compile_flags를 자동 튜닝 캐시 키에 포함시켰고(#3338), Ascend는 Bisheng 빌드 아이덴티티를 바이너리 캐시 키에 포함시켰습니다(#3691). 같은 커널이 서로 다른 컴파일 옵션이나 컴파일러에서 다른 바이너리를 생성하는데, 캐시가 소스와 형상만으로 키를 만들면 잘못 적중하여 측정된 성능이 실제 컴파일 결과와 맞지 않습니다. 두 개의 독립 노선이 같은 날 같은 일을 했다는 것은 이런 문제가 실제 사용에서 이미 충분히 고통스럽다는 뜻입니다.

부록. 자료 및 검증 설명

출처 검증표

출처 검증 결과
GitHub tile-ai 조직(29개 저장소) 모니터링 기간 내 푸시된 저장소 7개: tilelang, TileOPs, TileOPs-nightly, TileOPs.github.io, tilelang-hygon, TileFoundry, tilelang.github.io
병합 확정값(is:merged) 20건, 건별로 병합 시간이 모니터링 기간 내에 있음을 대조 확인(tilelang 13, TileOPs 4, TileFoundry 2, tilelang-hygon 1)
심사 중 확정값(is:open 업데이트) 45건, 업데이트 시간순으로 전량 대조 확인
신규 issue 7건, 건별로 본문을 확인했으며 그중 2건은 Ascend 저장소의 일일 테스트 봇 보고서
릴리스 페이지 및 PyPI TileOPs v0.1.0은 10-27 13:15에 게시되었고 PyPI 아티팩트 업로드는 10-27 07:45; tilelang 최신 버전은 여전히 0.1.15(09-30), tilefoundry 최신 버전은 0.0.4rc1(10-25)
tilelang-ascend 저장소 모니터링 기간 내 커밋 없음, 다만 PR과 일일 테스트 issue 활동은 활발
tilelang-metax 저장소 조직 뷰에 푸시 타임스탬프가 있으나 여섯 개 브랜치 모두 모니터링 기간 내 커밋이 없어 실질 변경 없음으로 처리
tilelang-musa 저장소 마지막 푸시 09-30, 모니터링 기간 내 커밋 없음
채택 측(TileKernels / FlashQLA) 모두 모니터링 기간 내 푸시 없음(10-14)
외부 백엔드 저장소 Hexagon 백엔드 15건 커밋; CANN 연산자 저장소, PPU 래퍼, 모델 스니펫 저장소 모두 모니터링 기간 내 작업 있음, 2.10 및 3.6 참조
Google News 중영문 다중 쿼리 정확 쿼리 적중 0건, 반환된 항목은 전부 9월 말부터 10월 중순에 걸쳐 있어 모두 모니터링 기간 밖
Hacker News 모니터링 기간 내 GameMaker 재구현 항목 1건, TileLang과 무관하여 제외
arXiv 키워드 검색 결과 모니터링 기간 내 신규 논문 없음, 최근 논문은 2026-07-24

전체 출처 목록

tilelang Bisheng 빌드 ID 캐시 키 반영(#3691) — https://github.com/tile-ai/tilelang/pull/3691