tile-ai

@DeepSeekHarness This is too long to read, but we have an easy repository to see this applied. At TileLang: tile-ai/tvm is a fork of apache/tvm and it's 208 commits ahead of and 814 commits behind. Let's apply the agent and super agent and get these two repositories closer together!

Got another patch upstream! tilelang#3286

Reality check on DeepSeek open-sourcing their inference engine (deepinfer): Why you aren’t running it next week on consumer/workstation gear Dev Twitter (yes I don't like to call it X) is melting down these days over the recent news that DeepSeek open-sourced their inference engine, with half the timeline assuming this means "git clone and serve production DeepSeek locally today". If you are sitting on hardware like 4xRTX PRO 6000 rig and think you are deploying deepinfer next week, you are in for a brutal wake-up call. Serving the weights is not the bottleneck here. Most of us are already running them fine via vLLM or SGLang. The bottleneck is that deepinfer is gated on three low-level kernel libraries that flat-out do not compile on SM120, and no magical cmake flag is going to bail you out. Before you burn an entire weekend debugging build failures, here is the actual reality and the technical path required if you want a real shot at this. Step 1: The Sanity Check (Don’t skip this)Before touching anything else, test whether your target kernel primitive even compiles on your hardware: git clone --depth 1 --branch v0.1.15 https://github.com/tile-ai/tilelang cd tilelang python examples/gemm_sm120/sm120_nvfp4_blockscaled_gemm.py - If it builds clean: You have a viable engineering project. - If it fails: Close your terminal, keep your SGLang/vLLM stack running, and go touch grass. There is zero shame in it. Step 2: Expert Requantization Assuming your kernel build succeeded, the checkpoint will still not load as-is. The native checkpoint scales use UE8M0 across 32x32 blocks. SM120 block-scaled MMA expects NVFP4 with 16-element vectors and E4M3 scaling factors. You must requantize the expert weights to match the instruction set architecture. Once completed, run an empirical perplexity and loss check across standard validation sets before trusting a single token of output. Step 3: Pick Your Distributed Poison (EP vs. TP) Without native NVLink or P2P support on RTX-class boards, you face an architectural trade-off: 1. Expert Parallelism (EP All-to-All): Without hardware P2P, every dispatch and combine pass traverses host RAM over PCIe. This completely eliminates the memory throughput advantages that DeepEP was designed to deliver. 2. Tensor Parallelism (TP4): Cheaper and simpler to stabilize, but you are no longer running their intended architectural topology. You are essentially running standard TP with deepinfer’s Python orchestration layer glued on top. Step 4: Gutting and Replacing the Kernel Layer You cannot run their low-level compute stack verbatim. You have to surgically separate the runtime engine from the vendor-specific kernels: - What you can keep: The high-level Python coordination, C-extension interfaces, batch scheduler, EPLB, prepare_tensor_input, and CUDA graph runners. These components are largely architecture-agnostic. - What you must gut and rewrite: DeepGEMM and FlashMLA need full replacement. DeepSelect is a Top-K kernel, which will cost you at least a full day of implementation and verification. Step 5: Realistic Throughput Expectations (Tokens/sec, Not Vibes) Measure performance strictly in measured decode speed rather than hype: - Host RAM offloading: At batch size 1, expect roughly ~7 tok/s. You will feel the PCIe bus latency on every single cold expert activation. Prefill throughput will look acceptable; decode is where it falls apart. - NVMe/SSD offloading tier: Expect sub-1 tok/s decode performance. The Takeaway Getting this operational is an immense amount of engineering effort just to end up with what is effectively a less mature, bespoke version of vLLM. Even if you possess genuine SM100 enterprise silicon, SM100 only solves the compilation target; it does not give you an off-the-shelf production server. Clone-and-go is the design philosophy of vLLM. DeepSeek’s release is: clone, compile for an hour, and then write the missing KV cache offloading layer yourself. This is not a slight against DeepSeek. The release is an exceptional piece of systems engineering and systems architecture documentation. However, it is an architectural blueprint showing how they operate at massive scale on their specific hardware topology; it is not a turnkey inference binary for local workstations.

So because the TileLang TVM fork is so old one has to imagine workarounds for metal.cooperative_tensor tilelang#3552

So TileLang has a custom fork of Apache TVM tile-ai/tvm and they haven't rebased since May. Interestingly, I looked into this for metal.cooperative_tensor which was added to Apache TVM tvm#19423 by someone from DeepSeek... and the same fix awaits since April tvm#34 ?!

Another TileLang compiler Metal emit bug: tilelang#3542 Interestingly, there's not Metal IR, it just text emit.

Desenvolvimento de Kernels de Alta Performance com TileLang O que é O Tile Language (tile-lang) é uma linguagem de domínio específico (DSL) projetada para simplificar a criação de kernels de alto desempenho para GPU, CPU e aceleradores. Utilizando uma sintaxe baseada em Python integrada a uma infraestrutura de compilador desenvolvida sobre o TVM, o tile-lang permite que você se concentre na produtividade sem perder o controle das otimizações de baixo nível essenciais para a performance de hardware. Para quem serve A ferramenta é voltada para desenvolvedores e engenheiros que criam ou otimizam kernels de alta performance, como GEMM, Dequant GEMM, FlashAttention e LinearAttention. O ecossistema oferece suporte multiplataforma para diferentes arquiteturas de hardware, incluindo compatibilidade com caminhos de hardware da NVIDIA (CUDA), AMD (ROCm), Apple (Metal com suporte a tensor cooperativo), processadores tradicionais (via backend LLVM) e NPUs, como o Huawei Ascend 950. Como começar Para iniciar o desenvolvimento utilizando o TileLang, você pode adotar as seguintes etapas: - Acesse a documentação: Verifique os guias oficiais de instalação para configurar o seu ambiente de desenvolvimento. - Estude os exemplos: O projeto disponibiliza diversos códigos de exemplo práticos, abordando desde operações básicas até implementações complexas de atenção esparsa. - Pratique com desafios: Utilize o "TileLang Puzzles" para aprender os conceitos da linguagem de forma prática. - Configure o ambiente de desenvolvimento: Instale a implementação do TileLang LSP (Language Server Protocol) no seu editor de código para contar com diagnósticos precisos, informações sobre layouts e dicas visuais dos formatos de buffers durante a programação. Instalação pip install tilelang 🔗 Repositório oficial: tile-ai/tilelang Veja em 1 minuto {% embed youtube.com %} 📚 Mais conteúdos do canal: gist.github.com