PyTorch y el equipo ads-ranking de Meta soltaron un kernel de atención específico de Blackwell esta semana — TLX Block Attention — y la historia más grande es la capa debajo de él. **TLX (Triton Language Extensions)** es el puente DSL entre la productividad de Triton y el control nivel CUTLASS crudo en Blackwell, exponiendo los nuevos tensor cores async tcgen05, descriptores TMA y TMEM (Tensor Memory de 256KB-por-SM) como primitivos Triton como `tlx.async_dot`, `tlx.async_descriptor_load`, `tlx.local_trans`, más sincronización mBarrier para pipelines de warp producer-consumer. Repo: github.com/triton-lang/triton-ext. Esta es la capa donde la mayoría de builders escribiendo kernels Blackwell en 2026 vivirán una vez que se estabilice.
El kernel mismo apunta a self-attention sparse de bloques fijos — bloques de 64 tokens, patrón block-diagonal, conocido en tiempo de compilación. Esa forma es específicamente para modelos ads-ranking y recommendation de Meta, no atención LLM. Como el patrón se conoce en compile time, el kernel elimina el loop de iteración multi-tile de Flash Attention, factores de corrección online-softmax, round-trip de logsumexp HBM, y preprocessing Di separado — cada tile Q atiende exactamente a un tile K/V, un solo GEMM, sin corrección necesaria. El forward pass usa 15 warps por CTA en un pipeline especializado (1 load / 1 QK-MMA / 4 softmax / 1 PV-MMA / 8 epilogue); el backward usa 20 warps a través de 7 etapas. TMEM triple-buffered en forward (~169KB / 256KB), double-buffered en backward (~162KB / 256KB). B200, BF16, sparsity=70% — forward 0.98ms vs 1.81ms de Flash Attention v2 (1.85×), backward 2.36ms vs 5.89ms (2.50×), total 2.31×. Precisión numérica supera FA v2 por 53% en max dQ diff.
El rotary backward fusionado es el segundo highlight y el patrón generalizable. Attention backward standalone 1.56ms más rotary backward 4.88ms = 6.44ms no-fusionado; fusionado en un solo kernel que mantiene dV en FP32 dentro de TMEM/registros, aplica el conjugado rotary in-place, luego hace un solo store global BF16 = 1.82ms. **3.54× más rápido.** La lección es portable más allá de workloads ads: cuando tienes valores intermedios FP32 en registros/TMEM, hacer tu math de epilogue en FP32 y store BF16 una vez elimina round-trips a través de memoria global que de otra manera dominan. Ese es el tipo de insight que los builders pueden aplicar a otros kernels de ops fusionados incluso sin TLX o Blackwell.
Lunes por la mañana: este kernel tal como se envió es útil para ti si shippas modelos ad-ranking, recsys o feature-interaction con atención block-diagonal en GPUs B200/B300 — clona facebookresearch/ads_model_kernel_library y benchmarkea. Si eres un builder LLM, el kernel no aplica (causal, sliding-window, y patrones sparse arbitrarios están explícitamente excluidos), pero el DSL TLX mismo es la parte a observar — es cómo se escribirán los kernels Triton Blackwell-aware, y la mayoría de los primitivos arquitectónicos (warp specialization, descriptores TMA, acumuladores TMEM) se generalizan a cualquier forma de atención que tu stack necesite. Límites honestos: solo-Blackwell (sm_100+), sin fallback Ampere/Hopper, head_dim hardcoded a 64 o 128, tamaño de bloque 64 fijo, licencia no declarada en el blog (verifica el repo). Para atención de forma LLM con estas técnicas, el port Blackwell de Flash Attention 3 y sus sucesores será el watch item del próximo trimestre.