PyTorch e o time ads-ranking da Meta soltaram um kernel de atenção específico do Blackwell essa semana — TLX Block Attention — e a história maior é a camada abaixo dele. **TLX (Triton Language Extensions)** é a ponte DSL entre a produtividade do Triton e o controle nível CUTLASS cru no Blackwell, expondo os novos tensor cores async tcgen05, descritores TMA e TMEM (Tensor Memory de 256KB-por-SM) como primitivos Triton como `tlx.async_dot`, `tlx.async_descriptor_load`, `tlx.local_trans`, mais sincronização mBarrier pra pipelines de warp producer-consumer. Repo: github.com/triton-lang/triton-ext. Essa é a camada onde a maioria dos builders escrevendo kernels Blackwell em 2026 vai morar quando se estabilizar.
O kernel em si mira self-attention sparse de blocos fixos — blocos de 64 tokens, padrão block-diagonal, conhecido em tempo de compilação. Essa forma é especificamente pra modelos ads-ranking e recommendation da Meta, não atenção LLM. Como o padrão é conhecido em compile time, o kernel elimina o loop de iteração multi-tile do Flash Attention, fatores de correção online-softmax, round-trip de logsumexp HBM, e preprocessing Di separado — cada tile Q atende a exatamente um tile K/V, um único GEMM, sem correção necessária. O forward pass usa 15 warps por CTA num pipeline especializado (1 load / 1 QK-MMA / 4 softmax / 1 PV-MMA / 8 epilogue); o backward usa 20 warps através de 7 estágios. TMEM triple-buffered no forward (~169KB / 256KB), double-buffered no backward (~162KB / 256KB). B200, BF16, sparsity=70% — forward 0.98ms vs 1.81ms do Flash Attention v2 (1.85×), backward 2.36ms vs 5.89ms (2.50×), total 2.31×. Precisão numérica supera FA v2 em 53% no max dQ diff.
O rotary backward fundido é o segundo destaque e o padrão generalizável. Attention backward standalone 1.56ms mais rotary backward 4.88ms = 6.44ms não-fundido; fundido num kernel único que mantém dV em FP32 dentro de TMEM/registradores, aplica o conjugado rotary in-place, depois faz um único store global BF16 = 1.82ms. **3.54× mais rápido.** A lição é portável além de workloads ads: quando você tem valores intermediários FP32 em registradores/TMEM, fazer tua math de epilogue em FP32 e store BF16 uma vez elimina round-trips através de memória global que de outra forma dominam. Esse é o tipo de insight que builders podem aplicar a outros kernels de ops fundidos mesmo sem TLX ou Blackwell.
Segunda de manhã: esse kernel como entregue é útil pra você se shipa modelos ad-ranking, recsys ou feature-interaction com atenção block-diagonal em GPUs B200/B300 — clone facebookresearch/ads_model_kernel_library e benchmarkeie. Se você é um builder LLM, o kernel não se aplica (causal, sliding-window, e padrões sparse arbitrários são explicitamente excluídos), mas o DSL TLX em si é a parte pra observar — é como kernels Triton Blackwell-aware vão ser escritos, e a maioria dos primitivos arquiteturais (warp specialization, descritores TMA, acumuladores TMEM) se generalizam pra qualquer forma de atenção que tua stack precise. Limites honestos: só-Blackwell (sm_100+), sem fallback Ampere/Hopper, head_dim hardcoded a 64 ou 128, tamanho de bloco 64 fixo, licença não declarada no blog (verifique o repo). Pra atenção formato LLM com essas técnicas, o port Blackwell do Flash Attention 3 e seus sucessores será o watch item do próximo trimestre.