TileLang 데일리 리포트 (2026-10-30)
모니터링 기간: 2026-10-29 07:00 ~ 2026-10-30 07:00 (베이징 시간, 24시간).
출처: GitHub tile-ai 조직 저장소별 점검(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 문서 사이트 / 야간 미러 저장소) + 병합 및 심사 중 실값 검색(병합 12건, 심사 중 업데이트 54건) + 신규 issue 전수 확인(4건) + 메인 브랜치 커밋별 점검(tilelang 5, tilelang-ascend 5, TileOPs 1, TileFoundry 1, 문서 사이트 5) + 릴리스 페이지 및 tag 점검 + 서드파티 백엔드 및 생태계 저장소 점검(저장소 10개) + Google News 중영문 6개 쿼리(프록시 경유) + Hacker News + arXiv.
이번 호 인덱스
- 오늘의 핵심: Ascend 저장소에 다섯 건의 커밋이 병합되어 알고리즘 측 연산자를 보완(GL M-5.3 kpool 압축, Chunk Scaled Dot KKT, W4A4 GEMM 수정, depthwise 컨볼루션 예시, 성능 프로파일링 skill), 메인 저장소는 전체 warp 리덕션을 redux.sync로 전환하고 autotuner의 독립 구성을 동시 컴파일로 변경, 24시간 내 12건 병합, 심사 중 업데이트 54건, 신규 issue 4건
-
- 핵심 프로젝트 진행 상황
- 1.1 메인 저장소: 전체 warp 리덕션을 redux.sync로 전환, 32 lane FP32 4개의 최댓값 연산이 20개의 shuffle에서 4개의 하드웨어 리덕션으로 변경(#3692)(10-29)
- 1.2 메인 저장소: autotuner의 독립 JIT 구성을 동시 컴파일로 변경, 락은 동일 키 후보에만 유지(#3476)(10-29)
- 1.3 메인 저장소: Ascend 버퍼 뷰의 레이아웃과 포인터 단위가 더 이상 별칭을 통해 전파되지 않음(#3623)(10-29)
- 1.4 메인 저장소: GM에서 L1로의 복사와 경계 초과 패딩을 하나의 스케줄 작업으로 병합, 불필요한 MTE2 펜스 제거(#3543)(10-29)
- 1.5 메인 저장소: CI가 ROCm과 Metal 작업을 호환되는 PyTorch wheel로 다시 고정(#3706)(10-29)
- 1.6 신규 결함 3건: 패딩이 있는 공유 버퍼의 배치 복사 시 형상 불일치 보고, 음수 스케일 반복 그래프의 역방향 인덱스 오컴파일, 뷰 비트 폭의 int32 오버플로(#3704 / #3705 / #3707)(10-29)
- 1.7 심사 중인 수정 3건: 비트 폭을 int64로 변경, CPU 스칼라 GEMM이 대상 영역만 초기화, 캐시 비활성화 시 두 계층의 프로세스 내 캐시 우회(#3708 / #3709 / #3710)(10-29)
- 핵심 프로젝트 진행 상황
-
- 멀티 백엔드 적응(Ascend / Hygon / Moore Threads / MetaX)
- 2.1 Ascend: GLM-5.3 희소 어텐션 indexer의 kpool 압축 연산자 병합, 단일 kernel 순수 Vector 경로(#1910)(10-29)
- 2.2 Ascend: 청크 게이티드 선형 어텐션의 Chunk Scaled Dot KKT 전방 연산자 병합, 정적 4블록 전개, GM 워크스페이스 제로(#1803)(10-29)
- 2.3 Ascend: W4A4의 int4 곱하기 int4 GEMM 결과가 전부 틀리는 문제 수정 병합, 근본 원인은 타입 크기 허위 보고와 전용 전치 명령의 조용한 비활성화(#1731)(10-29)
- 2.4 Ascend: depthwise_conv_2d 예시 재제출, 동시 배치를 피하기 위해 매니페스트화된 테스트 채널로 변경(#1920)(10-29)
- 2.5 Ascend: 연산자 성능 프로파일링 skill 추가, msprof op 출력으로 최적화 방향 결정(#1931)(10-29)
- 2.6 Ascend 심사 중: NPUIR 언어 계층의 뷰와 브로드캐스트 수정 3건(#1962 / #1960 / #1961)(10-29)
- 2.7 Ascend 심사 중: PTO 백엔드의 비트 단위 XOR 부호 없는 바이트 워크스페이스 뷰 수정(#1959)(10-29)
- 2.8 Hygon: HCU 2건 심사 중 업데이트 있음, 커뮤니티에 DCU 융합 연산자 저장소 등장, Moore Threads와 MetaX는 모니터링 기간 내 푸시 없음(10-29)
- 멀티 백엔드 적응(Ascend / Hygon / Moore Threads / MetaX)
-
- 생태계와 채택자(TileOPs / TileFoundry / TileKernels / FlashQLA / 서드파티)
- 3.1 TileOPs: 릴리스 이미지 병합, 연산자 라이브러리와 TileFoundry가 처음으로 같은 이미지에 포함(#2793)(10-29)
- 3.2 TileOPs 심사 중: GEMM이 전력 소비 상한에 따라 그리드 그룹화와 상주 블록 수를 재계산, q-b의 DRAM 읽기가 168 MB에서 94 MB로 감소(#2794)(10-29)
- 3.3 TileFoundry: AMX 런타임 연산자를 39개의 정적 TIR 계약으로 변경, 호환성 깨지는 변경(#279)(10-29)
- 3.4 채택자: TileKernels와 FlashQLA 메인 브랜치는 모니터링 기간 내 푸시 없음(10-29)
- 3.5 Ascend 측에 독립 커뮤니티 연산자 저장소 ops-tilelang 등장, 모니터링 기간 내 푸시 있음(10-29)
- 3.6 PTO-ISA 연동 두 저장소 활발: 문제 저장소가 PTOAS 0.70으로 이전, 벤치마크 저장소에 SV2 청크 스케일링과 열별 분석 추가(10-29)
- 3.7 관찰: T-HEAD 명의의 tilelang 브랜치 저장소에 모니터링 기간 내 푸시 있음, 브랜치 이름이 PPU의 행렬 명령 레이아웃을 가리킴(10-29)
- 생태계와 채택자(TileOPs / TileFoundry / TileKernels / FlashQLA / 서드파티)
-
- 커뮤니티, 튜토리얼 및 활동
- 4.1 문서 사이트가 내장 리덕션의 적용 경계를 명령 페이지에 명시, redux 병합과 동기화(10-29)
- 4.2 개발 도구 저장소가 다운스트림 추적 추출 후 3건 마무리, 업스트림 심사 중인 도구화 변경과 같은 방향(10-29)
- 4.3 신규 저장소 tilelang-knowledge 하루 4건: 연산자 설계 사례 33건, 출처가 있는 지식 카드 20장, 최적화 패턴 7건(10-29)
- 4.4 도해 강좌 부속 저장소와 분산 레이아웃 실험 저장소가 같은 날 신규 생성(10-29)
- 4.5 뉴스 측: Google News 중영문 쿼리 6조, Hacker News와 arXiv 모두 모니터링 기간 내 실질적 적중 없음(10-29)
- 커뮤니티, 튜토리얼 및 활동
-
- 트렌드 관찰
- 부록, 자료 및 검증 설명
오늘의 핵심
오늘의 병합량은 어제의 24건에서 12건으로 줄었으나 구조는 더 집중되었다: Ascend 전용 저장소가 5건, 메인 저장소가 5건, TileFoundry와 TileOPs가 각각 1건을 차지했다. Ascend 저장소의 이번 배치는 전부 알고리즘 측 연산자이며, 모두 업스트림 모델에 직접 사용할 수 있는 부류다. GLM-5.3-Flash 희소 어텐션 indexer의 kpool 압축 연산자는 차원별 softmax 풀링, BF16 경계, 정규화 Hadamard-128 회전과 E4M3 양자화를 하나의 순수 Vector kernel에 압축했으며, 각 블록이 64개의 pool을 커버하고 두 개의 벡터 서브코어가 각각 32개를 담당하며, 호스트 측 재배열이 전혀 없이 호출자의 3차원 레이아웃에 따라 스트라이드 복사를 수행한다; E4M3 인코딩은 순수 정수 비트 연산으로 구현되었으며(round-half-even, 서브노멀 포함, 448로 포화), FP8 타입에 의존하지 않고 근사도 하지 않는다. Chunk Scaled Dot KKT는 블록 게이트 선형 어텐션 포워드 내의 A = gate(K K^T)로, 정적 4블록 전개, GM 워크스페이스 제로, 호스트 연산자 제로이며, L0와 L1 세 가지 사례의 일치율은 모두 1.0이다.
메인 저장소의 오늘 주제는 컴파일과 발사 계층이다. 전체 warp 리덕션은 redux.sync로 전환되었다: 메인 브랜치의 네 개 독립 32레인 FP32 최댓값이 CUDA 13.1, sm_103a에서 20개의 shuffle로 전개되었으나, 변경 후에는 네 개의 redux.sync.max.f32가 되었고 shuffle은 포함되지 않는다; 정수 리덕션은 SM80 이상에서 하드웨어 리덕션을 사용하며, FP32의 최댓값/최솟값은 sm_100a, sm_100f, sm_103a, sm_103f에서만 사용 가능하고, 좁은 타입(CUTLASS의 FP16, BF16 포함)은 FP32로 승격되며, FP64는 여전히 shuffle 경로를 유지한다. 저자는 엔드투엔드 가속을 주장하지 않으며 명령어 수준 대조만 제시했다. autotuner의 동시 컴파일은 전체 구간에서 보유하던 컴파일 락을 분리했다: 각 커버리지 구성은 dataclasses로 독립적인 JITImpl을 복제하며, 1단계 키를 공유하는 후보만 직렬화되고, 디버그 출력의 디스크 쓰기는 여전히 직렬화된다. 저자가 제시한 크로스 백엔드 검증은 Ascend 1797건, ROCm 3220건, Metal 95건, CUDA 5500건이다.
기록할 만한 또 다른 세부 사항이 있다: 오늘 새로 열린 역방향 인덱스 오컴파일 결함(#3705)은 T-HEAD 명의의 tilelang 브랜치 저장소에서 런타임 재현되었으며, 이는 이미 조직 차원의 팀이 이 업스트림 코드를 자사 하드웨어 스택의 기반으로 삼아 추적하고 있음을 보여주며, 단순한 외부 의존성이 아님을 나타낸다.
1. 핵심 프로젝트 진행 상황
1.1 메인 저장소: 전체 warp 리덕션이 redux.sync로 전환, 네 개 32레인 FP32 최댓값이 20개 shuffle에서 4개 하드웨어 리덕션으로 (#3692)
10-29 15:08에 병합, 4개 파일 +337/-35. 이전의 일반적인 스레드 간 리덕션(예: T.reduce_max)은 그룹 전체가 정렬된 풀 warp여도 shuffle 트리로 전개되었다: 메인 브랜치 837486a에서 네 개의 독립 32레인 FP32 최댓값이 sm_103a + CUDA 13.1 컴파일 시 20개 shuffle, 0개 redux로 생성되었으나, 변경 후에는 네 개 redux.sync.max.f32, 0개 shuffle이 되었다. 방식은 스칼라 및 배치 전체 리덕션에 이미 존재하는 warp_reduce 헬퍼를 재사용하는 것이다: 호출자는 스레드 지오메트리(warp 정렬 여부, 단위 스트라이드, 풀 warp)만 판단하고, 아키텍처, 데이터 타입과 연산자 디스패치는 헬퍼에 위임한다. 정수 리덕션은 SM80 이상에서 하드웨어 redux를 사용하며, 더 작은 정수는 32비트로 승격된다; FP32의 최댓값/최솟값은 sm_100a, sm_100f와 sm_103a, sm_103f에서만 사용 가능하다; 좁은 부동소수점은 FP32로 승격된 후 헬퍼에 진입하며, FP64는 정밀도를 낮추지 않고 여전히 shuffle 경로를 사용한다. 서브 warp, 스트라이드, 패킹 리덕션 및 지원되지 않는 연산자와 타깃은 모두 shuffle로 폴백되며, 크로스 warp 공유 메모리 교환과 배리어는 변경되지 않는다. 부수적으로 컴파일 실패 한 건도 수정되었다: 사용자 정의 부호 없는 정수 reducer가 인라인 디스패치가 빈 타입을 반환하여 컴파일되지 않던 문제가 이제 shuffle 폴백으로 처리된다. PTX 회귀는 실제 명령어와 shuffle 부재를 단언하며, 런타임은 부호 있는/없는 경계값, 배치 미초기화, NVRTC 실행, 비유한값, 부호 있는 0과 FP64 정확한 결과를 커버한다. 저자는 엔드투엔드 가속을 명시적으로 주장하지 않는다. 기존의 유산 문제 하나도 설명에 명시되었다: 배치 리덕션에서 batch=4, 스레드당 2개 출력, batch=8, 스레드당 4개 출력 시 FP16과 BF16 실패가 존재하며, 메인 브랜치 원본 템플릿에서도 동일하게 재현되고, 실패는 하강 단계에서 발생하여 본 커밋과 무관하다.
출처: https://github.com/tile-ai/tilelang/pull/3692
1.2 메인 저장소: autotuner의 독립 JIT 구성이 동시 컴파일로 전환, 락은 동일 키 후보에만 유지 (#3476)
10-29 16:24에 병합되었으며, 2개 파일에서 +54/-46입니다. 메인 브랜치의 _make_jit_compile_func는 전체 컴파일 기간 동안 하나의 잠금을 유지하여, 기본 컴파일 옵션만 사용하는 구성조차 직렬화되어 프런트엔드, 다운스트림, 디바이스 컴파일이 모두 대기열에 쌓였습니다. 변경 방식은 각 재정의 구성마다 dataclasses.replace로 후보 JITImpl을 생성하되 백엔드, 타깃, 출력 선택, 로그 수준, 컴파일 플래그, 디버그 옵션은 원래 값을 유지하는 것입니다. 기본 옵션은 여전히 기존의 지연 호출 경로를 사용하여 커널 캐시와 객체 정체성을 보존합니다. 공유 1단계 키(get_tir에서 사용하는 인자 바인더로 계산되며 위치 인자, 기본값, 텐서 주석이 있는 팩토리를 보존)가 동일한 후보만 직렬화합니다. 디버그 파일의 공유 출력은 여전히 컴파일 내부에서 직렬화되며 debug_root_path, IR 덤프 설정, 레이아웃 시각화를 재정의합니다. 검증 데이터는 크로스 백엔드 파이프라인에서 확인되었습니다. Ascend 작업은 3be11be6과 메인 브랜치 c58a86f0을 병합했으며 Python 3.12.14, torch_npu 2.10.0, CANN 9.3.0 환경에서 1797개 검사가 통과했고, AscendC 예제 50개 전부 통과, PTO 예제 44개 통과 6개 스킵입니다. ROCm 7.2 작업은 gfx942에서 3220개 통과, 2765개 스킵입니다. Metal 작업은 macOS arm64에서 95개 통과, 5개 스킵입니다. CUDA 작업은 H100 PCIe, cu130에서 5500개 검사 통과, 491개 스킵, 예제 76개 통과 14개 스킵입니다. CuTeDSL 작업은 75개 통과 17개 스킵 2개 예외적 통과입니다.
출처: https://github.com/tile-ai/tilelang/pull/3476
1.3 메인 저장소: Ascend 버퍼 뷰의 레이아웃과 포인터 단위가 더 이상 별칭을 넘어 전파되지 않음 (#3623)
10-29 13:17에 병합되었으며, 23개 파일에서 +1139/-438입니다. T.view 또는 T.reshape로 생성된 여러 버퍼가 동일한 저장 영역을 공유하면서 서로 다른 레이아웃을 요구할 수 있는데, 이전에는 공유 데이터 변수를 기준으로 레이아웃을 일괄 해석하여 잘못된 Cube 매핑이나 DMA 주소가 선택될 수 있었고, 오래된 타일링 형상이 버전 오프셋을 잃거나 타입이 있는 포인터 오프셋을 잘못된 단위로 해석할 수 있었습니다. 이번 변경은 레이아웃을 개별 Buffer 뷰에 귀속시키며, 레이아웃 추론, 프랙탈 저장 정규화, 다중 버퍼 물질화와 복사 및 행렬 곱 다운스트림 전반에 적용됩니다. 뷰 정체성은 레이아웃 매핑과 조밀 L0 간격 적용 이후까지 유지되며, 베어 포인터 오프셋은 각자의 요소 단위로 재매핑됩니다. 저장 영역을 넘는 별칭의 레이아웃은 더 이상 전파되지 않으며, 각 뷰는 자신의 주석, 연산자에서 추론된 레이아웃 또는 일반 주소 지정을 사용합니다. 동일한 버전이 있는 할당에서 Cube 뷰는 이제 서로 다른 MajorK와 MajorMN 레이아웃을 사용할 수 있으며, 조건은 논리 스트라이드가 조밀하고, 베이스 오프셋이 0이며, 피연산자 타입이 호환되고, 패딩 후 저장 크기가 같아야 합니다. 다중 버퍼 별칭은 하나의 물리적 슬롯 크기를 공유하며 각 뷰의 요소 단위로 환산하고 인덱스 매핑은 전파하지 않습니다. 뷰는 기존 바이트를 재해석할 뿐이므로 기록자는 소비자가 기대하는 표현으로 기록해야 하며, 내장 명령에 뷰를 선택해 전달해야 할 경우 액세스 포인터 인터페이스를 사용할 수 있습니다.
출처: https://github.com/tile-ai/tilelang/pull/3623
1.4 메인 저장소: GM에서 L1로의 복사와 경계 초과 패딩을 하나의 스케줄 작업으로 병합하고 불필요한 MTE2 펜스 제거 (#3543)
10-29 11:39에 병합되었습니다. 전역 메모리에서 L1로의 복사와 그 경계 초과 패딩은 동일한 L1 데이터를 함께 준비하는데, 이전에는 패딩이 별도 스케줄 작업으로 분리될 수 있어 동기화 삽입 시 복사와 패딩 사이에 불필요한 MTE2 파이프라인 펜스가 추가되었습니다. 이번 변경은 각 복사와 그 패딩을 동일한 Ascend 작업으로 편성하여 소비자가 패딩을 포함한 완전한 데이터를 대기하도록 합니다. FP8 MX의 K 꼬리 패딩도 동일하게 병합하고 작업 내 조건부 패딩을 유지하며, 기존의 명시적 작업 스코프는 그대로 유지되고 중첩을 도입하지 않습니다. 기존의 MX 지연 패딩 큐는 삭제되어 각 피연산자의 패딩이 자신의 복사 바로 뒤에 따릅니다. 일반 경계 초과 패딩과 MX 꼬리 블록 두 종류를 포함한 47개 집중 테스트가 통과했습니다.
출처: https://github.com/tile-ai/tilelang/pull/3543
1.5 메인 저장소: CI가 ROCm 및 Metal 작업을 호환되는 PyTorch wheel로 다시 고정 (#3706)
10-29 21:38에 병합되었습니다. ROCm CI는 인덱스 변수를 통해 일치하는 PyTorch ROCm 소스를 노출하지만, 워크플로 수준의 “안전하지 않은 최적 매칭” 정책이 여전히 더 늦은 PyPI 후보를 선택하도록 했습니다. #3534의 ROCm 작업에서 리졸버가 사용 가능한 2.14.1 + ROCm 7.2 빌드 대신 PyPI의 PyTorch 2.15.0(CUDA 13.2)을 선택했으며, 공식에서 러너에서 두 가지 해석 결과를 직접 재현했습니다. 동일한 2.15.0 배포는 Metal 작업도 중단시켰습니다. macOS 26.6.2의 ARM64 러너에서 2.15.0은 Metal 2.0을 노출하지 않아 Apple 반가상 GPU를 거부했으며, 동일 이미지에서 이전에 통과한 것은 2.14.1이었습니다. 변경 방식은 세 가지입니다. 구성된 안정 ROCm 버전에 따라 백엔드 변수를 설정하여 uv가 PyTorch를 ROCm 소스에 유지하고 나머지 의존성은 최적 매칭을 사용하도록 합니다. 야간 PyTorch 설치를 첫 번째 인덱스로 전환하여 야간 작업이 구성된 소스를 벗어날 수 없도록 합니다. 테스트 스위트 시작 전에 PyTorch가 HIP 빌드인지, CUDA 빌드가 아닌지, 러너 GPU에 액세스할 수 있는지 확인하는 빠른 실패 검증 단계를 추가합니다.
출처: https://github.com/tile-ai/tilelang/pull/3706
1.6 신규 결함 세 건: 패딩이 있는 공유 버퍼가 일괄 복사로 갈 때 형상 불일치 보고, 음수 스케일 반복 그래프가 역방향 인덱스로 오컴파일, 뷰 비트 폭이 int32에서 오버플로 (#3704 / #3705 / #3707)
세 이슈 모두 정적 검사의 경계를 가리킨다. 첫째, #3704는 마지막 차원이 16 정렬이 아닌 전역 텐서(예: K=100)를 패딩이 있는 공유 버퍼(BK가 cdiv(K,16) 곱하기 16으로 112)에 복사할 때, 이 복사가 일괄 복사 경로로 라우팅되면 하향 변환이 곧바로 공유 블록과 전역 블록의 형상 불일치(112 대 100)를 보고한다고 한다. 위치는 CUDA 복사의 일괄 하향 변환 지점이며, 원소별 경로에서의 패딩 버퍼는 정상인데, 패딩은 행렬 곱 정렬을 만족시키는 흔한 작성법이므로 일괄 복사와 함께 쓰는 어떤 커널이라도 밟을 수 있다. 둘째, #3705는 평탄화 버퍼 단순화가 음수 스케일 반복 그래프에서 역방향 인덱스를 오컴파일한다고 보고한다. 재현 환경은 v0.1.15와 T-HEAD 명의의 tilelang 브랜치 저장소 소스 빌드이며, 보고는 동시에 업스트림 메인 브랜치와 해당 TVM 리비전에서 영향받는 분할 함수가 변하지 않았음을 확인했으므로 업스트림 공통 경로에 속한다. 셋째, #3707은 뷰와 재형성이 버퍼 비트 폭을 확인할 때 int32로 곱하여 2의 31제곱을 넘으면 즉시 되돌아간다고 보고한다. Qwen3.8의 lm_head 가중치는 248320 곱하기 2560 바이트의 NVFP4로 총 5,085,593,600비트인데, 비트 폭 곱이 790,626,304으로 되돌아가고, 이어서 정적 뷰 비트 수를 int32 상수로 변환할 때 리터럴 초과 예외를 던진다.
출처: https://github.com/tile-ai/tilelang/issues/3704
출처: https://github.com/tile-ai/tilelang/issues/3705
출처: https://github.com/tile-ai/tilelang/issues/3707
1.7 세 건의 수정 심사 중: 비트 폭을 int64로 변경, CPU 스칼라 GEMM은 대상 영역만 초기화, 캐시 끌 때 두 계층의 프로세스 내 캐시 우회 (#3708 / #3709 / #3710)
위 세 결함과 어제의 캐시 의미론 의제에 대응하여 수정이 같은 날 심사에 들어갔다. 비트 폭 건은 정적 차원 곱을 Python 정수 곱으로 바꾸고, 기호화 차원을 int64로 변환하며, int32를 초과하는 리터럴을 int64 상수로 변환한다. 회귀 사례는 해당 이슈의 버퍼(248320 곱하기 5120의 뷰 형상)를 그대로 취하며, 메인 브랜치에서는 오류가 나고 수정 후에는 통과하며, 기호화 차원의 동형 뷰도 커버한다. CPU 건은 스칼라 행렬 곱이 누산을 0으로 초기화할 때 하위 버퍼 전체를 초기화하는 문제를 수정한다. 4 곱하기 4 출력 버퍼를 7로 초기화한 뒤 1부터 3 행렬의 부분 영역을 호출하면, 원래 구현은 영역 밖 12개 원소까지 함께 0으로 만들지만, 수정 방식은 출력 영역 자체의 오프셋과 범위로만 초기화하며, 회귀 사례는 이에 따라 전부 1인 2 곱하기 2로 작성했다. 캐시 건은 어제 기록한 #3699를 수정한다. 커널 캐시를 끄면 캐시 상태만 뒤집혔는데, jit 래퍼는 실제 컴파일 전에 두 개의 프로세스 내 캐시를 더 조회하므로 블록 내 일반 호출은 같은 프로세스에서 앞서 컴파일된 커널을 얻었다. 수정 방식은 호출 시 캐시 스위치를 읽고(환경 변수 비활성화도 함께 커버), 비활성 기간에는 두 계층의 프로세스 내 캐시를 건너뛰며, 블록을 나온 뒤 복원한다. 테스트는 수정 전 1 실패 4 통과에서 5 전부 통과로 바뀌었다.
출처: https://github.com/tile-ai/tilelang/pull/3708
출처: https://github.com/tile-ai/tilelang/pull/3709
출처: https://github.com/tile-ai/tilelang/pull/3710
2. 멀티 백엔드 적응 (Ascend / Hygon / Moore Threads / MetaX)
2.1 Ascend: GLM-5.3 희소 어텐션 indexer의 kpool 압축 연산자 병합, 단일 kernel 순수 Vector 경로 (#1910)
10-29 15:41에 병합되었다. 연산자는 각 pool(pool_size가 4)에 대해 점수 텐서와 위치 인코딩을 더한 뒤 차원별 softmax 풀링을 하고, 풀링 벡터를 BF16 경계로 통과시킨 뒤 정규화된 Sylvester Hadamard-128 회전(128의 제곱근으로 나눔)을 곱하고, 마지막으로 각 벡터의 FP32 2의 거듭제곱 절대 최댓값 스케일링으로 E4M3 양자화를 한다. E4M3 행과 스케일 인자는 연속 임시 영역에 기록되고, 호스트 어댑터가 페이지 캐시로 산개시킨다(키 캐시는 부호 없는 바이트로 E4M3를 전달하고, 스케일 캐시는 제자리에서 FP32로 기록). 구현은 단일 kernel 순수 Vector 경로이며, 개발자 모드에서 자동 동기화, 자동 벡터와 행렬 코어 병합, 메모리 계획 세 가지 스위치를 켜고, 각 블록은 64개 pool을 커버하며 두 벡터 서브코어가 각각 32개를 독립 병렬 처리한다. 호스트 재배열은 없고, kernel은 호출자의 3차원 레이아웃대로 스트라이드 읽기를 하며, 위치 인코딩은 풀링에 접어 넣는다. E4M3 인코딩은 순수 정수 비트 연산으로, round-half-even, 서브노멀 포함, 448로 포화하며, FP8 타입을 사용하지 않고 근사도 하지 않는다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1910
2.2 Ascend: 분할 게이트 선형 어텐션의 Chunk Scaled Dot KKT 전방 연산자 병합, 정적 네 블록 전개, GM 작업 영역 제로 (#1803)
10-29 15:01에 병합되었습니다. 블록 단위 내적의 스케일링된 점곱 전방 계산이 추가되었습니다. 블록 게이티드 선형 어텐션의 전방 계산에서 A는 게이트 게이팅된 K 곱하기 K 전치와 같습니다. 단일 컴파일러 데코레이션 kernel로 구현되었으며, 네 개의 블록을 정적으로 전개하여 온칩 융합을 수행하고, 전역 메모리 작업 공간 없음, 호스트 연산자 호출 없음, 개발자 모드의 네 가지 스위치를 모두 켭니다. 테스트는 L0, L1, L2와 경계 네 단계로 계층화되며 명령줄에서 단계를 선택할 수 있습니다. L0는 세 가지 상황(게이팅 있음, 게이팅 없음, 소규모 게이팅 있음)을 커버하고, 정밀도는 PyTorch CPU의 FP32 금본과 대조하며 혼합 허용 오차 이중 게이트를 사용합니다. 보고서에서 제시한 테스트 형상은 배치 1, 시퀀스 길이 32768, 헤드 32, 헤드당 차원 128, 블록 64, 게이팅 활성화이며, 입력과 출력은 BF16, 누적은 FP32입니다. L0와 L1 각각 세 예시가 모두 통과했고, 일치율 1.0000, 최대 절대 오차는 0에서 0.25 사이입니다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1803
2.3 Ascend: W4A4의 int4 곱하기 int4 GEMM 결과 전면 오류 문제 수정 병합, 근본 원인은 타입 크기 허위 보고와 전용 전치 명령이 조용히 비활성화된 것 (#1731)
어제까지 심사 중이던 이 건이 오늘 15:08에 병합되었으며, #438에 해당합니다. 근본 원인은 세 층의 오류가 겹친 것입니다. 첫째, 프랙탈 기하가 타입 크기 1 기준으로 유도되었으나, 니블 패킹 타입은 각 요소가 4비트만 차지하고 바이트당 두 요소를 담으므로 32바이트 C0 블록의 올바른 요소 수는 32가 아니라 64입니다. 이로 인해 모든 복사 스트라이드, 반복 횟수, 바이트 계산이 전부 두 배씩 어긋나 데이터가 L0A와 L0B에서 잘못된 위치에 쌓였습니다. 둘째, B 피연산자의 전치 로드가 조용히 무시되었습니다. 범용 로드 경로는 zN에서 nZ로의 변환에 전치 플래그를 설정하지만, 4비트 전용 명령은 전치를 하드코딩으로 끄고 해당 플래그를 무시하므로 B는 행렬 곱 명령이 기대하는 레이아웃을 절대 얻지 못합니다. 올바른 경로는 전치를 포함한 전용 로드이며, 이는 CANN 헤더 파일에 문서화가 전혀 없어 프로브 kernel로 디바이스에서 실측하여 매개변수를 역추론해야 합니다(한 번 반복으로 전체 블록을 옮기고, 반복 횟수는 N 나누기 64, 소스 스트라이드는 K를 16으로 올림한 뒤 64로 나눈 값, 목적 간격은 3). 또한 이 버전은 실제로 2048바이트 블록 단위로 동작하며, 범용 문서의 512바이트 프랙탈 단위와 다릅니다. 셋째, 벡터와 전역 메모리 간 바이트 계산 역시 요소당 1바이트가 아닌 0.5바이트로 계산되어 각 행 복사 길이가 두 배로 길어져 인접 행 데이터가 섞였습니다. 계산 명령 자체는 항상 올바르며, 단지 잘못 배치된 데이터를 충실히 곱했을 뿐입니다. 증상은 99.6% 불일치, 출력이 음수 501과 627인데 올바른 값은 641과 1785입니다. 보고서는 전역 메모리의 int4 패킹 규약도 명확히 기술했습니다. 바이트당 두 니블, 낮은 니블이 짝수 요소, 높은 니블이 홀수 요소이며, r번째 행은 r 곱하기 N 나누기 2의 바이트 위치에서 시작하고 CANN 자체 형식과 일치합니다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1731
2.4 Ascend: depthwise_conv_2d 예제 재제출, 동시 배치를 피하기 위해 매니페스트화된 테스트 채널로 전환 (#1920)
10-29 15:43에 병합되었습니다. 이는 #1915의 재제출입니다. 원래 건이 병합된 후 야간 전체 CI가 실패하여 #1919로 롤백되었으며, 근본 원인은 L0 최대 형상 케이스가 8로우 동시 실행되는 레거시 스크립트 배치에서 알 수 없는 오류를 유발한 것입니다. 이번 건은 해당 케이스를 4분의 1로 축소하고, 연산자를 지속적 통합의 연산자 테스트 매니페스트에 등록하여 Pytest 채널로 전환했으며, 이미 마이그레이션된 다른 25개 연산자와 동일한 경로를 따릅니다. 동시 배치에서 완전히 철수했고 kernel 본체는 원래 건과 바이트 단위로 동일합니다. 연산자 자체는 fp16, fp32, bf16 및 임의의 커널 크기, 스트라이드, 패딩, 팽창을 지원하며, 스트라이드 1, 2, 4는 양자화 영역 경로로, 나머지는 스칼라 내적 폴백으로 갑니다. 구현은 패딩과 컨볼루션 이중 kernel에 1 곱하기 1 고속 경로를 더한 것이며, 순수 Vector 경로로 그리드는 형상에 따라 자동 적응하고 커널 수를 하드코딩하지 않습니다. 보고서의 몇 가지 실측 제약은 기록할 만합니다. 패딩 단계 소요 시간은 대략 문장 수 곱하기 35나노초에 해당하며 바이트 수와 무관하므로 마지막 밴드는 절단으로 마무리합니다. 공유 메모리 측 복사 대상 오프셋이 32바이트 정렬이 아니면 벡터 코어의 0x26 예외가 발생합니다. 동기화는 수작업 챔피언 모델을 사용하며(컨볼루션 구간은 자동 동기화 끔), 채널 간에는 깊이 1의 인접 이동과 계산 이벤트 프로토콜을 사용하며, 깊이 2의 프로토콜이나 중간에 전체 펜스를 삽입하면 교착 상태가 발생합니다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1920
2.5 Ascend: 연산자 성능 프로파일링 skill 추가, msprof op 출력으로 최적화 방향 설정 (#1931)
10-29 17:17에 병합되었습니다. 연산자 성능 프로파일링 개발 프로세스 자산이 새로 추가되었습니다. 성능 최적화 전에 msprof의 연산자 수준으로 한 번 실행하고, 출력에 따라 다음 최적화 방향을 판단합니다. 작성자는 덧셈 요소별, 리덕션 최솟값 예제, 동적 양자화 세 연산자에서 프로세스가 사용 가능함을 검증했습니다. 이러한 자산의 가치는 매번 프로파일링 기준을 새로 탐색하는 대신 “먼저 측정하고 나서 변경”을 프로세스에 명문화하는 데 있습니다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1931
2.6 Ascend 심사 중: NPUIR 언어 계층 세 건의 뷰와 브로드캐스트 수정 (#1962 / #1960 / #1961)
세 건이 같은 날 검토에 들어갔으며 서로 의존한다. 첫째, #1962는 로컬 통합 버퍼의 뷰 디스크립터를 수정한다. 브리지 계층이 공유 데이터 변수로 버퍼를 조회할 때 디스크립터가 적용되지 않은 할당 메모리 참조를 반환할 수 있어, 동일 유형 오프셋 별칭은 할당 영역 접두부를 읽게 되고, 다른 유형 별칭은 할당된 요소 유형을 지닌 채 복사 강하로 진입하며, 동일 유형 재형상은 잘못된 차원을 지닌 채 리덕션으로 진입할 수도 있다. 수정 방식은 통합 뷰 획득 헬퍼를 도입하고, 동일 유형의 비항등 뷰는 검사가 포함된 재해석 변환을 사용하며, 다중 요소 유형 접근의 할당은 단일 블록 바이트 할당에 타입 지정 뷰 연산을 더해 공통 저장 정체성을 유지하는 것이다. 둘째, #1960은 뷰가 원본 오프셋을 잃는 문제를 수정한다. 뷰가 데이터 변수로부터 버퍼를 재구성할 때 요소 오프셋을 잃어, 오프셋 별칭이 원본 첫 요소 주소가 아닌 할당 오프셋 0에서 시작하게 된다. 수정 방식은 동일 유형에 대해서는 오프셋을 유지하고, 바이트 수준 유형 변경 시 요소 폭에 따라 환산하며 나누어떨어짐을 요구하고, 정렬되지 않거나 지원되지 않는 변환은 즉시 거부하며, 저장소나 데이터 복사를 도입하지 않는 것이다. 셋째, #1961은 단항 벡터 브로드캐스트를 접근 영역 기준으로 추론하도록 변경한다. 기존에는 원본 버퍼 형상으로 브로드캐스트 차원을 추론하여 8 곱하기 8 소스를 24 곱하기 8 대상의 행 8에서 16에 쓰면 실패했다. 수정 방식은 생성된 메모리 참조 유형 기준으로 추론하도록 바꾸고, 접근 구간의 형상과 기존 호환성 검사를 유지하는 것이다. 세 건 모두 독립적인 클린 베이스라인 빌드와 비트 단위 검증 설명을 첨부했다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1962
2.7 Ascend 검토 중: PTO 백엔드 비트별 XOR의 부호 없는 바이트 워크스페이스 뷰 수정 (#1959)
10-29 16:36에 검토에 들어갔으며 #1721을 닫는다. PTO 백엔드에서 비트별 XOR의 임시 워크스페이스가 필요한 2차원 뷰가 아닌 원시 1차원 바이트 영역으로 명령에 전달되었다. 생성 시 목적 측의 2차원 형상을 워크스페이스의 형상 정보에 복사했으나 슬라이스로 표시하지 않아, 부호 없는 바이트(워크스페이스는 타입 재해석을 하지 않음)에 대해 1차원 블록 디스크립터가 해석되어 명령 실행 시 범위를 벗어나 읽고 쓰게 되었다. 워크스페이스 자체는 원본 요소 수에 따라 1차원 영역으로 할당되며, 용량은 정렬된 2차원 블록에 충분하고 실패는 용량이 아닌 디스크립터에 있었다. 수정 방식은 목적 형상을 상속한 후 2차원 뷰를 강제로 구성하고 슬라이스로 표시하는 것이며, 테스트에는 부호 없는 및 부호 있는 바이트의 여러 조합(2차원 자동 임시, 명시적 1차원 임시, 1차원, 정렬된 비영 오프셋 슬라이스에 센티널 추가, 부호 있는 바이트)을 보강하고 16비트 회귀도 유지했다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1959
2.8 Hygon: HCU 두 건 검토 중 업데이트, 커뮤니티에 DCU 융합 연산자 저장소 등장; Moore Threads와 MetaX는 모니터링 기간 내 푸시 없음 (10-29)
Hygon 측은 여전히 검토 중이다. 일반 복사의 공유 메모리 수직 인터리브 특성 브랜치(#30)와 분산 관련 건(#19) 모두 모니터링 기간 내 업데이트가 있었으며, 메인 브랜치 최신 커밋은 여전히 10-27의 버퍼 영역 그룹화 특성에 머물러 있다. 커뮤니티 측에서는 같은 날 DCU 방향의 저장소가 등장했는데, Hygon 플랫폼의 융합 연산자와 TileLang 컴파일러 변경 사항을 함께 보존하고 있으며 C++ 위주, 약 15 MB 규모로, 과거 커밋 이력에서 어텐션 연산자 지원과 가변 길이 지원을 확인할 수 있다. 이는 개인 실습 저장소이며 Hygon 공식은 아니다. Moore Threads 적응 저장소(tilelang-musa)와 MetaX 적응 저장소(tilelang-metax)는 모니터링 기간 내 푸시가 없었고, 최신은 각각 09-30과 10-26이다.
출처: https://github.com/tile-ai/tilelang-hygon/pull/30
출처: https://github.com/SealSoul117/TileLang_DCU_FuseCmpComm
3. 생태계와 채택 측 (TileOPs / TileFoundry / TileKernels / FlashQLA / 서드파티)
3.1 TileOPs: 릴리스 이미지 병합, 연산자 라이브러리와 TileFoundry가 처음으로 같은 이미지에 포함 (#2793)
10-29 11:18에 병합되었으며 9개 파일 +123/-49. 러너 이미지에 릴리스 대상이 추가되었다. 개발 이미지 위에 연산자 라이브러리와 스케줄러의 PyPI 릴리스 버전을 설치하고 두 버전 번호로 명명한다. 이미지 태그는 tileops0.1.0에 tilefoundry0.0.3, 그리고 cuda13.2, torch2.13, tilelang0.1.12 세 가지 환경 식별자를 더한 형태이며, 컨테이너 이미지 저장소의 연산자 라이브러리 릴리스 경로에 게시된다. 추가 스크립트는 이미지를 당겨 GPU에서 스택 검사, 베이스라인 검사, 패키징 테스트와 선택적 벤치마크 파일을 실행하는데, 휠에 테스트와 벤치마크가 포함되지 않기 때문에 모두 이미지 자체 버전에 대응하는 태그에서 가져온다. 문서는 빠른 시작을 업데이트했고 러너 문서에는 빌드, 검증, 푸시 단계를 보강했다. 동시에 sgl-kernel 베이스라인을 제거했다. 더 이상 이미지에 설치되지 않으며, 관련 벤치마크는 이전에 계속 건너뛰고 있었다.
출처: https://github.com/tile-ai/TileOPs/pull/2793
3.2 TileOPs 검토 중: GEMM 전력 상한 기준 그리드 그룹화 및 상주 블록 수 재계산, q-b의 DRAM 읽기 168 MB에서 94 MB로 감소 (#2794)
10-29 21:44에 병합 검토에 들어갔으며, 파일 2개 +35/-12로 백투백 호출의 성능 이슈를 이어받는다. 두 가지 변경: 첫째, A가 2차 캐시 파티션에 들어갈 수 있고 B는 들어갈 수 없을 때, 그리드 그룹화가 하나의 그룹을 전체 행에 걸치게 하여 각 B 블록이 16행 그룹마다 한 번이 아니라 메모리에서 한 번만 읽도록 한다. DeepSeek V3 프리필의 q-b 형상(4096 x 24576 x 1536, NT 레이아웃)에서 메모리 읽기가 168 MB에서 94 MB로 줄었다. 좁은 블록의 분할 경계는 여전히 16행 그룹 제약을 따르며 선택은 변하지 않는다. 둘째, 상주 블록 수를 폭이 균일한 이중 협력 그룹 그리드에서 프로세서 수를 직접 쓰는 대신 블록 수, 총 블록 수, 프로세서 수로 올림 계산한 상주 수로 변경한다. 가장 바쁜 블록은 어차피 같은 수의 블록을 실행해야 하고, 각 웨이브에서 더 적은 수의 서로 다른 패널을 읽을 수 있기 때문이다. 핑퐁, 좁은 블록, 혼합 폭 그리드는 여전히 프로세서 수를 사용한다. 테스트는 H200, CUDA 13.2, PyTorch 2.13, TileLang 0.1.12에서 연산자 테스트를 실행하여 212개 통과, 24개 건너뜀.
출처: https://github.com/tile-ai/TileOPs/pull/2794
3.3 TileFoundry: AMX 런타임 연산자를 39개의 정적 TIR 계약으로 변경, 파괴적 변경에 해당(#279)
10-29 15:59에 병합, 파괴적 변경(커밋 제목에 느낌표 포함). 동기는 AMX에 명시적이고 타입이 지정된 연산자 선언을 두어 에이전트가 동일한 계산 설명에서 매칭하고 평가할 수 있게 하는 것이다. 방식은 24개의 AMX 런타임 연산자 패밀리를 파라미터 영역, 피연산자와 레이아웃 패턴, 부작용 집합, M1~M4 대상 능력을 갖춘 39개의 정적으로 등록된 계약 버전으로 기술하는 것이다. Python 평가기, 원시 C 시뮬레이터를 골드 스탠더드로 추가하고, 59개의 대표 원자 테스트 케이스로 전체 39개 정적 연산자를 커버하며, 고계 및 중계 중간 표현 테스트 케이스는 실제 2, 4, 8, 12 스레드 영역에서 실행되어 후보의 수용과 거부를 검증한다. 고계 후보는 연산자 내 매칭 선언에서 발견되며 대상에 대한 능력 추론을 유지한다. 명명은 AMX 네임스페이스로 통일하고, 별칭은 등록 진입점이 보유하며, 등록 이름은 통합 조회로 제공된다. CPU 실행은 스레드 토폴로지를 사용하고, AMX 레지스터 상태를 논리적으로 사유하며, 스레드 수는 AMX 엔진 수가 아니다. 검증은 46개의 집중 테스트 통과이며, 57개 원시 C 결과 비교 1건과 라이프사이클 검사 2건을 포함한다.
출처: https://github.com/tile-ai/TileFoundry/pull/279
3.4 채택자: TileKernels와 FlashQLA 메인 브랜치는 모니터링 기간 내 푸시 없음(10-29)
알려진 두 채택자의 메인 브랜치는 모두 10-14에 멈춰 있으며 모니터링 기간 내 새 푸시가 없다. DeepSeek의 연산자 저장소(2018 스타)와 Tongyi 측의 FlashQLA(742 스타)이다. 당일 생태계 측의 신규 신호는 주로 이 두 내부 사용자가 아닌 커뮤니티 자체 저장소에서 나왔다.
출처: https://github.com/deepseek-ai/TileKernels
출처: https://github.com/QwenLM/FlashQLA
3.5 Ascend 측에 독립 커뮤니티 연산자 저장소 ops-tilelang 등장, 모니터링 기간 내 푸시 있음(10-29)
이 저장소는 10-20에 생성되었고 10-29 13:06에 푸시가 있었다. Ascend NPU를 위한 커뮤니티 고성능 연산자 저장소를 표방하며, kernel 구현, 자동화 테스트 케이스, 표준 워크로드, 성능 베이스라인을 통일된 방식으로 관리하여 연산자 자산을 재사용 가능하고 검증 가능하며 지속적으로 진화 가능하게 하는 것이 목표이다. tile-ai 조직이 아닌 Ascend 생태계의 포크 조직 산하에 있어, Ascend 측 연산자 공급이 공식 저장소에서 커뮤니티 협업 면으로 한 겹 넘치고 있음을 보여준다. 공식 tilelang-ascend와의 관계는 아직 불명확하며, 수록 시 생태계 관찰로만 다루고 공식 행동으로 보지 않는다.
출처: https://github.com/Atlas-Ascend-Fork/ops-tilelang
3.6 PTO-ISA 부속 이중 저장소 활발: 문제 저장소가 PTOAS 0.70으로 마이그레이션, 벤치마크 저장소에 SV2 청크 스케일링과 열별 분석 추가(10-29)
PTO-ISA 조직 산하 두 부속 저장소 모두 모니터링 기간 내 커밋이 있었다. 문제 저장소는 PTO의 브로드캐스트 로드를 PTOAS 0.70으로 마이그레이션했으며, 직전 주에는 양자화 스케일링 기호 명확화와 문제 설명 “답이 아닌 자신의 풀이를 가리키게” 수정 작업을 했다. 벤치마크 저장소는 두 건의 분석 커밋을 추가했다: SV2의 청크 스케일링과 열별 분석, 그리고 하나의 근거(라이브러리 내 절대 최대값 명령어 변형 수 31.39, 런타임 명령어 처리량과 루프라인 분석 첨부, 통합 버퍼 기준이 아님을 명시). 이 두 저장소는 공식 저장소와 교수 및 측정 부속 관계이며, Ascend 측 방법론의 방증으로 볼 가치가 있다.
출처: https://github.com/PTO-ISA/tilelang-puzzles-ascend
출처: https://github.com/PTO-ISA/tilelangGym
3.7 관찰: T-HEAD 산하 tilelang 포크 저장소 모니터링 기간 내 푸시 있음, 브랜치 이름이 PPU의 행렬 명령어 레이아웃을 가리킴(10-29)
T-HEAD 조직(공식 웹사이트는 개발자 포털 도메인) 산하에 tilelang 브랜치 저장소가 있으며, 2026-09-22에 생성되었고 업스트림 메인 저장소의 브랜치로, 13개 브랜치를 공개하고 있다. 메인 브랜치와 v0.1.11부터 v0.1.14까지의 버전 미러 외에도 PPU 행렬 공유 레이아웃 수정, PPU language 패키지, 런타임 컴파일 기능 및 여러 결함 수정 브랜치가 보인다. 해당 저장소의 메인 브랜치 최근 커밋은 10-22에 멈춰 있지만, 저장소 푸시 시간은 10-29 19:27(베이징 시간)로 갱신되었다. 저장소 내에는 설명 문서도 릴리스 기록도 없으므로 관찰 단서로만 활용할 수 있다: 전담 팀이 TileLang을 자체 하드웨어 스택 기반으로 삼아 브랜치 개발을 진행하고 있으며, 방향은 PPU의 행렬 명령 레이아웃과 관련되어 있다. 오늘의 #3705 결함 보고는 바로 이 저장소의 소스로 빌드하여 재현한 것이다.
출처: https://github.com/t-head/tilelang-for-sail
4. 커뮤니티, 튜토리얼 및 활동
4.1 문서 사이트가 내장 리덕션의 적용 경계를 명령 페이지에 명시, redux 머지와 동기화(10-29)
문서 사이트 모니터링 기간 내에 다섯 건의 자동 업데이트가 있었으며, 그중 명령 페이지에 리덕션 설명 단락이 새로 추가되어 오늘 머지된 하드웨어 리덕션 개편과 정렬되었다: CUDA의 스레드 간 리덕션은 warp 정렬되고 단위 스레드 스트라이드인 완전한 warp 그룹을 warp 리덕션과 동일한 소스의 헬퍼에 위임하며, 헬퍼는 SM80 이상의 정수 합산, 최댓값, 최솟값 및 비트 연산에서 하드웨어 리덕션 명령을 선택하고 더 작은 정수는 32비트로 승격한다. 스칼라 FP32의 최댓값/최솟값은 sm_100a, sm_100f와 sm_103a, sm_103f에서만 사용되며, 스칼라 FP16과 BF16은 FP32로 승격된다. 동일한 규칙이 더 큰 리덕션의 마지막 구간과 언패킹된 배치 리덕션에도 적용된다. 서브 warp, 스트라이드, 패킹, FP64 및 지원되지 않는 타입 연산자 조합은 shuffle 구현을 유지하며, warp 간 공유 메모리 교환과 펜스는 변경되지 않는다. 폴백 경계를 문서에 명시하는 것이 “하드웨어 리덕션 지원”만 기술하는 것보다 사용자에게 더 유용하다.
출처: https://github.com/tile-ai/tilelang.github.io
4.2 개발 도구 저장소, 하향 추적 추출 후 세 건 마무리 완료, 업스트림의 심사 중 도구화 변경과 동일 방향(10-29)
독립 개발 도구 저장소는 10-20에 메인 저장소에서 추출한 하향 추적을 완료한 후, 모니터링 기간 내에 세 건의 마무리를 보충했다: 추적 활성화 시 순환 임포트 방지, 지속적 통합을 계측 플러그인을 지원하는 버전으로 고정, 핵심 호환 별칭 로드 전에 TileLang 먼저 초기화. 이는 업스트림에서 심사 중인 “하향 추적을 개발 도구 저장소로 추출”하는 변경과 동일한 방향으로, 도구화가 더 이상 업스트림의 브랜치 실험이 아니라 독립적으로 릴리스 가능한 패키지로 형성되었음을 보여준다.
출처: https://github.com/easy-tilelang/tilelang-devkit
4.3 신규 저장소 tilelang-knowledge 하루 네 건: 33개 연산자 설계 사례, 출처 포함 지식 카드 20장, 최적화 패턴 7개(10-29)
해당 저장소는 10-29에 생성되었으며, 당일 네 건의 커밋으로 에이전트 지향 지식 기반을 구성했다: 먼저 저장소 구조, 이어서 33개 연산자 설계 사례 정리(각 항목에 고정 출처 포함), 그다음 출처 포함 프로그래밍 지식 카드 20장 보충, 마지막으로 증거 기반 GPU 최적화 패턴 7개 문서화. 저장소 자체 설명의 용도는 에이전트가 연산자를 마주했을 때 어떤 경험이 적용 가능한지 판단하고, 코드를 어떻게 수정해야 하는지, 어떻게 검증해야 하는지, 실패 후 어디서 찾아야 하는지를 알 수 있게 하는 것이다. 이는 TileFoundry의 연산자 계약, Ascend 저장소의 프로파일링 skill과 같은 종류의 시도이다: “어떻게 수정하고 어떻게 검증하는가”를 검색 가능한 자산으로 작성하는 것.
출처: https://github.com/qiuuuuu622/tilelang-knowledge
4.4 그림 해설 강좌 동반 저장소와 분산 레이아웃 실험 저장소 동일 날 신규 생성(10-29)
두 개인 저장소가 동일 날 생성되었다: 하나는 그림 해설 강좌 동반 저장소로, 독창적 장면과 애니메이션으로 각 지식 포인트를 완벽히 설명하고 영상과 함께 배포한다고 자술하며, 다른 하나는 분산 레이아웃 실험 저장소로, 커밋 설명은 분산 TileLang 레이아웃 실험 추가이다. 오늘 푸시가 있는 모델 이식 저장소와 Hexagon 백엔드 저장소(저장소 푸시 시간은 10-30 06:10으로 갱신되었으나 메인 브랜치 최근 커밋은 여전히 08-24의 의존성 업그레이드)를 더하면, 커뮤니티의 교육과 방언 백엔드 두 방향에서의 자발적 투입이 계속되고 있다.
출처: https://github.com/zachshi-ai/tilelang-course
출처: https://github.com/Leewin239/tilelang-distributed-layout
4.5 뉴스 측: Google News 6개 중영문 쿼리, Hacker News 및 arXiv 모두 모니터링 기간 내 실질적 적중 없음(10-29)
6개 쿼리는 정확한 단어+2일, 중국어+Ascend, 중국어+연산자, 조직명+2일, 영어+GPU, 중국어+컴파일러+7일을 포함한다. 중국어+Ascend만 한 건을 반환했으며, 이는 임상 인공지능 진전의 전재 기사로 연산자 언어와 무관하고, 나머지는 모두 빈 반환이다. Hacker News 최근 4일의 관련 적중은 모두 모니터링 기간 밖에 있다(가장 최근은 10-02의 Ascend 프로그래밍 스택 전재 토론). arXiv는 TileLang으로 검색 시 새 논문이 없다. 뉴스 측은 연속 여러 날 적중이 없으며, 데일리 리포트 본문은 계속 저장소와 생태계 저장소를 기준으로 한다.
5. 트렌드 관찰
Ascend 측의 리듬은 이미 「백엔드 이식」에서 「연산자 보충」으로 넘어갔다. 어제 기록한 것은 NPUIR 백엔드가 메인 브랜치에 병합된 것과, 검수 기준이 504개 자체 테스트 및 71 대 72의 예제였다. 오늘 병합된 다섯 건 중 네 건은 연산자와 예제이며, 업스트림 모델의 요구에 따라 배치되었다: GLM-5.3의 인덱서 압축, 블록 게이팅 어텐션의 내적, 4비트 양자화 행렬 곱, 양자화 비전 모델의 심층 합성곱. 이 연산자들의 공통점은 양자화와 융합을 호스트 측에 맡기지 않고 kernel 본체에 직접 기록했다는 것이다(E4M3 순수 정수 인코딩, 4비트 패킹 규약, 이중 kernel + 1승 1 패스트 패스). 동시에 두 건은 부수적으로 엔지니어링 기준을 세웠다: 연산자 성능 프로파일링을 프로세스에 포함시키고, 동시 배치로 인한 불안정 테스트 케이스를 목록화된 테스트 채널로 전환했다. 이는 Ascend 측의 주요 모순이 「컴파일이 가능한가」에서 「연산자의 수치와 성능이 안정적으로 검수될 수 있는가」로 변하고 있음을 보여준다.
메인 저장소의 오늘 세 건(하드웨어 리덕션, 동시 컴파일, 경계 초과 복사 그룹화)은 각각 명령어 선택, 컴파일 지연, 발사 품질에 대응하며, 모두 컴파일러 커널 문제이지 언어 계층 기능이 아니다. 이는 어제의 신규 결함 방향과 일치한다: 경계 초과, 별칭, 비트 폭 같은 정적 검사의 경계가 하나씩 열리고 있다. 오늘 세 건의 신규 결함 중 두 건(패딩 버퍼가 있는 배치 복사, 뷰 비트 폭 오버플로)이 실제 모델 형상에서 트리거되었으며, 이는 TileLang의 사용 범위가 경계 조건에 부딪힐 정도의 규모로 깊어졌음을 의미한다: 248320 x 5120의 NVFP4 가중치, K가 16의 정수 배가 아닌 어텐션 헤드 차원.
생태계 측에서 오늘 가장 주목할 점은 「fork화」다. Ascend 커뮤니티에 독립 연산자 저장소가 등장했고, T-HEAD 명의로 PPU 브랜치가 있는 TileLang 브랜치 저장소가 있으며, PTO-ISA는 문제 은행과 벤치마크를 모두配套 저장소로 만들었고, 별도의 개발 도구 저장소는 다운스트림 추적을 독립 패키지로 배포했다. 이러한 움직임의 공통 전제는 TileLang을 자사 하드웨어 스택의 기반으로 취급한다는 것이다: 자사 명령어 집합에 맞춰 레이아웃을 변경하고, 자사 테스트 체계로 검수하며, 자사 지식 형태로 경험을 축적해야 한다. 도메인 특화 언어에게 이는 스타 수보다 더 강한 채택 신호이지만, 업스트림이 「브랜치와 업스트림의 동기화」라는 장기 비용에 직면해야 함을 의미하기도 한다: 오늘 #3705 결함은 브랜치 사용자가 실제 빌드에서 먼저 부딪힌 것이다.
부록, 자료 및 검증 설명
출처 검증표
| 출처 | 검증 결과 |
|---|---|
| GitHub tile-ai 조직 저장소별(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 문서 사이트 / 야간 미러 저장소) | 모니터링 기간 내 푸시 있음: tilelang, TileOPs, TileFoundry, tilelang-ascend, tilelang-hygon, tilelang.github.io, TileOPs.github.io, TileOPs-nightly; tilelang-metax(10-26), tilelang-musa(09-30), TileRT, tilescale, DeepStack, tilelang-puzzles는 모니터링 기간 내 메인 브랜치 커밋 없음 |
| 병합 진실 검색(is:pr is:merged) | 12건: tilelang 5, tilelang-ascend 5, TileFoundry 1, TileOPs 1 |
| 심사 중 진실 검색(is:pr is:open updated) | 54건, Ascend 문서 검증 배치, NPUIR 세 건 뷰 수정 및 TileOPs의 GEMM 스케줄링 변경 포함 |
| 신규 issue 확인 | 4건: #3704, #3705, #3707 세 건 결함, 나머지 1건은 Ascend 저장소 일일 테스트 로봇 실패 알림 |
| 메인 브랜치 커밋 수준 검증 | tilelang 5건, tilelang-ascend 5건, TileOPs 1건, TileFoundry 1건, 문서 사이트 5건; tilelang과 두 저장소 병합 수 일치 |
| 릴리스 페이지 및 tag | tilelang 최신은 여전히 v0.1.15(09-30 릴리스), 모니터링 기간 내 신규 버전 없음; TileOPs 최신은 v0.1.0(10-27); TileFoundry 최신은 0.0.4rc1(10-25); tilelang-ascend 최신은 v0.1.2.000(09-09) |
| 서드파티 및 생태계 저장소 | 10개 저장소逐一 확인: Ascend 커뮤니티 연산자 저장소, PTO-ISA 이중 저장소, 개발 도구 저장소, 지식 저장소, 도해 강좌 저장소, 분산 레이아웃 저장소, DCU 융합 연산자 저장소, 모델 이식 저장소, Hexagon 백엔드 저장소, T-HEAD 브랜치 저장소 모두 모니터링 기간 내 푸시 있음(Hexagon 메인 브랜치는 신규 커밋 없음, 저장소 푸시 시간만 업데이트됨) |
| 채택자 | TileKernels와 FlashQLA 메인 브랜치는 10-14에 멈춰, 모니터링 기간 내 푸시 없음 |
| Google News(중영 6개 그룹) | 중국어 + Ascend 임상 인공지능 전재 기사 1건만 있으며 주제와 무관하여 수록하지 않음 |
| Hacker News | 최근 4일 관련 항목 모두 모니터링 기간 밖 |
| arXiv | TileLang 검색 결과 신규 논문 없음 |
전체 출처 목록
- [1] tilelang 전체 warp 리덕션에 redux.sync 사용(#3692) — https://github.com/tile-ai/tilelang/pull/3692
- [2] tilelang autotuner 독립 설정 동시 컴파일(#3476) — https://github.com/tile-ai/tilelang/pull/3476
- [3] tilelang Ascend 버퍼 뷰 레이아웃과 포인터 단위(#3623) — https://github.com/tile-ai/tilelang/pull/3623
- [4] tilelang Ascend 범위 초과 복사와 채우기 그룹화(#3543) — https://github.com/tile-ai/tilelang/pull/3543
- [5] tilelang CI PyTorch wheel 고정(#3706) — https://github.com/tile-ai/tilelang/pull/3706
- [6] tilelang 채우기가 있는 공유 버퍼 배치 복사 결함(#3704) — https://github.com/tile-ai/tilelang/issues/3704
- [7] tilelang 음수 스케일 반복 그래프 역방향 인덱스 결함(#3705) — https://github.com/tile-ai/tilelang/issues/3705
- [8] tilelang 뷰 비트 폭 int32 오버플로 결함(#3707) — https://github.com/tile-ai/tilelang/issues/3707
- [9] tilelang 비트 폭 int64 변경 수정(#3708) — https://github.com/tile-ai/tilelang/pull/3708
- [10] tilelang CPU 스칼라 GEMM 대상 영역만 초기화(#3709) — https://github.com/tile-ai/tilelang/pull/3709
- [11] tilelang 캐시 비활성화 시 프로세스 내 캐시 우회(#3710) — https://github.com/tile-ai/tilelang/pull/3710
- [12] tilelang-ascend GLM-5.3 kpool 압축 연산자(#1910) — https://github.com/tile-ai/tilelang-ascend/pull/1910
- [13] tilelang-ascend Chunk Scaled Dot KKT 연산자(#1803) — https://github.com/tile-ai/tilelang-ascend/pull/1803
- [14] tilelang-ascend W4A4 int4 GEMM 수정(#1731) — https://github.com/tile-ai/tilelang-ascend/pull/1731
- [15] tilelang-ascend depthwise_conv_2d 예제(#1920) — https://github.com/tile-ai/tilelang-ascend/pull/1920
- [16] tilelang-ascend 연산자 성능 프로파일링 skill(#1931) — https://github.com/tile-ai/tilelang-ascend/pull/1931
- [17] tilelang-ascend NPUIR 로컬 UB 뷰 디스크립터(#1962) — https://github.com/tile-ai/tilelang-ascend/pull/1962
- [18] tilelang-ascend NPUIR 뷰 소스 오프셋(#1960) — https://github.com/tile-ai/tilelang-ascend/pull/1960
- [19] tilelang-ascend NPUIR 단항 브로드캐스트 추론(#1961) — https://github.com/tile-ai/tilelang-ascend/pull/1961
- [20] tilelang-ascend PTO 비트 단위 XOR 워크스페이스 뷰(#1959) — https://github.com/tile-ai/tilelang-ascend/pull/1959
- [21] tilelang-hygon HCU 일반 복사 LDS swizzle(#30) — https://github.com/tile-ai/tilelang-hygon/pull/30
- [22] TileLang DCU 융합 연산자 저장소 — https://github.com/SealSoul117/TileLang_DCU_FuseCmpComm
-
[23] TileOPs 릴리스 이미지(#2793) — <ht
- [24] TileOPs GEMM 그리드 그룹화 및 상주 블록 수(#2794) — https://github.com/tile-ai/TileOPs/pull/2794
- [25] TileFoundry AMX 정적 TIR 계약(#279) — https://github.com/tile-ai/TileFoundry/pull/279
- [26] Ascend 커뮤니티 연산자 저장소 ops-tilelang — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [27] PTO-ISA tilelang-puzzles-ascend — https://github.com/PTO-ISA/tilelang-puzzles-ascend
- [28] PTO-ISA tilelangGym — https://github.com/PTO-ISA/tilelangGym
- [29] T-HEAD 소속 tilelang 브랜치 저장소 — https://github.com/t-head/tilelang-for-sail
- [30] tilelang 문서 사이트 — https://github.com/tile-ai/tilelang.github.io
- [31] 개발 도구 저장소 tilelang-devkit — https://github.com/easy-tilelang/tilelang-devkit
- [32] 지식 저장소 tilelang-knowledge — https://github.com/qiuuuuu622/tilelang-knowledge
- [33] 도해 강좌 부속 저장소 tilelang-course — https://github.com/zachshi-ai/tilelang-course
- [34] 분산 레이아웃 실험 저장소 — https://github.com/Leewin239/tilelang-distributed-layout
- [35] Hexagon 백엔드 저장소 tilelang-hexagon — https://github.com/xwhzz/tilelang-hexagon
- [36] DeepSeek 연산자 저장소 TileKernels — https://github.com/deepseek-ai/TileKernels
- [37] Qwen FlashQLA — https://github.com/QwenLM/FlashQLA
- [38] tilelang 릴리스 페이지 — https://github.com/tile-ai/tilelang/releases
- [39] TileOPs 릴리스 페이지 — https://github.com/tile-ai/TileOPs/releases
- [40] TileFoundry 릴리스 페이지 — https://github.com/tile-ai/TileFoundry/releases
- [41] Google News 검색 — https://news.google.com/