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

출처: GitHub tile-ai 조직 저장소별 점검(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 문서 사이트 / 언어 서비스 저장소 / 야간 벤치마크 저장소) + 머지 및 심사 중 실값 검색(머지 30건, 심사 중 업데이트 49건) + 신규 issue 전수 확인(11건) + 커밋 수준 diff 및 릴리스 페이지 점검 + PyPI 아티팩트 상태 점검 + 채택 측 저장소 점검(TileKernels / FlashQLA) + 커뮤니티 신규 저장소 검색 + Google News 중영문 다중 쿼리(프록시 경유) + Hacker News + arXiv + 산업 매체 보도.


이번 호 인덱스

  • 오늘의 핵심: 24시간 동안 30건 병합, 49건 심사 중 업데이트, 신규 이슈 11건; 메인 라인은 「스트림과 배리어 시맨틱」 관련 3건의 사일런트 오류에 집중; TileOPs 첫 릴리스는 여전히 PyPI 마지막 단계에서 정체; Hygon 2건 대형 병합
    1. 핵심 프로젝트 진행 상황
      • 1.1 메인 저장소: out_idx 출력이 호출자 스트림에서 할당되도록 변경, H100에서 재현된 5회 중 3회 손상이 0으로 감소 (#3649) (10-26)
      • 1.2 메인 저장소: NVRTC 디바이스 측 컴파일을 TVM-FFI 실행 경로에 연결, 컴파일 키 항목 추가 (#3285) (10-26)
      • 1.3 메인 저장소: 배치 TMA 로드의 cache hint에 템플릿 인자 하나 누락, DeepSelect의 TMA 경로 컴파일 실패 (#3683) (10-26)
      • 1.4 메인 저장소: 리덕션의 FP8/FP4 비유한 상수가 잘못된 리터럴로 출력, 서브바이트 누산기를 조기 거부로 변경 (#3189) (10-26)
      • 1.5 메인 저장소: SM120 블록 스케일링 GEMM에서 K 원자 명시적 전개, FullCol 구성 51% 향상 (#3286) (10-26)
      • 1.6 메인 저장소: JIT 출력 형상의 스칼라 파라미터를 할당 전에 해석 (#3602) (10-26)
      • 1.7 메인 저장소: sm8x 희소 메타데이터의 오퍼랜드 dtype이 폴백 없이 유입되지 않도록 변경 (#3335) (10-26)
      • 1.8 메인 저장소: DeepSelect Top-K 예제 병합, 129개 구성 중 torch.topk 대비 최대 8.00배 (#3304) (10-26)
      • 1.9 심사 중: TIRx 리팩터링 규모 1194개 파일로 확대, 패치되지 않은 업스트림 버전에 고정 (#3666) (10-26)
      • 1.10 신규 결함: ThreadSync 3종 배리어 오정렬, 심사 중 수정 3건 동반 (#3680 / #3682 / #3681) (10-26)
      • 1.11 신규 결함: sm_100 이상 리덕션 batch와 레이아웃 충돌, 심사 중 수정 1건 존재 (#3667 / #3684) (10-26)
      • 1.12 심사 중: 전치 대상 형상 검증, 영역화 kill 재배치 (#3677 / #3670) (10-26)
      • 1.13 종료됨: 퍼징 테스트 시리즈의 컴파일 타임 크래시 6건 (10-26)
    1. 멀티 백엔드 적응 (Ascend / Hygon / MetaX)
      • 2.1 Ascend 메인 저장소: 버전 지정 온칩 버퍼의 정밀 주소 지정 인터페이스 (#3607) (10-26)
      • 2.2 Ascend 심사 중: 수동 스케줄링 FlashMLA, BF16 프리필 및 FP8 디코딩 (#3678) (10-26)
      • 2.3 Ascend 심사 중: 온칩 메모리 오버플로 조기 거부, E4M3의 UB에서 L1 복사, 뷰 레이아웃과 포인터 단위 (#3655 / #3639 / #3623) (10-24~26)
      • 2.4 Ascend 심사 중: 배치 스케줄링 선택 가능, 크로스 코어 동기화에 배치 모드 요구, L0C 이중 복사 직접 전환 (#3675 / #3614 / #3468) (10-26)
      • 2.5 Ascend 저장소: 인스톨러의 SHMEM 세그먼트 이중 수정, 빈 껍데기 wheel이 더 이상 거짓 성공하지 않음 (#1943 / #1946) (10-26)
      • 2.6 Ascend 저장소: W4A4의 int4 행렬곱 결과 오류, 근본 원인은 니블 패킹이 1바이트로 처리된 것 (#1731) (10-26)
      • 2.7 Ascend 저장소: 2차원 초과 GM에서 UB 복사 시 DMA 행 수 기준으로 선행 차원 폴딩 (#1805) (10-26)
      • 2.8 Ascend 신규 결함: Cython 래퍼가 호출마다 분석기를 생성, 약 0.28밀리초 (#1945) (10-26)
      • 2.9 Hygon: 레지스터 파이프라인과 warp 분기 병합, 28개 파일 (#18) (10-26)
      • 2.10 Hygon: 네이티브 B16 저장 및 동적 WASP GEMM 병합, scale 버퍼 듀얼 뱅크 심사 중 (#28 / #29) (10-26)
      • 2.11 MetaX: 조직 뷰에 푸시 기록 있음, 6개 브랜치 모두 모니터링 기간 내 신규 커밋 없음 (10-26)
    1. 생태계 및 채택 측 (TileOPs / TileFoundry)
      • 3.1 TileOPs: 첫 릴리스 여전히 마지막 단계, 아티팩트는 여전히 2025-10-09 개발 버전 (#2738) (10-26)
      • 3.2 TileOPs: D128 밀집 전방 key 분리, prefill-rope 31.04마이크로초에서 17.09로 감소 (#2777) (10-26)
      • 3.3 TileOPs: D128 영속 전방 스레드별 재작성 및 LSE 출력 (#2774) (10-26)
      • 3.4 TileOPs: 융합 덧셈의 LayerNorm 행 CTA 간 분할, 정밀도 결함은 별도 신규 이슈 (#2768 / #2773) (10-26)
      • 3.5 TileOPs: 원소별 및 리덕션 4건, clamp의 bf16 행 5.9% 속도 향상 (#2766 / #2767 / #2769 / #2770) (10-26)
      • 3.6 TileOPs: 호출 텐서의 정렬 판단을 연산자 계층으로 상향, 3건 마무리 (#2772 / #2778 / #2780) (10-26)
      • 3.7 TileOPs: 오프셋 저장 3종 수정, SM90 그룹 GEMM 및 출력 ABI (#2761 / #2765 / #2762) (10-26)
      • 3.8 TileOPs 심사 중: FP8 쿼리의 페이지드 프리필을 e4m3 기준 value 곱셈으로 변경 (#2775) (10-26)
      • 3.9 TileOPs 심사 중: 희소 장시퀀스의 GD

N 프리필 분할 스캔(#2779)(10-26)

  • 3.10 TileFoundry: 동적 공유 메모리를 커널의 실제 뷰에 따라 크기 지정(#271)(10-26)
  • 3.11 TileFoundry: 전체 윈도우를 오퍼랜드 자체로 전달, sm_89의 GEMM 컴파일 복구(#273)(10-26)
  • 3.12 TileFoundry 신규 결함: 분석 측에서 공유 메모리 피크를 중복 계산하고 TIR로 전달하지 않음(#272)(10-26)
    1. 커뮤니티, 튜토리얼 및 이벤트
      • 4.1 언어 서비스: 실험적 CUDA 디버깅 출시, cuda-gdb 및 Nsight 연동(10-26)
      • 4.2 문서 사이트: 하루에 8건 업데이트(10-26)
      • 4.3 기여 프로세스 거버넌스 4연속: PR 설명 규범, 문서 전용 변경 시 백엔드 테스트 생략, Agent 제목 확인, 리뷰 봇 노이즈 감소(#3671 / #3672 / #3673 / #3674)(10-26)
      • 4.4 커뮤니티 신규 저장소: RISC-V 커뮤니티 허브 저장소, CANN 연산자 저장소, Hexagon 백엔드 저장소, 디버깅 도구 저장소(10-26)
      • 4.5 산업 측: Ascend와 DeepSeek 생태계 격차 축소 보도(10-26)
    1. 트렌드 관찰
  • 부록, 자료 및 검증 설명

오늘의 핵심

24시간 내 tile-ai 조직에 30건이 머지되었으며(tilelang 14, TileOPs 14, tilelang-hygon 2), 심사 중 업데이트 49건, 신규 이슈 11건 개설, 12건 종료. 어제 72시간 모니터링 기간의 78건 규모가 하루 만에 30건으로 줄었으나 리듬은 정상이며, 내용의 중심은 분명히 「성능 정리」에서 「스트림과 배리어 의미론」으로 이동했다: 메인 저장소에 머지된 출력 할당 수정(#3649)은 H100에서 다섯 번 실행 중 세 번 출력이 손상되던 문제를 0으로 낮췄는데, 원인은 세 어댑터가 호출자의 stream을 받아들이면서도 현재 스트림에서 out_idx 출력을 할당하고, 캐시 할당자가 할당 스트림 기준으로 정렬해 재사용하기 때문에 다른 스트림에 대기 중인 작업이 커널 결과를 덮어썼기 때문이다. 같은 날 드러난 세 건의 ThreadSync 결함(#3680)은 펜스가 분기 발산 경로에 배치된 문제, 분기 안의 펜스가 외부의 대기 중 의존성을 잘못 지우는 문제, 그리고 스레드 카운트 기반 루프에서의 펜스 위치를 지적했으며, #3682와 #3681 두 건의 심사 중 수정이 함께 있으며, tvm 포크 저장소의 한 건은 불리언 비트 연산을 논리 연산으로 번역해 증명기가 더 이상 10000 스텝 예산 내에서 포기하지 않게 했다(#88). 릴리스 측은 진전이 없다: TileOPs의 첫 릴리스 추적 티켓(#2738)은 TestPyPI 시험 실행과 첫 tag만 남았고, PyPI 최신은 여전히 2025-10-09의 0.0.1.dev2이며 저장소에는 tag도 release도 없다. 백엔드 측 Hygon은 하루에 두 건의 대형 변경을 머지했다: 레지스터 파이프라인에 warp 발산 추가(#18, 28개 파일 +5312/-64)와 네이티브 B16 저장에 동적 WASP GEMM 추가(#28); Ascend 메인라인은 버전이 있는 온칩 버퍼를 갖는 정밀 주소 지정 인터페이스를 머지했다(#3607, 14개 파일 +982/-26). 심사 중 가장 큰 것은 여전히 #3666으로, 규모가 어제의 1000여 개에서 1194개 파일 +106648/-104497로 확대되었다.


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

1.1 메인 저장소: out_idx 출력을 호출자 스트림에서 할당하도록 변경, H100에서 재현되던 5회 중 3회 손상 제로화(#3649)

10-26 22:36에 머지, 5개 파일 +183/-31, #3648 수정. cython, nvrtc, cutedsl 세 어댑터 모두 stream=을 받아 해당 스트림에서 실행하지만, out_idx 출력은 torch.empty로 현재 torch 스트림에서 할당되었다; 캐시 할당자는 「할당 스트림」 기준으로만 정렬해 재사용하므로, 여전히 대기 중인 작업이 있는 메모리 블록이 유휴 블록으로 반환되고 그 작업이 이후 커널 결과를 덮어썼다. 수정 방법은 output_allocation_stream(stream, device)를 새로 추가하는 것이다: 호출자 스트림이 현재 스트림과 다르면 호출자 스트림을 반환하고, 그렇지 않으면 빈 값을 반환한다; 핸들이 0이면 기본 스트림에 매핑되는데, 구버전 PyTorch의 ExternalStream(0)이 기본 스트림이 아닌 풀링 스트림을 반환하기 때문이다. H100 80GB에서의 재현은 4/5, 5/5, 5/5 손상에서 세 어댑터 각각 0/5로 바뀌었고, 신규 테스트 파일 10개 항목이 통과했다. 작성자는 wait_stream을 추가하는 방식를 명시적으로 반대했다: 그것은 매 호출마다 시작 스트림을 현재 스트림의 모든 작업 뒤에 정렬시키게 된다.

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

1.2 메인 저장소: NVRTC 디바이스 측 컴파일을 TVM-FFI 실행 경로에 연결, 컴파일 키 항목 추가(#3285)

10-27 00:02에 병합되었으며, 12개 파일 +133/-133입니다. 새로운 tl.cuda_compiler="nvrtc"가 추가되어 기존 TVM-FFI 호스트 런타임과 함께 동작합니다. NVRTC는 디바이스 측 CUDA 소스를 cubin으로 컴파일하고, 공유 CUDA 런타임 모듈이 이 바이트를 임포트합니다. 호스트 측 바인딩, 스칼라 변환, 조건, 루프 및 런치 준비는 여전히 TVM-FFI가 생성하며, 기본값은 여전히 nvcc이고 기존 execution_backend="nvrtc" 어댑터는 원래 구현을 유지합니다. 바이너리 캐시 키에 컴파일러 이름과 NVCC/NVRTC 버전이 추가되어 타깃 세부 사항 및 유효 컴파일 옵션과 함께 해시에 참여합니다. 아키텍처 토큰인 90a와 100f는 원형 그대로 유지되며, 호환되지 않는 타깃은 컴파일 전에 진단을 보고합니다. 성공적으로 생성된 각 NVRTC program은 finally에서 소멸됩니다. 이 PR의 CI는 H20에서 5506개 통과, 491개 스킵을 기록했고, 예제는 76개 통과, ROCm 7.2 채널은 3220개 통과, Ascend 채널은 1797개 항목 통과했습니다.

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

1.3 메인 저장소: 배치 TMA 로드의 cache hint에 템플릿 인자가 하나 부족해 DeepSelect의 TMA 경로가 컴파일되지 않음(#3683)

10-27 00:11에 병합되었으며, 2개 파일 +20/-11입니다. eviction_policy="evict_first" 또는 "evict_last"인 배치 TMA 로드가 main에서 컴파일 실패합니다. 코드 생성이 CacheHintSm90 템플릿 인자를 하나 내보내는데, 베어 포인터 버전의 tl::tma_load 오버로드는 배리어 타입만 받기 때문입니다. 이 결함은 DeepSelect의 TMA 경로를 직접 막았으며, GB200 + CUDA 13.1에서 해당 예제의 ring 및 replay 케이스로 재현되었습니다. 수정 방법은 #3296에서 캐시 힌트 부분을 추출하여 누락된 템플릿 인자를 보충하고 선택된 정책을 PTX 명령에 전달하는 것입니다. 공유 메모리 정렬 부분은 여전히 #3296에 남아 있습니다. 수정 후 동일 머신에서 DeepSelect 전체 파일 144개 항목 통과, TMA 복사 테스트 22개 항목 통과, CUDA 12.9의 SM90 컴파일 검사에서도 세 가지 정책의 PTX 모두에 마땅한 오퍼랜드가 포함됨을 확인했습니다(#3296의 나머지 변경은 아직 병합되지 않음).

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

1.4 메인 저장소: 리덕션의 FP8/FP4 비유한 상수가 불법 리터럴을 출력, 서브바이트 누산기를 조기 거부로 변경(#3189)

10-26 15:44에 병합되었으며, 7개 파일 +204/-4, #2998 수정입니다. T.reduce_max와 T.reduce_min이 출력이 FP8 E4M3일 때 불법인 inff 리터럴을 출력했습니다. 이제 비유한 E4M3, E8M0, FP4 상수는 기존 변환 생성자를 통해 유효한 CUDART_INF_F와 CUDART_NAN_F 표현식을 전달하며, 포화 및 NaN 동작은 그대로 유지됩니다. 같은 변경으로 일반 및 지연 리덕션에서 서브바이트 수치 누산기(FP4, int4 등)를 조기에 거부하고 명확한 진단을 제공합니다. FP4 입력에 FP16 또는 FP32 누산은 여전히 지원되며, 불리언 리덕션은 면제되고, E8M0 리덕션 지원은 이번 범위에 포함되지 않습니다. B300 + CUDA 13.1에서 28개 집중 테스트 전부 통과, 전체 리덕션 테스트 파일 150개 통과, 나머지 세 가지 배치 검증 실패는 main에 이미 존재하던 것입니다.

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

1.5 메인 저장소: SM120 블록 스케일링 GEMM의 K 원자 명시적 전개, FullCol 구성 51% 향상(#3286)

10-26 16:36에 병합되었으며, 4개 파일 +308/-2입니다. #3257이 T.gemm_blockscaled와 T.mma_gemm_blockscaled를 도입한 후의 마무리 작업입니다. A 오퍼랜드가 레지스터 분할에서 올 때, K 원자 루프가 원래 T.serial을 사용하여 매 분할 접근이 런타임 인덱스였고, nvcc는 자체 전개하지 않는 파이프라인화 및 warp 특수화 커널에서 분할을 로컬 메모리에 남겨 MMA 경로를 오버플로시켰습니다. T.unroll(..., explicit=True)로 변경한 후 모든 접근이 컴파일 타임 상수가 됩니다. 두 번째는 SFA와 SFB가 모두 공유 메모리에 있을 때, 기존에는 MMA 쌍마다 공유 메모리를 세 번 읽었지만 이제 lane마다 K 원자마다 컴팩트 선택 팩을 한 번만 적재하며, 행 우선 및 블록 K 우선 두 레이아웃을 모두 커버합니다. RTX PRO 6000 Blackwell에서 143개 통과, 4개 스킵, 출력은 float32 참조와 비트 단위로 일치합니다. 레지스터 A 벤치마크에서 FullRow 2.7% 향상, Square 11.1%, FullCol 51%(348.5에서 526.7 TFLOPS), ptxas 보고 스필이 664바이트 및 484바이트에서 320과 160으로 감소했습니다. 작성자는 또한 block_K = 128에 행 우선 공유 scale을 조합한 경로가 main에서 이미 검증을 통과하지 못하며 이번 변경과 무관하고 추후 별도 조사가 필요하다고 표시했습니다.

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

1.6 메인 저장소: JIT 출력 형상의 스칼라 파라미터가 할당 전에 분석됨 (#3602)

10-26 10:06에 병합, 7개 파일 +242/-3, #3406 수정. Cython 및 PTO 래퍼는 이전에 출력을 할당할 때 명시적 스칼라 PrimFunc 파라미터로 표현된 출력 차원을 분석하지 못했음; 이제 어댑터가 “출력 차원에서 스칼라 파라미터 위치로”의 매핑을 전달하고, 래퍼는 실인자 값을 직접 사용하며 중복된 시작 파라미터를 추가하지 않음. 회귀 테스트는 두 백엔드, 영차원 및 비영차원, 스칼라 및 다차원 출력, 출력 순서, 디바이스 배치 및 시작 파라미터 개수를 커버함. 동일 계열 결함의 후속 두 건은 여전히 심사 중: #3636은 (n, 2*n) 같은 복합 차원 표현식을 처리하고, #3638은 동일 이름이지만 서로 독립적인 동적 심볼이 이름으로 병합되는 문제를 처리함.

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

1.7 메인 저장소: sm8x 희소 메타데이터의 피연산자 dtype이 더 이상 폴백 없이 통과하지 않음 (#3335)

10-26 16:54에 병합, 1개 파일 +12/-5. 희소 메타데이터의 레이아웃 생성자는 아키텍처별로 디스패치하여 MMA 피연산자 dtype과 메타데이터 버퍼 폭을 페어링함; sm90 브랜치는 지원되지 않는 dtype을 거부하지만, sm8x 브랜치는 두 개의 화이트리스트 검사로 작성되어 폴백이 없어, 다른 표기법은 모두 아래의 배치 로직으로 떨어지며 그 로직은 버퍼 폭에서만 group과 interweave를 추론함. 그 결과 float8_e4m3fn(torch의 정규 표기법)이 16비트 메타데이터로 허용되고 fp16의 group 32 및 interweave 4 규약이 적용되었음; 인식되지 않는 모든 표기법 including 오타도 임의 폭에서 허용되었음. 이제 지원되지 않는 dtype은 직접 오류를 발생시킴.

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

1.8 메인 저장소: DeepSelect Top-K 예제 병합, 129개 구성에서 torch.topk 대비 최대 8.00배 (#3304)

10-26 16:53에 병합, 4개 파일 +1265/-0. 예제는 DeepSeek의 DeepSelect를 TileLang으로 재작성했으며 인라인 PTX를 포함하지 않음; 작성자는 129개 구성, 10개 사용 불가를 실측하여 RTX 5090(170개 SM) + CUDA 12.9 + PyTorch 2.7에서 반복 입력과 캐시 클리어 두 가지 방법으로 실행함. 반복 입력 기준에서 BF16, 배치 6, N이 16384, K가 512인 구성의 지연은 5.059마이크로초 대 torch.topk의 40.448로 8.00배 가속; N이 1048576인 구성은 15.503 대 100.254로 6.47배. FP32와 K가 4096인 조합은 해당 래퍼에서 사용 불가이며, K가 1인 네이티브 경로는 정렬 주소 오류가 있어 작성자는 표에서 가로줄로 “사용 불가”를 표시하고 0으로 표기하지 않음. 이 예제가 의존하는 캐시 힌트 수정은 바로 같은 날 병합된 #3683임.

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

1.9 심사 중: TIRx 리팩터링 규모가 1194개 파일로 확대, 패치되지 않은 업스트림 버전에 고정 (#3666)

10-25에 개설, 10-26에 세 번째 업데이트, 현재 1194개 파일 +106648/-104497. 방안은 TileLang을 내장 TVM에서 독립 배포되는 TIRx로 마이그레이션하고, 고정된 업스트림 베이스라인에 제로 패치를 적용하는 것: 3rdparty/tvm 및 그 서브모듈 항목을 삭제하고, 3rdparty/에는 CUTLASS, HIP 헤더 파일 및 ignore 파일만 남음; 수정되지 않은 apache/tvm-tirx 소스는 9af5544d에 고정되며, apache-tvm-tirx==0.1.0a2와 apache-tvm-ffi==0.1.15.post0에 의존하고, 기존에 +tilelang.1 접미사가 붙던 통합 패치는 더 이상 필요하지 않음. 이는 현재 큐에서 가장 큰 규모의 건이며, 이후 TVM 측 수정(예: 같은 날 tvm 포크 저장소의 Z3 불리언 비트 연산)이 메인라인에 어떻게 진입할지를 결정함.

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

1.10 신규 결함: ThreadSync 세 가지 유형의 오정렬 배리어, 세 건의 심사 중 수정 동반 (#3680 / #3682 / #3681)

#3680은 10-26 19:56에 생성되어, tl.ThreadSync가 세 곳에서 배리어를 잘못 배치한다고 지적했다. 스레드별 값을 대상으로 한 가드를 블록 내 일관으로 간주하여, 배리어가 분기 안에 남아 분기를 건너뛴 스레드를 기다리게 된다. if의 특정 분기 안의 배리어가 외부 시퀀스의 미결 의존성을 지워, 블록이 해당 분기를 건너뛸 때 if 이전의 쓰기와 if 이후의 읽기가 충돌한다. 루프의 이동 횟수가 스레드별(예: 스레드 번호로 나눈 나머지)일 때, 배리어가 루프 본문 안에 남아, 부분 동기화의 하강이 「스레드 사각형을 증명할 수 없음」으로 실패한다. #3682(3개 파일 +39/-7 수준 변경, 12개 신규 테스트)는 사전 분석 한 차례로 스레드별 값을 가질 수 있는 버퍼를 수집하고, 동기화 요약을 분기와 루프의 도달 가능성에 따라 축소한다. 작성자는 이 변경이 경쟁을 보고할 뿐 제거하지 않으며, 분기 분할도 하지 않는다고 밝혔다. #3681은 또 다른 같은 기원의 크래시를 수정한다. sm_100부터 float32 루프가 256비트로 벡터화되어, 8개 lane으로 브로드캐스트할 때 생성된 표현식이 임시 객체 주소를 취하여 NVRTC가 「lvalue여야 한다」로 컴파일을 거부하던 것을, __float_as_uint에 시프트를 더해 직접 비트를 조립하도록 바꿨다. 함께 제공되는 tvm 포크 PR #88은 불리언 비트 연산을 논리 연산으로 번역하여, 증명기가 10000 스텝 예산 내에서 같은 결론에 도달할 수 있게 한다.

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

1.11 신규 결함: sm_100 이상 리덕션 batch와 레이아웃 충돌, 심사 중인 수정 1건 존재(#3667 / #3684)

#3667은 10-26 08:31에 생성되었다. T.reduce_*가 sm_100과 sm_120에서 batch를 N과 같게 지정하면 하강할 수 없는데, fragment 복사가 sm_100 이상에서 256비트 로드를 사용하여 스레드당 보유 출력 수가 줄어들어, sm_90에서 합법이던 batch가 하드 검증에 걸리기 때문이다. #3684는 10-27 05:44에 생성되어 수정 방법을 첨부했다. ReduceOp는 N의 batch를 초과하지 않는 최대 진약수를 취하고 경고를 주도록 바꾸며, finalize_reducer는 테스트가 동작을 고정하고 있으므로 엄격함을 유지한다. 새로 추가된 sm_100a와 sm_120a 하강 테스트는 main에서 실패하고 이 브랜치에서 통과하며, 실측기는 RTX 5070 Ti이다. 같은 리덕션 문제의 또 다른 건인 #3341(일괄 전체 리덕션의 스레드 범위 해석)은 여전히 심사 중이다.

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

1.12 심사 중: 전치 대상 형상 검증, 영역화 kill 재배치(#3677 / #3670)

#3677은 tl.transpose의 대상 형상 검증을 TransposeNode 생성자에 넣어, NormalizeToAccessRegion 이후에 실행되므로 Python 래퍼를 우회하여 직접 생성해도 커버된다. 동등성은 Python 차원 목록을 비교하는 대신 arith::Analyzer::CanProveEqual로 판정하며, (3, n+1)과 독립적으로 생성된 (n+1, 3)은 여전히 허용되고, 상수가 불일치하면 축 번호와 양측 길이로 오류를 보고한다. 이전에는 불일치하는 대상이 하강하여 범위를 벗어난 인덱스가 되어 잘못된 결과나 범위를 벗어난 쓰기를 생성했다. CPU 대상 13개 테스트가 통과한다. #3670은 공유 메모리 활성도 분석에 남아 있던 scope_level 메커니즘을 교체한다. 이 값은 SeqStmt를 포함한 For 본문에서만 증가하고 While에서는 결코 증가하지 않아, 두 방향 모두 틀린다. While의 경우 kill 지점이 역방향 엣지를 넘지 못해 패커가 상호 배타적 영역을 겹치게 하여 sm90에서 실측 결과가 매 라운드 두 배가 되고, 형제 루프의 경우 kill이 전체 영역 경계를 넘어가게 한다. 새 규칙은 실제 외부 For 및 While 체인과 각 루프 구간의 연속 구간으로 계산하여, 경계 초과와 조기 중단 두 방향 모두 억제된다.

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

1.13 종료됨: 퍼징 시리즈의 컴파일 타임 크래시 6건(10-26)

10-26에 퍼징에서 나온 컴파일 타임 크래시 또는 잘못된 코드류 결함 6건을 일괄 종료했다. int16과 uint16 대상의 T.atomic_add는 nvcc에 16비트 오버로드가 없어 바로 중단된다(#2572). 불리언이 아닌 버퍼 요소를 if나 assert 조건으로 직접 사용하면 내부 is_bool() 단언으로 중단된다(#3009). 부동소수점 피연산자를 정수만 받는 div 및 mod 계열에 넣어도 동일하게 내부 단언이 발생한다(#3008). int2, uint2처럼 1, 4, 8, 16, 32, 64에 속하지 않는 정수 폭은 CUDA 코드 생성을 중단시킨다(#2996). reduce_max와 reduce_min을 FP8 또는 FP4 샤드에 기록하면 크래시가 발생한다(#2998). 그리고 전체 버퍼 목적지의 원자 리덕션이 0번 스레드 외의 모든 스레드를 조용히 누락시킨다(#2959). 마지막 건은 문서 측 수정도 함께 병합되었으며(#3340), 예제의 목적지를 전체 버퍼에서 명명된 요소로 바꾸어 예제 자체가 잘못된 작성법을 계속 가르치지 않도록 한다.

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


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

2.1 Ascend 메인 저장소: 버전이 있는 온칩 버퍼의 정밀 주소 지정 인터페이스(#3607)

10-26 10:04에 병합되었으며, 14개 파일에서 +982/-26입니다. T.annotate_buffer_addresses가 새로 추가되어, 스케줄링된 Ascend 파이프라인에서 UB, L1, L0A, L0B, L0C 및 SS 각 풀 내의 할당에 대해 정확한 바이트 주소 지정을 할 수 있게 합니다. 주소가 지정된 버퍼는 먼저 T.annotate_buffer_versions로 고정된 복사본 수를 선언해야 합니다. 검증은 버전 전개와 레이아웃 보정 이후에 수행되며, 누락되거나 지원되지 않는 복사본 수, 중첩, 정렬 불량, 동적 크기 및 용량 초과를 포함합니다. 할당자는 또한 SIMT의 UB 컨텍스트 예약과 AIV 사설 주소 공간을 위해 자리를 비워야 하며, L0의 scale factor 핸들은 해당 데이터 할당을 따릅니다. 부분 매핑은 주어진 구간만 고정하고, 나머지 버퍼는 정렬된 빈 공간에 배치되며 이미 점유된 주소를 재사용할 수 없습니다. 주소 지정 주석이 없는 풀은 기존 동작을 유지합니다. 서브코어 프로젝션 시에는 활성 저장소 복사본에 있는 주석만 유지됩니다.

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

2.2 Ascend 심사 중: 수동 스케줄링 FlashMLA, BF16 프리필 및 FP8 디코딩 (#3678)

10-26 19:31에 열렸으며, Ascend 950에서의 BF16 프리필과 컴팩트 FP8 디코딩을 다루고, 선택적 FP8 및 FP4 추가 KV cache, attention sink, 불법 인덱스 및 쿼리별 길이를 포함합니다. 어텐션 루프, 온라인 softmax, 역양자화 및 스케줄링은 모두 TileLang으로 표현되며, 소량의 intrinsic 헤더 파일만 단일 하드웨어 명령을 감쌉니다. 수동 구현은 AIC와 AIV 파이프라인 및 동기화를 명시적으로 제어하며, 영속화 버전은 T.Stage와 컴파일러가 관리하는 버퍼 버전 및 동기화를 사용하고, QK와 PV는 동일한 L0B 저장소 링의 서로 다른 레이아웃 뷰를 사용합니다.

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

2.3 Ascend 심사 중: 온칩 메모리 오버플로 조기 거부, E4M3의 UB에서 L1 복사, 뷰 레이아웃 및 포인터 단위 (#3655 / #3639 / #3623)

#3655는 MergeUBAllocations가 하강 단계에서 대상 물리 용량에 따라 오버플로를 거부하도록 합니다. 두 개의 독립 루프가 각각 합법적으로 32 KB의 L0A 샤드 두 개를 선택하고, 병합 후 64 KB가 정확히 가득 차며, 이전에는 물리 버퍼 끝을 넘어서는 포인터 오프셋을 크기 메타데이터 없이 내보냈지만, 이제는 짧은 진단을 제공합니다. #3639는 E4M3의 UB에서 L1 복사를 활성화합니다. ND 입력은 기존 scatter 및 후복사 파이프라인을 통해 자동 재배열되거나, 미리 배열된 NZ 입력으로 직접 복사되며, 기존 형상 검사를 재사용하고 풋프린트와 지연 분석에 FP8 요소 크기 처리를 추가합니다. #3623은 T.view와 T.reshape이 만든 버퍼가 각자의 레이아웃을 유지하도록 하여, 공유 데이터 변수를 통해 레이아웃을 해석할 때 Cube 매핑이나 DMA 주소를 잘못 선택하는 것을 방지하고, 원시 포인터 오프셋을 자신의 요소 단위로 재매핑합니다.

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

2.4 Ascend 심사 중: 배치 스케줄링 선택 사항, 크로스코어 동기화 배치 모드 요구, L0C 이중 복사 직접 변환 (#3675 / #3614 / #3468)

#3675는 Ascend에 선택적 배치 스케줄링(schedmode 1)을 추가합니다. 정상 모드에서는 스케줄러가 필요한 코어 리소스가 부분적으로 사용 가능하면 연산자를 시작하며, 여러 스트림의 전체 코어 동기화 연산자가 교차할 수 있습니다. 배치 모드에서는 모든 리소스가 준비될 때까지 기다린 후 전체를 시작하며, 전체 코어 동기화 시나리오에서는 반드시 이를 사용해야 합니다. #3614는 동일 메커니즘이 mode 0에서 나타나는 증상과 수정 방법을 설명합니다. #3468은 A5에서 FixPipe가 이중 대상 모드와 quant_pre를 동시에 켤 수 없는 제한을 처리하여, T.dual_copy가 L0C의 fp32에서 f16/bf16 변환을 직접 수행하도록 하여 Vector 변환 한 번과 fp32 UB 버퍼 한 블록을 절약하고, L0C에서 UB로의 fp32 대역폭도 절약합니다.

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

2.5 Ascend 저장소: 설치 프로그램의 SHMEM 섹션 이중 수정, 빈 껍데기 wheel이 더 이상 가짜 성공하지 않음 (#1943 / #1946)

두 건의 심사 중 PR이 동일한 결함(#1728)을 가리킵니다. install_ascend.sh --enable-shmem은 build.sh 실패 후 반환 코드를 확인하지 않고 계속 bdist_wheel을 실행하며, package_data의 와일드카드가 파일을 찾지 못해도 오류를 보고하지 않아, __init__.py만 있는 빈 껍데기 패키지를 만들고 Successfully installed shmem-1.0.0을 출력하며, 이십여 줄 후에 source ../../../install/set_env.sh에서 “shmem 환경 설정 실패”를 보고하여 문제 해결 방향을 환경 구성으로 유도합니다. #1946은 두 곳만 수정합니다. 반환 코드 검사를 추가하여 즉시 종료하고, set_env.sh를 스크립트 디렉토리 기반 절대 경로로 변경하며, shmem 섹션을 문자 그대로 추출하는 제어된 harness로 검증합니다(수정 전 10개 중 6개 실패 4개 통과, 수정 후 10개 모두 통과). #1943은 다른 경로를 취합니다. 기존 패키지를 제거하기 전에 네이티브 빌드, 환경 스크립트, wheel 빌드 및 wheel 존재 여부를 순서대로 확인하고, 10개의 CPU 제어 흐름 테스트를 추가하며, 이것이 설치 프로그램의 연속 실행과 진단을 수정하는 것이지 업스트림 SHMEM 컴파일 실패 자체를 포함하지 않는다고 설명합니다.

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

2.6 Ascend 저장소: W4A4의 int4 행렬 곱셈 결과 오류, 근본 원인은 니블 패킹이 1바이트로 취급된 것(#1731)

#1731은 10-26에 업데이트되었으며, 위치 파악과 재현이 상당히 완전하다. 전체 C++ 템플릿 스택이 sizeof(AscendC::int4b_t) == 1로 프랙탈 지오메트리를 추론하지만, 해당 타입은 니블 패킹이므로 32바이트 C0 블록의 올바른 원소 수는 32가 아닌 64이며, 이에 따라 세 계층의 이동 모두에서 2배 편차가 발생한다: L1에서 L0A/L0B로의 LoadData 파라미터가 잘못된 상수로 계산되어 데이터가 온칩에서 잘못된 위치에 적재된다; B 오퍼랜드의 전치 로딩은 load_cbuf_to_cb_s4에 의해 하드코딩되어 비활성화되고 플래그가 조용히 폐기되어, B는 mad_s4가 기대하는 nZ 레이아웃을 결코 얻지 못한다; UB와 GM 사이의 바이트 길이는 원소당 1바이트로 계산되어 각 행이 두 배로 복사되고 인접 행이 섞인다. 계산 명령 자체는 올바르므로 증상은 「크기는 정상, 수치는 전부 오류」(원 보고서 99.6% 불일치)이다. 수정은 모두 src/tl_templates/ascend/common.h와 두 곳의 변환에 적용되었으며, 새로운 W4A4 예제와 10건의 int4 회귀가 추가되어, 1024 큐브 엔드투엔드가 99.6% 불일치에서 불일치 없음으로 바뀌었고, 기존 gemm_v0 스위트 48건, l1_to_l0와 gm_to_ub 263건, 원소별 492건이 모두 통과했다.

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

2.7 Ascend 저장소: 2차원을 초과하는 GM에서 UB로의 복사가 DMA 행 수로 선행 차원을 접음(#1805)

#1805는 #1231을 수정한다. copy_gm_to_ub와 copy_ub_to_gm의 DMA 모델은 2차원 사각형인데, AscendCopy::Lower는 region의 활성 차원 중 마지막 두 개만 행과 열로 취하며, 활성 차원이 두 개를 초과하면 선행 차원이 조용히 폐기된다: 3D 블록 복사 A[t:t+16, :, :](region이 16, 4, 8로 총 512 원소)가 4와 8로 강등되어 DMA가 32개 원소만 옮기고, UB 샤드의 나머지 15개 토큰은 초기화되지 않아 이후 계산이 쓰레기 값을 읽는다. 수정 방법은 선행 차원을 DMA의 행 수로 접는 것이며, 두 방향 모두 동일하게 변경한다. 같은 날 Ascend 저장소에는 문서 사실 확인 PR 일괄 업데이트(sin과 cos, axpy, clamp, bitwise 계열, clear, fill)가 있었으며, 그중 clamp 건은 PTO에서 T.tile.clamp의 하한이 무효인 문제도 함께 수정했다.

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

2.8 Ascend 신규 결함: Cython 래퍼가 호출마다 분석기를 생성, 약 0.28밀리초(#1945)

10-26 15:52에 개설되었으며, 보고자가 계층별 위치 파악과 인과 검증을 제시했다. CythonKernelWrapper.forward()는 호출마다 무조건 TVM Analyzer()를 하나 생성하며, 로컬 중앙값은 약 0.28밀리초인데, 이는 동적 shape 분기에서 한 번만 사용된다; 정적 shape 커널의 dynamic_symbolic_map은 빈 사전이므로 simplify()는 결코 실행되지 않는다. 인과 검증 방법은 모듈의 Analyzer 전역을 「simplify 호출 시 오류를 던지는」 스텁으로 교체하는 것이었고, 모든 호출이 정상적으로 성공했으며 전체 호출 소요 시간이 442~467마이크로초에서 104~107마이크로초로 감소했고 커널 규모와 무관했다(1개 block과 32개 block이 동일한 수준). 실측 환경은 Ascend 910B3와 CANN 9.2.0이다. 이는 소규모 연산자를 밀집 호출하는 추론 경로에 고정 오버헤드이므로 수정 방향을 주목할 가치가 있다.

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

2.9 Hygon: 레지스터 파이프라인과 warp 발산 병합, 28개 파일(#18)

10-26 13:33에 병합되었으며, 28개 파일 +5312/-64, 10-08에 개설되어 2주 넘게 수렴되었다. 변경은 tilelang-hygon에 레지스터 파이프라인을 추가하여 기존 공유 메모리 파이프라인과 협력할 수 있게 하며, 대응하는 커밋과 배리어 명령이 파이프라인 단계에 따라 삽입된다; warp 수가 4를 초과하면 warp 발산이 활성화된다. 작성자가 제시한 결론은 example_gemm이 rocblas에 상응하는 성능에 도달했다는 것이다. 같은 저장소에서 당일 주의해야 할 점은 #18의 병합 순서이다: #18이 #28보다 먼저 병합되었으며, 이는 메인라인 변경을 브랜치에 병합한 후의 대규모 병합에 해당하므로 행 수가 #28보다 훨씬 크다.

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

2.10 Hygon: 네이티브 B16 저장과 동적 WASP GEMM 병합, scale 버퍼 듀얼 뱅크 심사 중(#28 / #29)

#28은 10-26 16:23에 병합되었고, 6개 파일 +340/-230, 세 가지 작업: MLS의 Let 파라미터 방출 순서 유지, GEMM의 warp 분할을 행렬 저장소로 전파, 패킹된 B16에서 B32 출력 별칭 거부; 동시에 WASP 차원을 동적으로 만들고, 축별 정렬 추론, pong의 K 꼬리 필터 수정, WASP 엔드투엔드 회귀와 기본 MLS 테스트 분리.#29는 10-26 17:13에 개설되어 리뷰 중: scale 버퍼의 행 계획에 LDS bank와 분리된 이중 bank 파티션을 부여하여 할당, scale 복사, MMAC 하강이 동일한 행 매핑을 공유하도록 하고, 물리적 구간 수명 주기와 m0 정렬 페이징을 유지하며, 유휴 wave의 작업을 두 bank에 분산하고, 수동 warp 특수화 하의 블록 스케일링 GEMM에 사용.

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

2.11 MetaX: 조직 뷰에 푸시 기록이 있으나, 6개 브랜치 모두 모니터링 기간 내 신규 커밋 없음 (10-26)

조직 페이지에는 tilelang-metax가 모니터링 기간 내 푸시 타임스탬프를 표시하지만(베이징 시간 10-26 15:08), 브랜치별로 대조한 결과 대응하는 커밋을 찾지 못했다: 기본 브랜치 main의 최신은 여전히 2026-09-29의 버전 번호 상향, dev 브랜치 최신은 2026-10-16의 한 번의 병합, 0.1.14 브랜치는 2026-09-17, 0.1.12와 race는 더 이전; 이벤트 스트림에도 모니터링 기간 내 푸시 이벤트가 없다. 이에 따라 「실질 변경 없음」으로 처리하고, 이번 호 항목에 포함하지 않으며, 그 성격을 추측하지도 않는다. 해당 저장소의 최근 공식 릴리스는 여전히 10-16의 0.1.15이다.

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


3. 생태계와 도입 주체 (TileOPs / TileFoundry)

3.1 TileOPs: 첫 릴리스는 여전히 마지막 칸, 아티팩트는 여전히 2025-10-09 개발 버전 (#2738)

첫 릴리스 추적 이슈가 10-26 갱신되었고, 8개 정리 항목 중 6개 완료: 네이밍과 매니페스트 파라미터, 패키징과 저장소 파일, Python 3.13, 공개 API와 확장점, 연산자 동작 규약, 입력 검증 생성, 문서 사이트. 남은 두 칸: 야간 벤치마크(깨진 두 기준선 수정 #2763 병합됨; 새 행을 모든 대조기 위로 끌어올리고 누적 스캔을 단일 기동으로 변경 #2766 병합됨, GQADenseWithLSEFwdOp의 D128 프로그램 #2774 병합됨, 그러나 MHA 행은 여전히 FA3의 0.996에 불과, 튜닝 계속)와 릴리스 자체(제3자 고지, PyPI 설치 문서, 개발자 의존성 그룹, CHANGELOG와 릴리스 절차 #2753, 그리고 maintainer가 실행해야 하는 TestPyPI 시험 실행과 첫 tag). PyPI와 저장소 확인: tileops 최신은 여전히 2025-10-09의 0.0.1.dev2, 총 세 버전, 저장소에 tag와 release가 전혀 없음. 첫 릴리스는 「문서 준비 완료, 아티팩트 미업로드」의 마지막 칸에 있으며, 이미 여러 날 지속되었다.

출처: https://github.com/tile-ai/TileOPs/issues/2738

3.2 TileOPs: D128 dense forward key 분할, prefill-rope 31.04마이크로초에서 17.09로 감소 (#2777)

10-26 22:45 병합, 2개 파일 +527/-124. D128 dense 호출에서 128행 타일이 SM을 유휴 상태로 만들면, 각 타일의 key를 min(SM 수 나누기 타일 수, 가장 짧은 타일의 key 스텝 수)份으로 분할하고, 분할이 가장 긴 타일에서 두 스텝의 key를 균등하게 덜어낼 수 있을 때만 분할; 각份은 해당 행의 스케일 최댓값과 합을 가진 미정규화 fp32 출력을 작성하고, 병합 커널이 최대 최댓값으로 가중하므로 상쇄 항과 큰 logit 모두 정밀도를 잃지 않음. 타일의 마지막 key 스텝에서, 어떤 warpgroup이 마지막으로 필요로 하는 key가 낮은 64개 key 내에 있으면, softmax는 그 64개에만 마스킹과 지수를 수행하고, warpgroup 0의 인과 대각 스텝은 반드시 적중. H200 1500 MHz에서 prefill-rope의 f16과 bf16이 31.04와 31.01에서 17.09와 17.09마이크로초로 감소, FA3 대비 2.324와 2.328배; qwen3-8b 블록 프리필이 31.30에서 17.79로 감소. 테스트 노드 순증 14.

출처: https://github.com/tile-ai/TileOPs/pull/2777

3.3 TileOPs: D128 영속 forward를 스레드별로 재작성하고 LSE 출력 (#2774)

10-26 18:22에 병합되었으며, 2개 파일에서 +656/-48입니다. 근본 원인은 GQADenseWithLSEFwdOp가 이전에 비영속 warp 특화 커널에서 실행되었기 때문인데, 영속 커널이 return_lse를 거부하여 야간 벤치마크 실측에서 FA3의 0.73~0.78에 그쳤습니다. 이번에는 TileLang의 스레드별 원시 연산(TMA 매핑, mbarrier, WGMMA 디스크립터, warp shuffle, stmatrix)으로 D128 영속 밀집 전방을 재작성하여 log-sum-exp를 직접 출력하고, 동일한 구현이 GQADenseFwdOp와 GQADenseWithLSEFwdOp를 동시에 지원합니다. 모든 key와 value가 L2의 절반을 초과할 때 한 헤드의 쿼리 블록을 인접 배치하여 동시 타일이 L2의 K와 V를 공유하도록 하며, 각 타일의 첫 QK는 이전 타일의 꼬리 구간에서 시작됩니다. 공개 인터페이스와 매니페스트는 변경되지 않았습니다. 병합 후 MHA 행은 여전히 FA3의 0.996이며, 추적 단일 주석 튜닝이 계속되는데 이는 #2777과 두 개의 독립적인 최적화 라인에 속합니다.

출처: https://github.com/tile-ai/TileOPs/pull/2774

3.4 TileOPs: 융합 덧셈의 LayerNorm 행 CTA 간 분할, 정밀도 결함은 별도 이슈로 (#2768 / #2773)

#2768은 10-26 15:45에 병합되었으며, 5개 파일에서 +525/-24입니다. 시작점은 야간 벤치마크의 회귀(llama-70b 디코드 행이 0.0037에서 0.0044밀리초로 상승)였는데, 조사 결과 코드 변경이 아니라 튜닝 재선택이었습니다. 튜닝 키 변경 후 FusedAddLayerNormKernel이 재튜닝되었고, 튜너가 L2 백투백 타이밍 후보에서 1024 스레드를 선택하여 콜드 스타트가 256 스레드보다 19% 느렸습니다. 새로 작성된 FusedAddLayerNormHeldKernel은 한 행(SM90에서 하나의 cluster CTA 몫)을 레지스터에 유지하고, 가장 바쁜 SM의 비용이 더 낮을 때 한 행을 최대 8개 CTA로 분할하며, 블록 크기는 전체 warp 수 중 몫 접근량을 나눌 수 있는 값으로 선택합니다. 디스패치는 “대기 행 수”와 “레지스터 255 초과” 두 조건이 서로를 거부하는 방식으로 선택하고, 튜닝은 수용된 빌드 몫을 채울 수 있는 폭만 제공하여 디코드 행이 다시 선택될 수 없도록 합니다. llama-70b 디코드는 2.368마이크로초로 flash-attn의 4.800보다 빠릅니다. 같은 변경에서 드러난 정밀도 문제는 별도로 #2773으로 분리되었습니다. 행이 0에서 멀어지면 fp32 합산의 평균과 분산이 정밀도를 잃고(상수 10000.1인 행은 정규화 후 편향이 아닌 약 0.7), 1e36 근처에서는 제곱합 오버플로로 전부 NaN이 됩니다. 작성자의 제안은 Welford 또는 두 번 스케일링 방식을 사용하는 것입니다.

출처: https://github.com/tile-ai/TileOPs/pull/2768

3.5 TileOPs: 원소별 및 리덕션 네 건, clamp의 bf16 행 5.9% 속도 향상 (#2766 / #2767 / #2769 / #2770)

네 건이 같은 날 병합되었으며 모두 “먼저 근본 원인을 찾고 수정”에 속합니다. clamp(#2767)는 자체 정밀도 내에서 비교하도록 변경되었는데, 근본 원인은 기존 구현이 각 원소를 fp32로 확장하고 각 경계마다 NaN 판정과 선택을 수행하여 원소당 약 5개 명령어가 필요했기 때문입니다. 이제 원소 쌍마다 HMNMX2.NAN 하나로 처리하여 16384 정방 bf16이 0.25704밀리초에서 0.25024로 감소, torch-compile 대비 1.0587배이며, 모든 16비트 모드, NaN, 양/음의 무한대, 부호 있는 0의 경계에서 torch의 fp32 비교와 비트 단위로 일치함을 검증했습니다. 누적 스캔(#2766)은 dtype=float32의 cumsum과 cumprod를 “먼저 확장한 복사본으로 스캔”에서 한 번의 실행으로 목표 정밀도에 직접 쓰도록 변경하여, hidden-state-f32 워크로드에서 전송량이 112MiB에서 48MiB로 감소하고 cumprod는 47.36에서 14.62마이크로초, cumsum은 47.86에서 14.62마이크로초로 감소했습니다. GELU 곱(#2769)은 공유 gelu()를 통해 fp32로 평가하고 한 번만 반올림하도록 변경되었는데, 기존 erf 경로는 NaN 선택 분기, 세 번의 16비트 반올림, 더 긴 의존성 체인을 가졌습니다. 소형 게이트 그리드의 블록 수는 56에서 14로 감소했습니다. 이항 연산(#2770)은 2바이트 피연산자의 스레드당 원소 수를 8에서 16으로 늘려 두 개의 16바이트 로드를 비행 중으로 유지하며, 34개 행의 기본 구성이 이에 따라 변경되어 MinimumFwdOp의 llama-7b 행이 8.42에서 8.00마이크로초로 감소했습니다.

출처: https://github.com/tile-ai/TileOPs/pull/2767

3.6 TileOPs: 호출 텐서의 정렬 판정을 연산자 계층으로 상향, 세 건 마무리 (#2772 / #2778 / #2780)

#2772 은 10-26 15:28 에 병합되었으며, 55개 파일 +373/-224: 실행기는 구현 실행 전에 저장 시작점이 16바이트 경계에 있지 않은 호출 텐서를 정렬된 저장소로 복사하고, 성공 후 기록된 해당 텐서들을 다시 복사하며, 호출자의 텐서를 반환한다. 복사된 텐서가 다른 실인자와 겹치는 경우(그리고 적어도 한쪽이 기록 대상인 경우), 복사가 커널 자체의 검사를 가리게 되므로 실행 전에 바로 거부한다. 주소 없는 저장소(가짜 텐서 또는 읽기를 거부하는 래퍼)는 복사하지 않는다. 커널 측 정렬 및 임시 저장 로직(MoE 그룹 GEMM, 인과 컨볼루션 상태, 융합 QK-norm RoPE 의 qkv 포함)은 전부 삭제되었다. 대가는 핫 패스에서 정렬 호출의 호스트 시간이 1~2퍼센트 포인트 증가하고, MoEPostPermuteFwdOp 가 out 을 가질 때 4.4퍼센트 포인트 증가한다. #2778 은 이 판단을 생성된 옵저버에 접어 넣고 계획 캐시가 복사를 결정하게 하며, 핫 패스는 그 결정만 읽고, 실행기의 기존 별도 주소 순회는 삭제되었다. #2780 은 10-26 에 열렸다: 현재 네 개의 커널만 out 과 입력의 겹침을 검사하고, 나머지 out 을 가진 연산자들은 조용히 경쟁하며, 각 검사마다 data_ptr 을 반복해서 읽는다(MoEPostPermuteFwdOp 는 한 번 호출에 12회 더 읽음). 방향은 목록에서 연산자별로 겹침 허용 여부를 선언하고, 연산자 계층에서 통일적으로 검사하는 것이다.

출처: https://github.com/tile-ai/TileOPs/pull/2772

3.7 TileOPs: 오프셋 저장소 3종 수정, SM90 그룹 GEMM 과 출력 ABI(#2761 / #2765 / #2762)

#2761 은 그룹 GEMM 의 MMA 구현이 시작점에 오프셋이 있는 연속 텐서를 받아들이게 하고, 벡터 접근을 정렬하며, 결과를 오프셋 출력에 다시 쓸 때 식별자와 건드리지 않은 전문가 행을 유지하고, 복사 전에 겹침을 검사한다. 원래 구현은 이런 뷰를 만나면 CUDA 주소 미정렬로 실패했으며, 새로 추가된 16개 회귀가 FP16 과 BF16 각각의 a, b, bias, out 오프셋을 커버하고, 컴팩트 및 마스크 두 가지 레이아웃을 갖춘다. #2765 는 같은 능력을 SM90 의 MoE 그룹 GEMM 에 보강한다: 저장 시작점이 16바이트 경계에 있지 않을 때 a, b 를 복사하고 out 을 임시 저장하며, vector_aligned 는 이제 “기록되는 텐서는 임시 저장하고 시작 후 다시 복사해야 한다”는 규칙을 명시하고, BF16 과 FP16 의 MMA 커널이 이를 사용하도록 변경되며, FP8 커널은 SM90 FP8 커널에서 상속한다. #2762 는 출력 ABI 를 수정한다: TVM FFI 어댑터가 피호출자가 출력을 할당한다고 선언할 때, TileLang 실행체가 직접 할당하고 반환하도록 하며, 다른 어댑터는 기존 할당 경로를 유지한다. 보고된 회귀는 기존 구현이 새 빌드에서 27개 Mamba 케이스 전부 실패했고, 후보 구현은 공개 호출 경로에서 통과한다는 것이다.

출처: https://github.com/tile-ai/TileOPs/pull/2765

3.8 TileOPs 심사 중: FP8 쿼리의 페이지드 프리필을 e4m3 으로 value 곱하기로 변경(#2775)

10-26 17:20 에 열렸고, 10-27 07:00 에도 여전히 업데이트 중이다. GQAPagedFP8WSKernel 은 FP8 쿼리의 value 곱을 e4m3 에서 수행하도록 변경한다: 변환을 담당하는 warp 가 각 value 타일을 K 주순서의 e4m3 피연산자로 전치하고(ldmatrix 전치, 워드 단위 1회 바이트 치환, stmatrix), 각 16 내에서 key 순서를 정렬하여 각 스레드의 점수가 바로 자신이 공급하는 확률이 되게 하며, 레지스터 내 패킹으로 확률을 e4m3 으로 반올림한다. softmax 는 FA3 의 FP8 형태를 따른다: 2의 8제곱 오프셋으로 작은 확률이 e4m3 의 서브노멀에 들어가지 않게 하고, 행 최댓값은 각 더 큰 점수에 따라 이동하며, key 타일은 마지막에서 거꾸로 스캔한다. H200 에서 llama-8b-fp8-prefill 은 369.3 에서 339.8 마이크로초로 감소했고(FA3 대비 1.000 에서 1.087), qwen3-fp8-ragged-odd 는 130.9 에서 123.9 로 감소했다. 저자는 정밀도 허용 오차가 2의 마이너스 4제곱에서 2의 마이너스 3제곱으로 완화되었음을 사실대로 표기했으며, 원인은 e4m3 확률 자체에 2의 마이너스 4제곱 절반 규모의 반올림이 있고 FA3 도 같은 종류의 입력에서 마찬가지로 초과하기 때문이다.

출처: https://github.com/tile-ai/TileOPs/pull/2775

3.9 TileOPs 심사 중: 희소 장시퀀스의 GDN 프리필 분할 스캔(#2779)

10-26 에 열렸다. SM89 에서 희소 장시퀀스 GDN 프리필 그리드가 기존 상태 스캔과 독립 출력 프로그램을 재사용한다: 적용 조건은 조밀하고 key 주순서인 D128 입력, 최소 16384 토큰, 그리고 상태 스캔 CTA 가 SM 수의 4분의 1을 넘지 않는 것이며, 다른 호출은 기존 구현을 유지한다. 분해 방식은 프로그램 캐시 식별자에 기록되고, 꼬리 블록과 초기 상태에 대한 상주 회귀가 보강되었다. 실측(CUDA 13, PyTorch 2.13, TileLang 0.1.15)에서 32768 길이 bf16 은 1.0384 밀리초 대 FLA 의 1.0656, 65536 은 2.1178 대 2.1786 이며, 16384 에서는 여전히 FLA 가 더 빠르다(0.4695 대 0.4409). 저자는 이득이 이 네 개의 value head 조합에서만 성립함을 명확히 밝혔고, 여덟 개와 열여섯 개 value head 의 기존 모델 형태는 이 경로를 선택하지 않으며, 후보가 보편적으로 FLA 보다 빠르지도 않다.

출처: https://github.com/tile-ai/TileOPs/pull/2779

3.10 TileFoundry: 동적 공유 메모리를 커널 실제 뷰에 따라 크기 결정(#271)

10-26에 개설되어 검토 중이며, 12개 파일 +340/-32. 합성된 호스트 진입점은 이전에 모든 TIR 커널에 dynamic_smem=0을 전달했는데, 정수 TensorView 오프셋으로 공유 메모리에 접근하는 커널은 이 때문에 동적 공유 메모리 없이 실행되었고, 48 KiB를 초과할 때도 cudaFuncAttributeMaxDynamicSharedMemorySize를 승격하지 않았다. 작성자는 목표 용량이 아니라 커널이 실제로 도달하는 최대 바이트 수에 맞춰 크기를 정하기로 했으며, 그 이유는 두 가지다. CUDA는 정적 공유 메모리와 동적 공유 메모리를 동일한 블록당 한도로 계산하는데, 해당 런타임은 자체적으로 여러 정적 공유 메모리(ProgramMetaData 또는 공유 할당 텐서)를 보유하므로 최대 용량을 요청하면 바로 시작되지 않는다. 또한 블록당 얼마를 요청하는지는 하나의 SM에 몇 개의 블록이 들어갈 수 있는지도 결정하는데, sm_89에서 99 KiB를 요청하면 512바이트만 사용하더라도 한 블록만 배치할 수 있다. 코드 생성 또한 첫 번째 문 앞에서 smem_base를 한 번 선언하도록 변경되었는데, loop_outputs에서 그것이 원래 첫 번째 루프의 블록 내에 위치하여 두 번째 루프에서는 보이지 않아 nvcc가 미정의 식별자를 보고했기 때문이다. H200과 sm_89에서 두 가지 GPU 테스트가 통과했으며, 패치하지 않으면 전자는 불법 메모리 접근, 후자는 컴파일 실패가 발생한다.

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

3.11 TileFoundry: 전체 창을 피연산자 자체로 전달, sm_89의 GEMM 컴파일 복구 (#273)

10-26에 개설되어 검토 중이며, 5개 파일 +54/-24. schedule finalize가 sm_89 GEMM에 대해 작성한 코드는 컴파일되지 않는다. 하강 과정에서 전체 누산기를 0에서 시작하여 전체 길이를 덮는 창 T.ptr_of(acc[0:0+16, 0:0+8])으로 감싸는데, CUDA 코드 생성은 분할 레지스터 타일의 포인터 슬라이싱에 대해 “일반 소스 레이아웃 필요”를 보고한다. 더 나아가 MMA는 레지스터 분할을 lhs_frag_dst[0:0+16, 0:0+32] 같은 날 슬라이싱으로 읽는데, 역시 처리 분기가 없다. 이러한 창들은 사실상 모두 전체 피연산자다. 수정은 하강 과정에 들어간다. ConvertHIRToTIR이 한 번 창이 해당 피연산자를 덮는지(시작이 정적 0이고 크기가 피연산자 형상과 같음) 판정하고, 창을 구성하는 세 곳이 이 판정을 공유하여 최종적으로 T.ptr_of(acc)를 작성하고 MMA는 분할을 직접 읽는다. TIR 쌍둥이가 있는 49개의 스케줄링 픽스처 중 29개의 finalize 출력이 이에 따라 변경되었고, sm_89 GEMM과 mesh 뷰가 컴파일 복구되었으며, 두 개의 wgmma 출력은 포인터 슬라이싱 거부를 넘어 다음 뷰에서 멈추는데, 그 지점의 레이아웃은 main에서도 동일하게 거부된다.

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

3.12 TileFoundry 신규 결함: 분석 측이 공유 메모리 피크를 중복 계산하고 TIR로 전달하지 않음 (#272)

10-26에 개설되었으며, 작성자는 #271을 검토하던 중 발견했고, 두 가지 격차 모두 동일한 결과를 가리킨다. 분석을 통과한 프로그램이 시작 시 거부되거나, 최종 TIR이 공유 메모리가 얼마나 필요한지 스스로 설명할 수 없다. 첫째, 메모리 분석은 smem을 facts.memory.shared.per_cta(sm_89에서 101376바이트)와 비교하는데, CUDA는 이 한도를 정적 공유 메모리와 동적 공유 메모리에 동시에 적용하며, 런타임 자체가 12바이트의 __shared__ ProgramMetaData를 보유한다. 이 점유는 runtime.cuh에만 기록되어 있고 분석에 알려지지 않아, 정확히 12바이트 이내 차이의 배치가 분석을 통과하고 시작 시 거부된다. 둘째, 분석은 CTA당 공유 메모리 최고 수위를 RegionMemoryMetadata.peaks에 기록하지만, ConvertHIRToTIR은 이를 산출된 TIR에 쓰지 않고, PrimFunction.metadata는 출력도 파싱도 하지 않아 finalize 이후 반드시 재계산해야 하며(#271이 바로 모든 정수 주소의 공유 메모리 뷰를 스캔하여 재계산함), 작성자가 작성한 호스트 진입점이 과소 요청해도 실행 전에 발견되지 않는다.

출처: https://github.com/tile-ai/TileFoundry/issues/272


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

4.1 언어 서비스: 실험적 CUDA 디버깅 출시, cuda-gdb와 Nsight 연결 (10-26)

tilelang-lsp는 10-26 17:11에 첫 디버깅 커밋(78a5761, 제목은 실험적 TileLang CUDA 디버깅 추가)을 푸시했다. 내용은 네 부분이다. 새로운 tlsp.debug 런처를 추가하여 Python 줄 번호 매핑, -G와 -lineinfo 컴파일 추가, 캐시 격리 및 생성 소스 캡처를 수행하고, Nsight의 cuda-gdb를 VS Code에 연결하여 Python 중단점, 시작 전 검사 및 생성된 CUDA 소스 탐색을 지원하며, 단위, 시작 구성 및 선택적 GPU 인수 테스트(대기 중단점, 디바이스 측 단일 단계, 스레드 검사 및 수치 커버)를 보완하고, 중영문 문서와 CI를 동시에 업데이트했다. 이는 해당 저장소가 8월 이후 처음으로 한 실질적 커밋이며, 이전 마지막 업데이트는 8-04의 CI 테스트였다.

출처: https://github.com/tile-ai/tilelang-lsp/commit/78a5761404

4.2 문서 사이트: 하루에 여덟 건 업데이트 (10-26)

tilelang.github.io가 모니터링 기간 내에 여덟 번 푸시했습니다(베이징 시간 07:32, 10:16, 10:21, 16:51, 17:25, 22:42, 익일 00:08, 00:16). 커밋 제목은 일관되게 Update docs로, 메인 저장소 병합에 따라 자동 동기화되는 문서 빌드에 해당합니다. 같은 날 서드파티 로컬라이제이션 저장소(tilelang-docs-l10n)의 푸시 한 건과 중국어 RISC-V 커뮤니티 허브 저장소의 최초 콘텐츠 커밋도 있었는데, 이는 문서 동기화 경로가 정상적으로 작동하고 있음을 보여줍니다.

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

4.3 기여 프로세스 거버넌스 4종: PR 설명 규범, 문서 전용 시 백엔드 테스트 스킵, Agent 제목 확인, 리뷰 봇 노이즈 감소(#3671 / #3672 / #3673 / #3674)

네 건이 같은 날 병합되었으며, 모두 “리뷰 비용”을 둘러싼 것입니다. #3671은 CONTRIBUTING.md에 PR 설명에는 문제, 방안, 트레이드오프가 포함되어야 하며, 증거는 그것이 뒷받침하는 주장 옆에 붙여야 하고, 결함 수정은 전후 동작 또는 회귀 테스트의 의도와 결과를 기술하며, 성능 변경은 비교 가능한 측정과 조건을 제시해야 한다고 명시했습니다. #3672는 문서 전용 PR에 대해 CUDA, ROCm, Metal, Ascend, CuTeDSL의 빌드와 테스트를 건너뛰도록 했는데, 기존 lint 작업의 체크아웃으로 병합 기준 차이를 확인하여 변경 경로가 모두 Markdown, reStructuredText 또는 docs/_static/ 하위일 때 건너뛰는 방식입니다. 화이트리스트는 의도적으로 좁혀서 실행 가능한 예제, 빌드 파일, 워크플로는 여기에 포함되지 않습니다. #3673은 커밋 스킬에 확인 단계를 추가했습니다. 최종 diff, 검증, PR 설명이 준비되면 먼저 사용할 제목 원문을 기여자에게 보여주고, 기여자가 확인하거나 교체하도록 합니다. #3674는 리뷰 봇의 자동 요약과 독스트링 커버리지 검사를 끄고, 공유 리뷰 지침을 추가하여 발견 사항을 현재 diff와 실제 호출 경로에 고정하고, 회귀와 기존 한계를 구분하며, 기존 구현을 먼저 비교한 뒤 수정 방안을 제시하도록 요구했습니다.

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

4.4 커뮤니티 신규 저장소: RISC-V 커뮤니티 허브 저장소, CANN 연산자 저장소, Hexagon 백엔드 저장소, 디버깅 도구 저장소(10-26)

모니터링 기간 내에 등록할 만한 외부 진전이 여러 건 나타났습니다. TongmingLAIC/tilelang-riscv는 10-23에 생성되고 10-26에 푸시되었으며, TileLang의 RISC-V 커뮤니티 허브를 표방하며 컴파일러 백엔드, AI 연산자, 개발 도구, 문서를 집결시키는 것이 목표입니다. 현재는 여전히 빈 골격이지만, 방향성은 국내 RISC-V AI 스택의 요구와 일치합니다. Atlas-Ascend-Fork/ops-tilelang은 CANN 커뮤니티의 Ascend NPU용 TileLang 연산자 저장소로, 커널 구현, 자동화 케이스, 표준 부하, 성능 베이스라인의 통합 관리를 강조하며 10-26에 푸시가 있었습니다. xwhzz/tilelang-hexagon은 Qualcomm Hexagon 백엔드를 다루며, 저장소는 32 MB이고 10-26에 푸시가 있었습니다. superAngGao/tilelang-debugger는 H200에서 소스 기준 위치 파악을 위한 디버깅 워크스페이스입니다. 그 외에 RuneFang/tvm_tilelang_cookbook(37 stars)과 localizethedocs/tilelang-docs-l10n의 정기 업데이트도 있었습니다.

출처: https://github.com/TongmingLAIC/tilelang-riscv

4.5 산업 측: Ascend와 DeepSeek 생태계 격차 축소 보도(10-26)

10-26에 한 매체가 SemiAnalysis의 10-22 보고서를 인용해, DeepSeek와 Ascend의 긴밀한 협력이 엔비디아와의 칩 하드웨어 및 소프트웨어 생태계 격차를 좁히고 있으며, DeepSeek가 2027년 또는 2028년에는 주로 Ascend 기술 스택에 의존할 것으로 전망했다고 전했습니다. 보도는 특히 이 판단이 DeepSeek가 Ascend 생태계를 겨냥한 소프트웨어 도구를 출시한 지 몇 주 후에 나왔으며, 그중에는 엔비디아 플랫폼과 일대일로 대응하는 소프트웨어 모듈 여섯 개가 포함된다고 언급했습니다. 이 소식 자체는 TileLang의 기능 변경과 관련이 없으며, 이번 호에서는 생태계 배경으로만 등록합니다. 이 여섯 모듈 중 고급 언어 컴파일 도구가 바로 tilelang-ascend이며, Ascend 저장소의 당일 인스톨러와 int4 행렬 곱 수정이 바로 같은 경로에 해당합니다.

출처: https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report


5. 동향 관찰

  1. 정확성 결함의 중심이 “성능 회귀”에서 “컴파일 시 보고되지 않고 특정 아키텍처나 저장 영역에서만 오류가 발생”하는 쪽으로 이동했습니다. 당일 메인라인의 세 건이 모두 이 유형입니다. 출력이 잘못된 스트림에 할당되거나(#3649, 기본이 아닌 스트림이면서 현재 스트림에 대기 중인 작업이 있을 때만 손상됨), TMA 캐시 힌트에 템플릿 인자가 하나 부족하거나(#3683, eviction_policy가 있는 일괄 로드에서만 실패), sm8x 희소 메타데이터에 폴백이 없는 경우(#3335, torch 정규 표기만 인식되고 fp16의 레이아웃 규약이 잘못 적용됨)입니다. 이러한 결함의 공통점은 하드웨어 또는 파라미터 전제 조건이 있어 일반적인 CI 범위를 벗어난다는 점이며, 따라서 점점 더 퍼징 테스트 시리즈에 의해 주도되고 있습니다. 이 장에서 당일 종료된 여섯 건 중 다섯 건이 해당 시리즈에서 나왔습니다.

  2. 동기화와 메모리 계획이 주전장이 되다. 당일 드러난 #3680의 ThreadSync 결함 3건, 이에 대응하는 #3682와 #3681 두 건의 수정, tvm 포크 저장소 #88의 증명기 수정, 메인 저장소 #3670의 영역화 kill 재배치, TileFoundry #272의 공유 메모리 피크가 TIR에 반영되지 않는 문제는 모두 같은 사실을 가리킨다. 즉 컴파일러의 동기화와 메모리 판단이 「휴리스틱」에서 「증명 가능」으로 전환되고 있다는 것이다. 이는 #3666이 왜 중요한지도 설명한다. 그것이 내장 TVM을 독립 패키징된 TIRx로 교체했기에, 이후 이러한 증명기 수정이 포크 저장소에 머무르지 않고 정상적인 릴리스 리듬으로 메인라인에 진입할 수 있게 된 것이다.

  3. 백엔드 적응이 「정밀 제어」 단계에 진입하다. Ascend의 T.annotate_buffer_addresses는 UB, L1, L0A, L0B, L0C, SS 각 풀을 바이트 단위로 주소 지정하고 중첩, 정렬, 용량 검증을 수행할 수 있게 한다. Hygon의 레지스터 파이프라인, warp 분기, scale 버퍼 듀얼 뱅크 파티션 역시 「돌아가는 것」에서 「하드웨어 구조에 맞게 배치하는 것」으로 나아가는 것이다. 이러한 역량이 최근 2주에 집중되었다는 것은 국산 백엔드가 이미 가용성 문턱을 넘어 성능 상한 문제를 다루기 시작했음을 의미한다.

  4. 릴리스 엔지니어링은 여전히 약점이다. TileOPs의 첫 출시는 문서, 매니페스트, 패키징, API 안정성, 입력 검증이 모두 준비된 후에도 TestPyPI 시험 실행과 첫 tag에서 수일째 막혀 있으며, PyPI 산출물은 2025-10-09에 머물러 있다. 196개의 스타를 보유하고 당일 14건이 병합된 연산자 라이브러리에게는 1퍼센트포인트를 더 최적화하는 것보다 설치 가능성이 외부 채택 속도를 더 크게 좌우한다. 추적 목록에서 유일하게 완료되지 않은 기술 항목은 MHA 행 0.996이며, 경영진은 이것이 계속해서 tag를 막도록 두어서는 안 된다.

  5. 외부 생태계가 동시에 확장되고 있다. RISC-V 커뮤니티 허브 저장소, CANN 연산자 저장소, Hexagon 백엔드 저장소가 같은 날 등장하거나 업데이트되었으며, 이는 TileLang이 「tile-ai 조직의 프로젝트」에서 여러 플랫폼이 각자 저장소를 만드는 생태계로 변모하고 있음을 보여준다. 국내 독자에게 RISC-V 방향은 특히 주목할 가치가 있다. 이는 이전의 AI 안경, 엣지 추론 하드웨어 노선과 직접적으로 관련되지만, 현재는 여전히 빈 골격에 불과하며 실질적 진전은 향후 실행 가능한 연산자와 백엔드가 나오는지를 보아야 한다.


부록, 자료와 검증 설명

출처 검증표

출처 검증 결과
GitHub tile-ai 조직(29개 저장소) 모니터링 기간 내 7개 저장소 푸시: tilelang, TileOPs, TileOPs-nightly, tilelang-hygon, tilelang-metax, tilelang-lsp, tilelang.github.io
병합 실측값(is:merged) 30건, 건별로 병합 시각이 모니터링 기간 내에 있는지 대조
심사 중 실측값(is:open 업데이트) 49건, 업데이트 시각 순으로 전수 대조
신규 issue 11건, 건별로 본문 확인
종료된 issue 12건, 그중 6건은 퍼징 테스트 시리즈
PyPI(tileops / tilelang) tileops 최신 0.0.1.dev2(2025-10-09), tilelang 최신 0.1.15(2026-09-30)
릴리스 페이지 tag 검증 TileOPs는 tag도 release도 없음, tilelang 최신 v0.1.15(2026-09-30)
tilelang-ascend 저장소 모니터링 기간 내 푸시 없음(마지막 10-23), 다만 PR과 issue 활동은 활발
tilelang-metax 저장소 조직 뷰에 푸시 타임스탬프가 있으나 6개 브랜치 모두 모니터링 기간 내 커밋 없음, 실질 변경 없는 것으로 처리
채택 측(TileKernels / FlashQLA) 모두 모니터링 기간 내 푸시 없음(10-14)
Google News 중영문 다중 쿼리 24시간 내 TileLang 컴포넌트 정확 매칭 없음, 산업 측 보도 1건은 4.5 참조
Hacker News 모니터링 기간 내 프런트엔드 항목 1건, TileLang과 무관하여 제외
arXiv 키워드 검색 결과 모니터링 기간 내 신규 논문 없음, 최근 1편은 2026-07-24
회원 및 채택 측 저장소 커뮤니티 검색 외부 저장소 6곳에서 모니터링 기간 내 푸시 있음, 4.4 참조

전체 출처 목록

/tilelang/pull/3468>

#2775) — https://github.com/tile-ai/TileOPs/pull/2775