TileLang 데일리 리포트 (2026-11-03)
모니터링 기간: 2026-11-02 07:00 ~ 2026-11-03 07:00 (베이징 시간, 24시간; 이전 기간 기준점은 11-02 07:00).
출처: GitHub tile-ai 조직 저장소별 점검(29개 저장소, 기간 내 푸시 9개) + 머지 및 심사 중 실값 검색(머지 18건, 심사 중 업데이트 44건) + 신규 issue 전수 확인(8건) + 메인 브랜치 커밋별 점검(tilelang 5, TileOPs 7, TileFoundry 1, tilelang-ascend 4, tilelang-hygon 1, tilelang-lsp 1, 문서 사이트 5) + 릴리스 페이지 및 tag 점검(5개 저장소) + 채택 측 2개 저장소 + Google News 중영문 검색(프록시 경유) + Hacker News + arXiv.
이번 호 인덱스
- 오늘의 핵심: Ascend 950의 영속화 FlashMLA 커널 착지(#3678, 1997줄 추가, 42회 커밋); 벡터화 계획의 실제 회귀 수정 1건(#3734); #line 지시자를 라인별 재앵커링으로 변경(#3703); TileOPs 하루 7건의 성능 및 테스트 인프라 병합; TileFoundry가 공유 메모리 정렬을 증명 가능한 분석 계층으로 구현(#283), 그리고 main의 lowering이 sm_89에서 7개 프로그램 중 4개 오류로 측정됨; 24시간 내 18건 병합, 심사 중 업데이트 44건, 신규 이슈 8건
-
- 핵심 프로젝트 진행 상황
- 1.1 메인라인: Ascend 950의 영속화 FlashMLA 커널, BF16 prefill과 컴팩트 FP8 decode 이중 경로(#3678)(11-02)
- 1.2 메인라인: cast 디커플링이 건너뛰어진 후 잔류한 512비트 벡터, SM100에서 ViT 전방 컴파일 실패(#3734)(11-02)
- 1.3 메인라인: #line 지시자를 라인별 재앵커링으로 변경, 디버거 매핑에 밀집 모드 선택 가능(#3703)(11-02)
- 1.4 메인라인: 어텐션 예제가 디바이스 공유 메모리 상한에 따라 tile을 추론, RTX PRO 6000에서 2통과 7실패에서 9전통과로 변경(#3702)(11-02)
- 1.5 메인라인: CUDA float32x8 비트 변환을 내장 bit cast로 변경, SASS 명령어 19개에서 18개로 감소(#3748)(11-02)
- 핵심 프로젝트 진행 상황
-
- 멀티 백엔드 적응
- 2.1 Ascend: WY Fast 전방 연산자 신규 추가, wy_fast_bwd_split 역방향 분할 이관(#1827 / #1828)(11-02)
- 2.2 Ascend: T.tile.fill과 T.tile.clear 문서 사실 검증 병합(#1634 / #1628)(11-02)
- 2.3 Ascend 심사 중: cython wrapper 캐시 심볼 형상 해석, 호스트 오버헤드 라인 이어받기(#1967)(11-02)
- 2.4 Ascend 심사 중: z3 버전 동기화를 메인라인에서 이관, 전체 저장소 메인라인 동기화 추진(#1965 / #1975)(11-02)
- 2.5 Ascend 심사 중: 리덕션 연산자 결함 2건, 첫 번째 Parallel 축 브로드캐스트를 스칼라 읽기로 변경(#1978 / #1968 / #1979)(11-02)
- 2.6 Hygon: gfx948을 gfx946 호환 처리, 우선 기존 명령어 테스트 가능하게 함(#32)(11-02)
- 2.7 Moore Threads, MetaX, mlir-ascend 모니터링 기간 내 푸시 없음(11-03)
- 멀티 백엔드 적응
-
- 생태계 및 채택사
- 3.1 TileOPs: 하루 7건 병합, 6건 성능, 1건 테스트 인프라(11-02)
- 3.2 TileOPs: warp 전용 decode의 softmax와 QK/PV 오버랩, 478회 스캔 호출 기하평균 0.952배(#2817)(11-02)
- 3.3 TileOPs: MoE 토큰별 수집 프리뮤테이션 라인, H200에서 387.8 TFLOPS(#2815)(11-02)
- 3.4 TileOPs: 16비트 원소별 범위 검사 포함 나눗셈 제거, 8바이트 오퍼랜드 스레드당 1벡터(#2814 / #2813)(11-02)
- 3.5 TileOPs: Norm 통계량을 행 내 원소 감산으로 변경, 1e36 근처 오버플로 제거(#2811)(11-02)
- 3.6 TileOPs: tileops_test 타깃으로 tests/ops가 CUDA 없이 실행(#2816)(11-02)
- 3.7 TileFoundry: 공유 메모리 정렬을 증명 가능한 분석 계층으로 구현, 83개 픽스처 계획 중 52승인 31거부(#283)(11-02)
- 3.8 TileFoundry 심사 중: CUDA lowering이 동기화 계획 발행 및 역검증, main에서 7개 프로그램 중 4개 오류(#285)(11-02)
- 3.9 채택사: FlashQLA, TileKernels 모니터링 기간 내 푸시 없음(11-03)
- 생태계 및 채택사
-
- 커뮤니티, 튜토리얼 및 이벤트
- 4.1 Ascend 동기화 정확성이 집중 버그 리포트 면으로 부상: 하루 8건 신규 이슈, 두 보고자가 각각 체계적으로 제출(11-02)
- 4.2 행 간 및 전치 뷰의 SIMD 로드가 DMA 의존성 상실, 근본 원인은 논리 영역 재구성(#3749)(11-02)
- 4.3 뉴스 측: Google News 중영문, Hacker News, arXiv 세 채널 모두 모니터링 기간 내 적중 없음(11-03)
- 커뮤니티, 튜토리얼 및 이벤트
-
- 트렌드 관찰
- 부록. 자료 및 검증 설명
오늘의 핵심
이날의 핵심 흐름은 단 하나였다. TileLang가 DeepSeek의 MLA 디코드 커널을 Ascend 950으로 옮겼고, 그것도 컴파일러가 스케줄링할 수 있는 방식으로 옮겼다. #3678은 한 번에 1997줄, 20개 파일, 42개 커밋을 추가했으며, BF16 prefill과 컴팩트 FP8 decode 두 경로를 모두 포괄하고, 선택적 FP8 및 FP4 추가 KV 캐시, attention sink, 불법 인덱스와 query별 독립 길이까지 포함한다. 구현 방식이 주목할 만하다. 어텐션 루프, 온라인 softmax, 역양자화와 스케줄링이 전부 TileLang로 작성되었고, T.Stage를 통해 컴파일러가 버퍼 버전과 동기화를 관리하게 하며, QK와 PV는 동일한 L0B 저장 링의 서로 다른 두 layout 뷰를 사용하고, L1 배치는 T.annotate_buffer_addresses로 직접 지정한다. prefill은 하나의 누산기 할당 안에서 네 개의 서로소 출력 슬라이스를 나누고, 정규화가 한 조각 끝나면 한 조각을 비운다. 이 작성 방식을 성립시키기 위해 컴파일러 측에서 동시에 세 가지 기능을 보강했다. L0C에서 UB로의 복사가 NZ 출력을 유지하고, 각 AIV 스트림이 자체 파이프라인 이전에 선택되며, 명시적으로 스케줄링된 단계가 물리적 발사 위상이 같을 때 소스 순서를 유지한다. 목표 형상은 매우 좁게 작성되었다. 헤드 수 64, Dqk와 Dv 모두 512, topk는 64의 배수다.
두 번째 흐름은 실제 컴파일러 회귀를 수정하는 것이다. #3734는 SM100의 ViT 순전파를 처리한다. Cannot convert type float32x16 to CUDA type. 경로는 길지만 각 단계는 명확하다. 벡터화 플래너가 256비트 BF16 저장을 위해 16개 lane을 선택하고 cast 제약을 지연시켜, DecoupleTypeCast가 계산과 메모리 접근을 분리하기를 기대했다. 그런데 #3579가 방금 “메모리 읽기 가드”를 디커플링에서 제외했고(이 가드를 반복하면 시맨틱이 바뀐다), 결국 최종 합법화 단계에서 여전히 지연 전략에 따라 폭을 선택해, 잔류 FP32 표현식이 512비트 벡터로 부풀려졌고 CUDA 백엔드가 이를 표현할 수 없었다. 수정 방식은 지연 제약을 IR 속성이 아닌 플래너의 파라미터로 만들고, 마지막 합법화 단계에서만 false를 전달해 가드의 FP32에서 BF16으로의 변환이 8개 lane으로 줄고, 성공적으로 디커플링된 BF16 복사는 여전히 16개 lane을 유지한다. 주소 cast도 데이터 벡터 타입 제약에서 분리해 기존 접근 분석으로 합법성을 판정한다.
세 번째 흐름은 디버깅 경험이다. #3703은 #line 지시어를 “각 TIR 문장의 첫 줄만 앵커링”에서 선택적 계층적 발사로 변경한다. 레벨 0은 비활성, 1은 원상 유지, 2 이상은 각 줄에서 생성된 비어 있지 않은 줄에 대해 재앵커링하며, 여러 줄로 펼쳐진 문장과 마무리 중괄호를 포함한다. 구현은 CUDA, HIP, CPU C, C host, Metal, Ascend가 공유하는 C 계열 출력 스트림에 위치하며, vendored된 TVM은 건드리지 않는다. 소스 스팬이 없는 문장과 함수는 명시적으로 <tilelang-generated>로 전환하고, 사용자 원문은 매크로 연속 줄과 문자열 리터럴이 깨지지 않도록 그대로 유지한다. 새 환경 변수 TILELANG_EMIT_LINE_DIRECTIVES를 Python 측 기본값으로 추가하고, 캐시 조회 이전에 pass 설정에 접어 넣는다. 로컬에서 42개 통과, 6개 건너뜀, CUDA와 HIP 사례는 작성자 머신에서 건너뛰었고, cuda-gdb의 실제 동작은 아직 직접 검증되지 않았다. 작성자 스스로 이 점을 명시했다. 같은 날 LSP 저장소에서도 한 건이 따라붙어 CUDA 디버깅을 새로운 줄 매핑에 정렬했다.
연산자 측의 리듬은 멈추지 않았지만 방향이 바뀌었다. TileOPs는 하루 7건 머지 중 6건이 성능, 1건이 테스트 인프라였고, 전부 동일 작성자로부터 나왔다. 각 건은 여전히 H200 실측 표와 체크리스트 회귀 기준을 첨부한다. 그중 가장 가치 있는 것은 #2816이다. 테스트 시에만 존재하는 tileops_test 타깃을 새로 추가해, PrivateUse1에서 각 연산자에 대한 공유 workload 참조를 계산하여 tests/ops의 스모크 테스트가 더 이상 CUDA 위에 서지 않아도 되게 한다. 1281줄, 59개 파일을 변경했고, 작성자는 경계도 명확히 밝혔다. 공유 참조와의 비교는 백엔드 무관 계약과 이식성만 증명할 수 있을 뿐, 실제 백엔드의 수치 정확성을 증명하지는 못하며, 그 라인은 여전히 #2611에 걸려 있다.
마지막은 TileFoundry의 숨은 흐름이다. #283은 “공유 메모리 접근을 어떻게 정렬해야 하는가”를 분석 계층으로 만든다. 동기화 패밀리가 lowering의 발사 순서에 따라 각 함수를 순회하고, 각 공유 메모리 접근에 대해 명령, 발사 mesh, 피연산자와 바이트 범위를 취하며, 충돌 접근 전에 commit/wait와 스코프 내 배리어를 계획한다. 83개 스케줄링 픽스처 중 52개를 계획하고 31개를 거부한다(24개는 TMA 복사, 6개는 wgmma 사용, 둘 다 완료 시맨틱이 아직 선언되지 않음, 1개는 스레드 경계가 다른 루프에 배리어를 삽입해야 함). 짝을 이루는 #285는 심사 중으로, 계획을 CUDA lowering으로 바꾸고 역방향 검증한다. sm_89 단일 카드에서 빌드 가능한 7개 프로그램 중 6개가 참조와 완전히 일치하고 synccheck와 memcheck가 모두 깨끗한 반면, main 방식으로 lower하면 같은 7개 중 4개가 결과가 틀린다. #283 설명에서 제시된 직접적 증상은 sm_89의 E4M3 GEMM 최대 편차 659다.
1. 핵심 프로젝트 진행 상황
1.1 메인라인: Ascend 950의 지속형 FlashMLA 커널, BF16 prefill과 컴팩트 FP8 decode 이중 경로 (#3678)
11-02 09:26에 병합되었으며, 1997줄 추가, 17줄 삭제, 20개 파일, 42회 커밋으로 이번 모니터링 기간 중 가장 큰 변경입니다. 커널은 Ascend 950을 대상으로 하며, prefill은 BF16, decode는 컴팩트 FP8을 사용하고, 선택적으로 FP8 및 FP4 추가 KV 캐시를 지원하며, attention sink, 불법 인덱스, 쿼리별 독립 길이를 지원합니다. 어텐션 루프, 온라인 softmax, 역양자화, 스케줄링이 모두 TileLang으로 표현되며, T.Stage와 컴파일러가 관리하는 버퍼 버전 및 동기화로 구동됩니다. QK와 PV는 L0B 저장 링의 서로 다른 두 layout 뷰를 공유하고, L1 배치는 T.annotate_buffer_addresses로 지정됩니다. prefill은 기존 누산기 할당 내에서 네 개의 서로소 출력 슬라이스를 분리하고 독립 동기화 이벤트를 배치하여, 정규화가 한 슬라이스 완료되면 즉시 하나씩 배출합니다. decode는 전체 행 출력 복사를 사용합니다. 컴파일러 측에서는 세 가지 지원이 보강됩니다. L0C에서 UB로의 복사는 NZ 출력을 유지하고, 각 AIV 스트림은 파이프라인 이전에 선택되며, 명시적 스케줄 단계가 물리적 발행 위상에서 동률일 때 소스 순서를 유지합니다. 목표는 헤드 수 64, Dqk와 Dv 모두 512, topk가 64의 배수이며, pair gather는 이미 기록된 제로 burst 코드 생성 우회 방식을 따릅니다. 정확성 및 벤치마크 러너, 버퍼 재사용 및 쿼리 경계 회귀, FlashMLA MIT 고지가 포함된 README가 함께 제공됩니다.
출처: https://github.com/tile-ai/tilelang/pull/3678
1.2 메인라인: cast 디커플링이 건너뛴 후 잔류하는 512비트 벡터, SM100에서 ViT 전방 컴파일 실패 (#3734)
11-02 09:05에 병합되었습니다. e8225d7f5에서 SM100 대상 ViT 전방은 가드된 FP32→BF16 출력 저장 지점에서 Cannot convert type float32x16 to CUDA type를 보고합니다. 플래너가 256비트 BF16 저장을 위해 16개 lane을 선택하고 cast 제약을 지연시키며, DecoupleTypeCast가 계산과 메모리 접근을 분리하기를 기대합니다. #3579가 메모리를 읽는 가드들을 디커플링에서 건너뛰게 한 이후, 후속 LegalizeVectorizedLoop가 동일한 전략으로 계속 지연시키면서 잔류 FP32 표현식이 512비트 벡터가 되어 CUDA 측에서 표현할 수 없게 되었습니다. 수정 방법은 defer_cast_constraints를 플래너의 C++ 파라미터(IR 속성이 아님)로 만들어 호출자가 단계별로 선택하도록 하는 것입니다. LayoutInference와 LowerTileOp는 true를 유지하고, 최종 합법화는 일괄 false를 전달하며, 계산된 폭을 명시적 vectorize_hint로 전달하여 재작성기가 기본 전략으로 재계획하지 않도록 합니다. 가드된 FP32→BF16은 8개 lane으로 줄어들고, 성공적으로 디커플링된 BF16 복사는 여전히 16개 lane을 유지하며, 가드 안전 검사는 유지됩니다. 주소 cast와 접근 포인터 파라미터는 데이터 벡터 타입 제약에서 제외되고, 연속 주소는 스칼라 기저 주소 더하기 스트라이드로 낮아지며, 기존 접근 분석이 합법성을 판정합니다. int64 인덱스의 정렬 벡터 접근과 4-lane 원자 사례는 영향을 받지 않습니다.
출처: https://github.com/tile-ai/tilelang/pull/3734
1.3 메인라인: #line 지시자가 라인별 재앵커링으로 변경, 디버거 매핑에 선택적 밀집 모드 추가 (#3703)
11-02 04:06에 병합되었으며, 682줄 추가, 11개 파일입니다. 기존 #line 발행은 하나의 TIR 문장이 생성하는 첫 줄만 앵커링하고, 후속 줄은 컴파일러 논리 줄 번호로 진행되어 cuda-gdb가 생성된 로드, 연산, 저장을 무관한 Python 줄에 귀속시키는 문제가 있었습니다. 합성 문장과 종료 중괄호는 진행된 매핑을 상속할 수 있고, 소스 스팬이 없는 함수는 이전 함수의 소스 파일을 상속할 수 있었습니다. 이번 변경은 tl.emit_line_directives를 계층화합니다. 0 또는 False는 끄고 기본값을 유지하며, 1 또는 True는 문장 수준을 유지하고, 2 이상은 비어 있지 않은 각 생성 줄에 대해 재앵커링합니다(여러 줄로 확장된 문장과 종료 중괄호 포함). 구현은 C 계열 출력 스트림에 위치하며 CUDA, HIP, CPU C, C host, Metal, Ascend를 포함하고, vendored TVM은 변경하지 않았습니다. 동시에 문장과 함수의 소스 스팬을 기록하고, 소스가 없는 영역은 명시적으로 <tilelang-generated>로 전환하며, 사용자 원문은 매크로 연속 줄과 문자열 리터럴 손상을 피하기 위해 그대로 보존합니다. Python 측 기본값으로 TILELANG_EMIT_LINE_DIRECTIVES를 추가하고, 캐시 조회 전에 pass 구성에 반영하며, 명시적 구성이 우선합니다(기본값이 불법일 때 포함). 로컬에서 Metal과 LLVM 빌드로 검증했습니다. 행 지시자와 compile-only 두 그룹 42개 통과, 6개 건너뜀. 지시자를 제거한 후 벡터화 덧셈과 조건 선택 두 픽스처는 레벨 0, 1, 2에서 산출물이 완전히 일치했습니다. 작성자는 CUDA와 HIP 사례가 로컬에서 건너뛰어졌으며 cuda-gdb 동작은 직접 테스트하지 않았다고 밝혔습니다.
출처: https://github.com/tile-ai/tilelang/pull/3703
1.4 메인라인: 어텐션 예제가 디바이스 공유 메모리 상한에 따라 tile 산출, RTX PRO 6000에서 2 통과 7 실패에서 9 전부 통과로 (#3702)
11-02 06:59에 병합되었으며, 345줄 추가, 13개 파일입니다. 여러 어텐션 예제의 비튜닝 경로가 Hopper 크기의 tile을 하드코딩하고 있어, sm_86, sm_89, sm_120에서 블록당 동적 공유 메모리 상한이 약 99 KiB인 환경에서 실행이 바로 실패했습니다. 오류 원문은 허용되는 동적 공유 메모리를 163840바이트로 설정하라는 요구였으나, 디바이스 0은 블록당 최대 101376바이트입니다. 이제 각 예제는 shared_memory_per_block_optin에서 tile을 추론하며, 방식은 기존 GQA decode 예제(example_gqa_decode.py)를 따릅니다. 데이터센터 구성은 변경되지 않았는데, 각 분기가 이전 단계 tile 요구량보다 낮을 때만 트리거되기 때문입니다. 약 99 KiB 디바이스에서의 적용 지점: attention sink의 네 예제는 각각 64×64 2단계 128스레드(순방향)와 64×32 1단계 128스레드(역방향)이며, example_gqa_bwd.py는 순방향 64×64, 역방향 64×32, example_mla_decode.py는 BLOCK_N 32, BLOCK_H 16을 사용합니다. 후자 두 개에는 threads 파라미터가 추가되었습니다. 64행 또는 16행 tile에서 T.gemm은 각 warp가 16행 전체를 독점하려면 4개의 warp가 필요하며, 256스레드는 lower 단계에서 바로 실패합니다(각각 acc_s와 acc_s_cast의 layout 충돌, 그리고 M=16 N=32에 유효한 warp 분할 없음 오류). 또한 flash-attn 2.x에는 sm_12x 커널이 없어 네 개의 varlen 예제가 원래 참조 단계에서 오류를 발생시켰는데, 이제 flash_attn_varlen_supported를 한 번 탐지하여 지원되지 않으면 순수 PyTorch의 varlen_attention_ref로 전환합니다(GQA, 우하단 정렬 causal 및 window_size 지원). 검증은 RTX PRO 6000 Blackwell(sm_120)에서 수행되었습니다: CUDA 9개 항목 전부 통과(기존 2개 통과 7개 실패), CuTeDSL도 동일하게 9개 항목 전부 통과. H100은 227 KiB 선택 상한이 있어 기존 tile을 그대로 사용하며, 역시 9개 항목이 통과했습니다.
출처: https://github.com/tile-ai/tilelang/pull/3702
1.5 메인라인: CUDA float32x8 비트 변환을 내장 bit cast로 변경, SASS 명령 19개에서 18개로 감소(#3748)
11-02 09:58에 병합되었으며, 20줄 추가, 3줄 삭제, #3679를 수정하고 #3681을 대체합니다. 변경은 CUDA에만 적용됩니다: float32x8의 비트 재해석을 비트 연산 시퀀스에서 __builtin_bit_cast(uint64_t, make_float2(v, v))로 교체하여 생성되는 SASS 명령 수가 19개에서 18개로 줄었습니다. 폭은 크지 않지만 컴파일러 백엔드에서 전형적인 “의미는 동일하고 산출물은 더 짧은” 유형의 마무리입니다.
출처: https://github.com/tile-ai/tilelang/pull/3748
2. 멀티 백엔드 적응(Ascend / Hygon / Moore Threads / MetaX)
2.1 Ascend: WY Fast 순방향 연산자 신규 추가, wy_fast_bwd_split 역방향 분할 이관(#1827 / #1828)
두 건 모두 11-02 02:33에 병합되었으며, 전후 차이는 1분 이내입니다. #1827은 “이관 및 신규 추가” 방식으로 WY Fast 연산자를 Ascend NPU 측으로 옮겼고, #1828은 wy_fast_bwd_split의 분할 역방향을 이관하여 보완했습니다. WY 표기법은 선형 어텐션 계열 모델에서 재귀 행렬 곱을 저랭크 형태로 단순화하는 데 흔히 쓰이는 수단으로, 순방향과 역방향 분할이 함께 갖춰짐으로써 이 연산자 체인이 Ascend에서 학습 가능한 수준으로 완성되었습니다. 두 PR 설명에는 템플릿 인사말만 있으며, 변경 범위는 제목과 커밋을 기준으로 하며, 본 보고서는 내부 구현 세부 사항을 추정하지 않습니다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1827, https://github.com/tile-ai/tilelang-ascend/pull/1828
2.2 Ascend: T.tile.fill 및 T.tile.clear 문서 사실 검증 병합(#1634 / #1628)
같은 배치로 병합된 것으로 T.tile.fill과 T.tile.clear의 API 동작에 대한 문서 측 사실 검증 두 건이 있습니다. #1628의 논의에는 유지보수 관례가 한 줄 남아 있습니다: 이 두 인터페이스는 서로 유사하므로, 한쪽을 수정할 때 다른 쪽도 대응 조정이 필요한지 함께 확인해 달라는 내용입니다. 이는 지속적으로 진행 중인 API 사실 검증 배치에 속하며, 같은 시기 심사 중인 항목으로는 T.tile.bitwise_not(#1656), T.tile.bitwise_and(#1651), T.tile.bitwise_lshift(#1659), T.tile.bitwise_or(#1655)가 있습니다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1634, https://github.com/tile-ai/tilelang-ascend/pull/1628
2.3 Ascend 심사 중: cython wrapper 캐시 심볼 형상 파싱, 호스트 오버헤드 라인 이어받기 (#1967)
11-02 업데이트. #1967은 cython wrapper에서 심볼 형상 파싱 결과를 캐시하여, 매 호출마다 발생하는 dispatch 오버헤드를 계속 낮추는 것을 목표로 하며, 결함 #1898에 대응한다. 이 라인과 지난 호에 머지된 #1948은 동일 목표의 두 가지 착지점이다: #1948은 매 호출 시 무조건 TVM Analyzer를 구성하던 고정 오버헤드를 지연 구성으로 바꾸었고(910B3에서 약 450마이크로초에서 약 106마이크로초로 감소), #1967은 심볼 경로에서 반복 파싱 자체의 비용을 처리한다. 같은 계열에는 심사 중인 #3757도 있으며, 메인라인 측에서는 동적 ABI와 공유 심볼 검증을 정렬한다: 출력 이름이 아닌 Var 객체를 키로 사용하여 동일 이름의 서로 다른 변수가 같은 심볼로 취급되는 것을 방지하고, 단일 차원의 stride는 폴백으로만 취급하여 값 결정을 앞서서 하지 않도록 한다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1967, https://github.com/tile-ai/tilelang/pull/3757
2.4 Ascend 심사 중: z3 버전 동기화를 메인라인에서 이관, 전체 저장소 메인라인 이관 동시 추진 (#1965 / #1975)
#1965는 11-02 12:18에 업데이트되어 메인라인의 z3 버전 insert sync를 Ascend 저장소로 이관했다. #1975는 같은 날 03:28에 업데이트되었으며, 여전히 TileLang-Ascend 전체를 메인라인 다중 백엔드 아키텍처로 이관하는 작업이다. 두 가지는 동일한 수렴 경로의 두 가지 측면이다: 하나는 동기화 알고리즘 구현을 교체하고, 다른 하나는 저장소의 골격을 교체한다. 이와 함께 메인라인 측 작업으로 #3752(z3 rlimit 설정)가 심사 중이며, #3735는 thread sync에서 모듈화된 인덱스의 비중복성을 증명하는 것을 제안한다—지난 호 리포트에서 ThreadSync 롤백이 남긴 증명 예산 문제를 ‘해석 자원 설정’과 ‘증명해야 할 형상 감소’ 두 방향에서 동시에 처리하고 있다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1965, https://github.com/tile-ai/tilelang-ascend/pull/1975
2.5 Ascend 심사 중: 리덕션 연산자 결함 두 건, 첫 Parallel 축 브로드캐스트를 스칼라 읽기로 변경 (#1978 / #1968 / #1979)
세 건 모두 11-02에 업데이트되었다. #1978은 AscendVidReduction이 행 분할 루프에서 전체 [1,N] UB에 대해 tile 연산자 크기를 잘못 절반으로 줄이는 문제를 수정하며(#1800 대응); #1968은 threads=2일 때 sigmoid와 round의 카운트를 절반으로 줄이며(#1897 대응); #1979는 첫 Parallel 축의 브로드캐스트 피연산자가 스칼라 읽기를 사용하도록 하여 VEC의 UB 정렬 함정을 우회한다(#1885 대응). 세 가지 모두 같은 종류의 문제에 속한다: Ascend 측 벡터화 경로는 형상과 정렬에 대한 가정이 CUDA 측보다 더 엄격하여, 경계 조건이 깨지면 컴파일 타임 또는 런타임 하드 오류가 발생한다.
출처: https://github.com/tile-ai/tilelang-ascend/pull/1978, https://github.com/tile-ai/tilelang-ascend/pull/1968
2.6 Hygon: gfx948을 gfx946 호환으로 처리, 우선 기존 명령 테스트 가능하게 (#32)
11-02 10:46에 머지됨. 이 변경은 gfx948을 gfx946 호환으로 처리하여 기존 MLS atoms, MMAC, barrier 및 컴파일 옵션을 재사용하며, 목적은 gather load를 연결하기 전에 기존 builtins가 새 모델에서 정상 동작하도록 하는 것이다. 이는 전형적인 2단계 접근이다: 먼저 새 아키텍처 번호를 기존 경로에 연결하여 테스트 가능한 베이스라인을 확보한 후, 실제로 새 명령이 필요한 부분을 보충한다. 해당 저장소의 이전 일반 복사 LDS swizzle 수정은 지난 모니터링 기간에 머지되었다.
출처: https://github.com/tile-ai/tilelang-hygon/pull/32
2.7 Moore Threads, MetaX, mlir-ascend 모니터링 기간 내 푸시 없음 (11-03)
tilelang-musa의 마지막 푸시는 09-30에, tilelang-metax는 10-26에, tilelang-mlir-ascend는 10-31 08:21(모니터링 기간 시작보다 이름)에 멈춰 있다. 따라서 이번 호의 백엔드 동향은 Ascend와 Hygon 양측에 집중되었으며, 그중 Ascend는 머지 18건 중 4건, 심사 중 업데이트 44건 중 대다수를 차지했다.
출처: https://github.com/tile-ai/tilelang-musa, https://github.com/tile-ai/tilelang-metax
3. 생태계 및 채택 측 (TileOPs / TileFoundry / TileKernels / FlashQLA)
3.1 TileOPs: 하루 7건 머지, 6건 성능, 1건 테스트 인프라 (11-02)
11-02 05:10부터 05:22 사이에 모두 7개의 커밋이 병합되었으며, 작성자는 모두 lcy-seso로 어텐션, MoE, 원소별 연산, 정규화, 테스트 목표, CI 계약의 6개 방향을 다룹니다. 이전 회차의 «하루 13개»와 비교하면 수량은 절반으로 줄었지만 구조가 바뀌었습니다. 이전 회차는 shape별 튜닝 마무리였고, 이번 회차에는 테스트 인프라(#2816)와 CI 계약(#2818) 두 가지 비성능 변경이 등장하여 연산자 라이브러리가 다중 백엔드와 유지보수성을 위한 구조적 준비를 시작했음을 보여줍니다. 각 커밋에는 여전히 H200 실측 표가 첨부되며, manifest에 이미 측정된 행의 회귀가 없음을 선언합니다.
출처: https://github.com/tile-ai/TileOPs/commits/main
3.2 TileOPs: warp 전용 decode의 softmax와 QK/PV 중첩, 478회 스캔 호출 기하평균 0.952배 (#2817)
11-02 05:21에 병합. 변경은 warp 전용 GQA decode의 softmax를 QK, PV 행렬 곱과 중첩시킵니다. 확률은 레지스터에 남고, K와 V는 독립 배리어를 사용하며 3슬롯 링 버퍼를 공유합니다. 선택 조건은 D가 128, 64개 key의 tile, KV 헤드당 한 행 tile, softcap 없음, 그리고 «CTA 하나당 SM 하나»의 그리드가 점유율과 스트림 길이 한계를 만족할 수 있는 경우이며, 그렇지 않은 호출은 계속 직렬 프로그램으로 처리됩니다. 9개 플래너 경계는 모두 호스트 측에서 고정 디바이스 사실로 검증되며, 실제 어텐션 텐서 할당이나 132 SM 카드가 필요하지 않습니다. 실측 환경은 H200, CUDA 13.2, PyTorch 2.13.0+cu132, TileLang 0.1.12이며, CUPTI 디바이스 사용 시간 기준으로 shape (16, 8, 4096, 64, 8, 128)에서 float16은 0.074610밀리초 대 fa3의 0.077490(1.039배, 3.654 TB/s), bfloat16은 0.074260 대 0.077220(1.040배, 3.671 TB/s)입니다. 영향을 받은 478회 스캔 호출의 기록된 기하평균 지연은 main의 0.952배입니다.
출처: https://github.com/tile-ai/TileOPs/pull/2817
3.3 TileOPs: MoE가 토큰별 수집으로 사전 치환, H200에서 387.8 TFLOPS (#2815)
11-02 05:19에 병합. tight 레이아웃의 MoE 입력이 토큰별 수집으로 변경되어 동일한 hidden 행이 여러 라우팅 간에 재사용됩니다. 토큰이 GPU를 채우지 못할 때는 라우팅을 서로 다른 블록으로 분할합니다. 동시에 tight 스캔에서 사용되지 않던 전문가 순서 소스 인덱스 할당과 쓰기를 삭제했습니다. 스케줄링에는 새 규칙이 추가되었습니다. 융합되지 않은 K 주순서 tight 그룹의 평균 행 수가 128보다 크고 192 이하이면 H200에서 ping-pong 64×256 tile을 선택하며, 그룹 크기를 명시한 호출은 여전히 비용 모델을 사용합니다. 실측 (4096, 256, 8, 7168, 2048)에서 bfloat16은 7.4417밀리초 대 flashinfer-cutlass의 7.6982(1.034배, 387.844 TFLOPS), float16은 7.5040 대 7.9869(1.064배, 384.624 TFLOPS)입니다. 기록된 라우팅 및 짝수 그룹 스캔은 두 dtype과 편향 모드를 모두 커버하며, manifest에 노이즈를 초과하는 회귀는 없습니다.
출처: https://github.com/tile-ai/TileOPs/pull/2815
3.4 TileOPs: 16비트 원소별 연산에서 범위 검사가 있는 나눗셈 제거, 8바이트 피연산자는 스레드당 한 벡터 (#2814 / #2813)
#2814는 11-02 05:18에 병합되었으며, bfloat16 Sigmoid와 16비트 SiLUAndMul이 4분의 1 스케일 역수를 사용하도록 변경하여 범위 검사가 있는 나눗셈을 우회하고, 해당 보조 함수를 SiLU와 공유합니다. 동시에 지수 연산에서 사용되지 않는 서브노멀 가드를 생략하고, 융합 게이트 출력은 스레드 간에 인접한 16바이트 벡터로 저장하며, gate가 -88에 가까울 때 무한값을 곱하는 미반올림 SiLU 활성화를 유지합니다. 실측 bfloat16 Sigmoid는 16777216개 원소에서 0.017504밀리초 대 torch-compile의 0.018656(1.066배, 3.834 TB/s), 268435456개 원소에서 0.248545 대 0.262529(1.056배, 4.320 TB/s)이며, SiLUAndMul은 (2048, 14336)에서 0.041856 대 0.045440(1.086배)입니다. #2813은 05:17에 병합되었으며, 8바이트 피연산자가 긴 런치에서 스레드당 한 벡터로 처리되도록 합니다. 변경 전 TileFoundry의 roofline 분석으로 이러한 비트 단위 OR 연산자가 모두 메모리 제약임을 확인했습니다. 4.8 TB/s의 HBM 대역폭 기준으로 각 workload의 하한을 제시하면, hidden-state-prefill의 int64는 41.9마이크로초, int32는 21.0마이크로초, llama-7b-ffn의 int64는 56.4, int32는 28.2, cnn-feat-broadcast의 int64는 42.8, linear-bias의 int64는 14.0, int32는 7.0입니다.
출처: https://github.com/tile-ai/TileOPs/pull/2814, https://github.com/tile-ai/TileOPs/pull/2813
3.5 TileOPs: Norm 통계량을 행 내 원소 기준으로 변경하여 1e36 근처의 오버플로 제거(#2811)
| 11-02 05:10에 병합됨. 근본 원인이 매우 직접적으로 기술되어 있다. 모든 평균 및 분산 커널이 원본 행에 대해 fp32 합산을 수행하는데, 합은 행의 오프셋에 따라 증가하며(대략 N 곱하기 | x | ), 1e36에 가까운 행은 직접 오버플로되어 모든 출력이 NaN이 된다. 오버플로되지 않더라도 평균은 오프셋 크기 수준에서 반올림되며, x − mean은 이 반올림을 각 원소에 전파하여 상수 행에서 계산된 분산이 반올림의 제곱과 같아져 bias가 아닌 플러스/마이너스 weight로 정규화된다. 수정 방법은 행으로 구성된 각 fp32 값에 불변식을 추가하는 것이다. 즉, 오프셋이 아닌 행의 분산도에 의해서만 제약받도록 하는 것이다. 구현상으로는 행 내의 특정 원소를 기준으로 삼아 다시 빼는 방식으로, warp 주도 커널은 warp의 첫 원소를, 샤딩 커널은 행의 첫 원소를 사용한다. 이는 이전 호 리포트의 #2812에 나열된 네 가지 Norm 연산자 결함과 동일한 유형의 문제에 대한 정면 수정이다. |
출처: https://github.com/tile-ai/TileOPs/pull/2811
3.6 TileOPs: tileops_test 타깃으로 tests/ops를 CUDA 없이 실행(#2816)
11-02 05:20에 병합됨. 이 모니터링 기간 내 단일 변경 중 가장 큰 규모로, 1281행 추가, 524행 삭제, 59개 파일이며 #2803을 닫는다. 테스트 시에만 존재하는 tileops_test 타깃을 새로 추가하여 PrivateUse1에서 각 연산자의 공유 workload 참조를 계산하고, tests/ops의 스모크 테스트가 이를 통해 실행된다. 지원되지 않는 참조 호출은 명시적으로 unvalidated로 나열되어 건너뛰며, CUDA에서만 성립하는 제외 항목도 동일한 타깃에서 검증된다. 연산자-워크로드 레지스트리는 workloads/calls를 통해 공유되며, 벤치마크 전용 바인딩, 복사 회계 및 입력 오버라는 benchmarks/_cases에 남는다. 또한 테스트 디바이스 생성기, 컨볼루션 지원을 보충하고 테스트 능력 표시를 수정하며, 참조 구현이 GQA 역방향의 scale과 빈 행 시맨틱을 준수하도록 한다. 작성자는 한계를 눈에 띄는 곳에 명시했다. 이 타깃이 테스트하는 것은 백엔드 독립적인 계약과 테스트 이식성이며, 공유 참조와의 비교만으로는 실제 백엔드의 수치적 정확성을 증명할 수 없다는 것이고, #2611은 여전히 열려 있다. 사용 가능한 참조가 없는 호출은 계속 unvalidated 형태로 가시화된다. 같은 배치의 #2818은 dispatch 계약을 workload 데이터 생성에서 분리하여, 계약 사례에 대해 형태만 인스턴스화하고 사용되지 않는 텐서 내용은 생성하거나 검증하지 않으며, 다섯 개의 호스트 구조 테스트 스위트를 compile-contract-gate에서 한 번만 실행하도록 한다.
출처: https://github.com/tile-ai/TileOPs/pull/2816
3.7 TileFoundry: 공유 메모리 정렬을 증명 가능한 분석 계층으로 구현, 83개 픽스처 중 52개 계획 31개 거부(#283)
11-02 02:09에 병합됨. 문제는 실질적이다. finalize가 공유 메모리 접근 사이에 commit, wait 또는 barrier를 발행하지 않아, 공유 메모리를 통해 임시 저장하는 스케줄러가 경쟁 상태에 빠지며, 카드에서 sm_89의 E4M3 GEMM이 torch 대비 최대 659만큼 편차가 발생한다. #267에서 정한 방향은 “HIR 상에서 접근 순서를 정렬하는 분석”과 “타깃에 따라 순서를 발행하는 lowering”이며, 최소 올바른 형태만 발행하고 증명할 수 없으면 거부하는 것이다. 이번 변경이 전달하는 것은 분석 부분이다. 새로 추가된 동기화 패밀리는 lowering의 발행 순서를 따라 각 함수를 순회한다. 각 공유 메모리 접근에 대해 analysis/emission.py에서 명령, 발행 mesh 및 피연산자를 가져오고, IterationScope에서 스코프를 가져오며, 메모리 패밀리에 기록된 배치에서 바이트 수를 가져온다. 충돌 접근 이전에 commit 한 번과 wait 한 번(미완료 cp.async 쓰기 대상) 및 “양쪽을 동시에 보유하는 가장 내부 스코프”를 커버하는 펜스를 계획하고, 먼저 실행되는 자식 노드의 활성 위치로 번호를 매긴다. 연산자는 새 레지스트리 visitor_registry/completion.py를 통해 효과가 어떻게 완료되는지 선언하며, 아무것도 선언하지 않은 것은 UNKNOWN으로 처리된다. 이 패밀리는 UNKNOWN 이후의 접근, 반환 시 아직 비행 중인 UNKNOWN 접근, 스레드 경계를 넘어 대기하지 않는 루프 내 펜스 삽입, 그리고 보존된 호출이며 피호출자가 공유 메모리를 보유하는 상황을 거부한다. 기본적으로 이 패밀리를 요청하는 호출자가 없으므로, finalize를 수행하는 75개 스케줄 픽스처의 바이트는 main과 완전히 동일하다. 83개 픽스처 중 52개 계획, 31개 거부(24개는 TMA 복사 사용, 6개는 wgmma 사용으로 완료 시맨틱 미선언, 1개는 스레드 관련 경계의 루프에 펜스 삽입 필요)이며, 코퍼스 사례는 이러한 거부를 항목별로 단언한다.
출처: https://github.com/tile-ai/TileFoundry/pull/283
3.8 TileFoundry 심사 중: CUDA lowering이 동기화 계획을 발행하고 역방향 검증, main에서 7개 프로그램 중 4개 오류(#285)
11-02 12:40 업데이트. 이 항목은 #283의 실행 측: CUDA lowering이 동기화 계획에 따라 단계를 실행하고, 실제 실행이 계획과 일치하는지 항목별로 검사하며, 불일치 시 LoweringError를 발생시킨다. 코퍼스상 결과: 기본 흐름은 14개 프로그램에 동기화를 더하고 31개를 거부; 이 14개의 TIR은 main과의 차이가 새로 추가된 구문뿐이며, 나머지 30개는 main과 완전히 동일; 핵심 visitor를 거친 후 main이 finalize할 75개 프로그램은 바이트 단위로 일치. 단일 sm_89에서 빌드 가능하고 동기화가 있는 7개 프로그램을 실행: 6개가 참조와 완전히 일치하고 synccheck와 memcheck 하에서 모두 깨끗; 동일한 7개를 main 방식으로 lower하면 4개가 잘못된 결과. 작성자는 동시에 main에 이미 존재하는 두 가지 문제를 보고: GEMM과 loop_outputs에서 racecheck가 동일한 cp.async의 여러 lane이 같은 주소에 쓰는 것을 보고 (HIR의 copy_async가 분할 대상을 받지 않아 각 레인이 전체 tile을 복사), loop_outputs는 또한 하나의 출력을 잘못된 요소에 쓴다.
출처: https://github.com/tile-ai/TileFoundry/pull/285
3.9 채택자: FlashQLA, TileKernels 모니터링 기간 내 푸시 없음 (11-03)
QwenLM/FlashQLA의 마지막 푸시는 10-30 03:02에 멈췄고, deepseek-ai/TileKernels는 10-14 05:49에 멈췄으며, 두 저장소 모두 이번 모니터링 기간에 커밋이 없다. TileKernels는 10-14 이후 여러 기간 연속으로 푸시가 없었고, 채택자의 가시적 활동은 이번 기간에 전부 자사 저장소인 TileOPs와 TileFoundry 두 곳에 집중되었다.
출처: https://github.com/QwenLM/FlashQLA, https://github.com/deepseek-ai/TileKernels
4. 커뮤니티, 튜토리얼 및 이벤트
4.1 Ascend 동기화 정확성이 집중 장애 보고 면으로 부상: 하루 8건의 신규 issue, 두 보고자가 각각 체계적으로 제출 (11-02)
모니터링 기간 내 메인 저장소에 8건의 issue가 추가되었으며, 번호는 #3740에서 #3749에 집중되어 있고 모두 Ascend AscendC와 PTO 두 백엔드의 코드 생성과 동기화 삽입을 지목한다. 그중 7건은 동일 보고자가 제출했으며, 환경은 TileLang 0.1.15, PyTorch 2.10.0+cpu, torch_npu 2.10.0이고, 주제는 동일한 프로그램이 두 백엔드에서 보이는 행동 분기: 생성 코드가 어떤 코어를 기동하는지 불일치 (AscendC는 mix(1,2), PTO는 vector), enable_auto_schedule=True일 때 AscendC가 동기화를 삽입하지 않음, 자동 스케줄링을 끄면 생성에 성공하지만 켜면 실패하고 오류 메시지가 불명확, 동일한 코드를 AscendC는 생성할 수 있지만 PTO는 실패. 이 보고 묶음의 공통점은 모두 재현 가능한 prim_func와 줄 단위 대조 생성 코드를 첨부했다는 점으로, 바로 회귀 케이스로 전환 가능한 형태이다. 여기에 심사 중인 #3735 (thread sync에서 모듈러 인덱스 비중복 증명)와 #3580 (술어와 승격 인덱스 동기화)을 더하면, Ascend 측 동기화 정확성은 현재 가장 집중된 한 축이 되었다.
출처: https://github.com/tile-ai/tilelang/issues/3747, https://github.com/tile-ai/tilelang/issues/3745
4.2 행 교차 및 전치 뷰의 SIMD 로드가 DMA 의존성을 잃음, 근본 원인은 논리 영역 재구성 (#3749)
같은 배치에서 기술적 깊이가 가장 높은 항목으로, 보고자 환경은 main 브랜치 e8225d7f5, Ascend950PR, CANN 9.2.0. 문제는 두 곳의 연속 SIMD 로드가 더 이른 GM에서 UB로의 복사에 대한 의존성을 잃는 것: 하나는 행 우선 로드가 행 중간에서 시작해 다음 행으로 넘어가는 경우, 다른 하나는 전치 뷰에 접근하며 물리적 연속 차원이 마지막 논리 차원이 아닌 경우. 근본 원인 경로는 매우 상세히 기술되어 있다: 생성된 tl.access_ptr이 물리적 로드 범위를 덮지만, MemoryAccessDetector::LogicalPointerRegion이 재구성한 논리 영역이 그중 일부를 누락하여, 의존성 분석이 입력 DMA와 로드가 무관하다고 판단하고, AutoSchedule이 입력 복사를 로드 이후로 옮기며, InsertSync도 MTE2에서 V로의 set/wait를 발행하지 않는다. 최소 재현에서 work.shape는 (2, 128)이고, 먼저 T.copy(A, work[1, 0:64])를 한 후, T.SimdVF() 안에서 T.simd.vld(work[0, 64])로 128개의 연속 반정밀도 요소를 읽는다: 생성된 읽기 서술자는 tl.access_ptr(work[0, 64], 128, 1)이고, 물리적으로 [64, 192)를 덮으며, 그중 [128, 192)가 바로 방금 DMA로 쓰인 부분인데, 재구성된 영역은 work[0:1, 0:128] 즉 물리적 [0, 128)로 쓰인 구간을 완전히 놓친다; 행 수는 ceil(128/128)로 계산되어 1행이 되고, 시작 열 64는 반영되지 않았다. 이 항목은 이번 모니터링 기간의 주력이 보강 중인 FlashMLA, PTO 벡터화와 같은 기반이며, 한 번 적중하면 조용히 잘못된 결과가 된다.
출처: https://github.com/tile-ai/tilelang/issues/3749
4.3 뉴스 측: Google News 중영문, Hacker News, arXiv 세 채널 모두 모니터링 기간 내 히트 없음 (11-03)
여섯 그룹의 Google News 쿼리(영문 TileLang, TileLang GPU kernel, tile-ai TileLang; 중문 TileLang 연산자, TileLang Ascend, MetaX 또는 Hygon 또는 Moore Threads)가 반환한 최신 관련 보도는 10-21의 SemiAnalysis 화웨이와 DeepSeek 기사이며, 나머지는 09-29부터 10-01까지의 DeepSeek 오픈소스 Ascend 컴포넌트 배치에 집중되어 있어 모두 본 모니터링 기간보다 이릅니다. arXiv에서 all:"TileLang"의 최신 히트는 2026-07-24의 TileSight이며, TileLang과 Ascend의 결합 쿼리는 결과가 없습니다. Hacker News의 후보 항목 중 TileLang과 관련된 것은 저장소 푸시 항목뿐입니다. 후보 목록 14건 중 12건은 GitHub 저장소 발견에서 나왔고, 2건은 무관한 Show HN과 LSP 저장소이며, 이번 호의 정보면은 전적으로 저장소 측으로 구성됩니다.
출처: https://news.google.com/, https://arxiv.org/
5. 트렌드 관찰
이날 가장 기록할 만한 판단은 Ascend가 더 이상 ‘적응 대상’이 아니라 ‘최우선 대상’이라는 것입니다. FlashMLA처럼 DeepSeek 자체 디코딩 경로의 핵심 커널이 Ascend 950을 직접 대상으로 메인라인에 작성되었으며, 사용된 것은 T.Stage, annotate_buffer_addresses 같은 명시적 스케줄링을 위한 저수준 기능이지 별도 저장소를 만들어 포팅한 것이 아닙니다. 컴파일러 측의 세 가지 지원(L0C에서 UB로 NZ 유지, AIV 스트림을 파이프라인 전에 선택, 명시적 단계를 소스 순서대로 발행)과 함께 보면, Ascend 경로가 ‘실행 가능’에서 ‘스케줄링 의도대로 실행 가능’으로 전환되고 있음을 알 수 있습니다.
두 번째 단서는 정확성의 중심이 하향 이동하고 있다는 점입니다. 지난 호 메인라인은 TIRX 마이그레이션을 위해 ThreadSync 정확성 개선 일체를 롤백했고, 이번 호에는 즉시 두 가지 보완이 등장했습니다. 메인라인 #3752는 z3로 리소스 상한을 조정하고, #3735는 모듈화 인덱스의 비교차성으로 증명해야 할 형상을 줄이는 방안을 제시합니다. 동시에 #3749는 ‘의존성 분석이 물리적 로드 범위를 보지 못하는’ 결함을 MemoryAccessDetector::LogicalPointerRegion 계층까지 파고들었습니다. 이를 TileFoundry라는 암선과 함께 보면——분석 계층은 83개 픽스처 중 52개의 공유 메모리 정렬을 증명할 수 있지만, lowering을 main 방식대로 하면 7개 프로그램 중 4개가 틀리게 계산됩니다——같은 신호가 세 번 나타납니다. TileLang의 현재 주요 리스크는 기능 부족이 아니라 ‘생성된 코드가 맞아 보이지만 실제로는 틀린’ 것이며, 팀이 선택한 방식은 정확성을 거부 가능한 증명으로 만드는 것이지 패치로 하나씩 막는 것이 아닙니다.
세 번째는 테스트와 계약이 독립적인 계층으로 분리되기 시작했다는 점입니다. TileOPs의 tileops_test는 스모크 테스트를 CUDA에서 떼어냈고, compile-contract-gate는 dispatch 계약과 workload 데이터 생성을 분리했으며, 둘 다 같은 시간 내에 머지되었습니다. 작성자는 동시에 ‘공유 참조는 실제 백엔드의 정확성을 증명할 수 없다’고 명확히 밝혔습니다. ‘무엇을 테스트할 수 있는가’와 ‘무엇이 증명되었는가’를 분리해 서술하는 이러한 습관은 성능 패치를 몇 건 더 머지하는 것보다 이 프로젝트가 프로토타입 단계를 벗어나고 있음을 더 잘 보여줍니다.
마지막으로 한 가지 대비에 주목할 필요가 있습니다. 연산자 측은 하루에 7건의 성능 머지가 있었고 그중 6건이 H200 실측 표를 동반했지만, 뉴스 측 세 채널은 24시간 동안 제로 히트였습니다. TileLang의 대외적 인지도는 여전히 ‘DeepSeek 오픈소스 Ascend 컴포넌트’ 같은 이벤트에 집중되어 있으며, 일상적 진전은 거의 저장소 안에서만 발생합니다. 이 프로젝트를 추적하는 사람에게 이는 데일리 리포트의 주체가 보도가 아니라 커밋과 PR일 수밖에 없음을 의미합니다.
부록: 자료와 검증 설명
출처 검증표
| 출처 | 검증 결과 |
|---|---|
| tile-ai 조직 (29개 저장소) | 모니터링 기간 내 9개 저장소 푸시 |
| 메인 저장소 tilelang | 메인 브랜치 5건 머지, 신규 이슈 8건, 문서 사이트 5건 자동 빌드 |
| TileOPs | 메인 브랜치 7건 머지, 사이트 1건 |
| TileFoundry | 메인 브랜치 1건 머지, 1건 심사 중 신규 |
| tilelang-ascend | 메인 브랜치 4건 머지, 6건 심사 중 업데이트 |
| tilelang-hygon | 1건 머지 |
| tilelang-lsp | 1건 머지 (CUDA 디버그 라인 매핑 정렬) |
| TileOPs-nightly | 1회 자동 푸시 |
| tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack | 모니터링 기간 내 푸시 없음 |
| 릴리스 페이지 (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) | 모니터링 기간 내 신규 버전 없음 |
| FlashQLA / TileKernels | 모니터링 기간 내 푸시 없음 |
| Google News (중영문 6개 쿼리) / Hacker News / arXiv | 모니터링 기간 내 실질적 히트 없음 |
전체 출처 목록
- [1] tilelang Ascend 950 FlashMLA 커널 (#3678) — https://github.com/tile-ai/tilelang/pull/3678
- [2] tilelang cast 분리 건너뛰기 후 벡터 폭 수정 (#3734) — https://github.com/tile-ai/tilelang/pull/3734
- [3] tilelang #line 라인별 재앵커링 (#3703) — https://github.com/tile-ai/tilelang/pull/3703
- [4] tilelang 어텐션 예제 공유 메모리 예산 (#3702) — https://github.com/tile-ai/tilelang/pull/3702
- [5] tilelang CUDA float32x8 비트 변환 (#3748) — https://github.com/tile-ai/tilelang/pull/3748
- [6] tilelang float32x8 비트 변환 결함 (#3679) — https://github.com/tile-ai/tilelang/issues/3679
- [7] tilelang 동적 ABI 및 공유 심볼 검증 (#3757) — https://github.com/tile-ai/tilelang/pull/3757
- [8] tilelang z3 rlimit 설정 (#3752) — https://github.com/tile-ai/tilelang/pull/3752
- [9] tilelang Ascend 모듈형 인덱스 비교차성 증명 (#3735) — https://github.com/tile-ai/tilelang/pull/3735
- [10] tilelang 행 간 및 전치 뷰 SIMD 로드 시 DMA 의존성 누락 (#3749) — https://github.com/tile-ai/tilelang/issues/3749
- [11] tilelang AscendC 및 PTO 백엔드 동기화 코드 오류 (#3747) — https://github.com/tile-ai/tilelang/issues/3747
- [12] tilelang 자동 스케줄링 시 동기화 미삽입 (#3745) — https://github.com/tile-ai/tilelang/issues/3745
- [13] tilelang 두 백엔드 시작 커널 불일치 (#3740) — https://github.com/tile-ai/tilelang/issues/3740
- [14] tilelang-lsp CUDA 디버그 라인 매핑 정렬 (71e762e0) — https://github.com/tile-ai/tilelang-lsp/commit/71e762e0
- [15] TileOPs warp 전용 decode softmax 오버랩 (#2817) — https://github.com/tile-ai/TileOPs/pull/2817
- [16] TileOPs MoE 토큰별 수집 (#2815) — https://github.com/tile-ai/TileOPs/pull/2815
- [17] TileOPs 16비트 원소별 범위 검사 제거 나눗셈 (#2814) — https://github.com/tile-ai/TileOPs/pull/2814
- [18] TileOPs 8바이트 피연산자 스레드당 벡터 하나 (#2813) — https://github.com/tile-ai/TileOPs/pull/2813
- [19] TileOPs Norm 통계량 행 내 원소 감소 (#2811) — https://github.com/tile-ai/TileOPs/pull/2811
- [20] TileOPs tileops_test 테스트 타깃 (#2816) — https://github.com/tile-ai/TileOPs/pull/2816
- [21] TileOPs dispatch 계약과 workload 분리 (#2818) — https://github.com/tile-ai/TileOPs/pull/2818
- [22] TileOPs Norm 결함 보고 (#2812) — https://github.com/tile-ai/TileOPs/issues/2812
- [23] TileFoundry 공유 메모리 정렬 분석 (#283) — https://github.com/tile-ai/TileFoundry/pull/283
-
[24] TileFoundry CUDA lowering 동기화 계획 방출 (#285) — <https://github.com/tile-ai/T
- [25] tilelang-ascend WY Fast 순전파 연산자(#1827) — https://github.com/tile-ai/tilelang-ascend/pull/1827
- [26] tilelang-ascend wy_fast_bwd_split 역전파(#1828) — https://github.com/tile-ai/tilelang-ascend/pull/1828
- [27] tilelang-ascend T.tile.fill 문서 검증(#1634) — https://github.com/tile-ai/tilelang-ascend/pull/1634
- [28] tilelang-ascend T.tile.clear 문서 검증(#1628) — https://github.com/tile-ai/tilelang-ascend/pull/1628
- [29] tilelang-ascend 기호 형상 파싱 캐시(#1967) — https://github.com/tile-ai/tilelang-ascend/pull/1967
- [30] tilelang-ascend 메인라인 z3 동기화 마이그레이션(#1965) — https://github.com/tile-ai/tilelang-ascend/pull/1965
- [31] tilelang-ascend 전체 메인라인으로 마이그레이션(#1975) — https://github.com/tile-ai/tilelang-ascend/pull/1975
- [32] tilelang-ascend 전체 블록 행 분할 tile 크기 절반 축소(#1978) — https://github.com/tile-ai/tilelang-ascend/pull/1978
- [33] tilelang-ascend threads=2 카운트 절반 축소(#1968) — https://github.com/tile-ai/tilelang-ascend/pull/1968
- [34] tilelang-ascend 첫 Parallel 축 브로드캐스트 스칼라 읽기(#1979) — https://github.com/tile-ai/tilelang-ascend/pull/1979
- [35] tilelang-hygon gfx948 호환(#32) — https://github.com/tile-ai/tilelang-hygon/pull/32