Diseño de núcleos de GPU de alto rendimiento con TileLang: Tensor-Core GEMM, Fused Softmax, FlashAttention y Autotuning

@tilelang.jit(out_idx=[-1]) def make_matmul(M: int, N: int, K: int, block_M: int = 128, block_N: int = 128, block_K: int = 32, num_stages: int = 3, subprocesos: int = 128, use_swizzle: bool = False, dtype: str = “float16”, accum_dtype: str = “float”): @T.prim_func def main(A: T.Tensor((M, K), dtype), B: T.Tensor((K, N), dtype), C: T.Tensor((M, N), dtype)): con T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=threads) as (bx, by): A_shared = T.alloc_shared((block_M, block_K), dtype) B_shared = T.alloc_shared((block_K, block_N), dtype) C_local = T.alloc_fragment((block_M, block_N), accum_dtype) if use_swizzle: T.use_swizzle(panel_size=10, enable=True) T.clear(C_local) para ko in T.Pipelined(T.ceildiv(K, block_K), num_stages=num_stages): T.copia(A[by * block_M, ko * block_K]A_compartido) T.copiar(B[ko * block_K, bx * block_N]B_compartido) T.gemm(A_compartido, B_compartido, C_local) T.copia(C_local, C[by * block_M, bx * block_N]) return main def smem_bytes(bloque_M, bloque_N, bloque_K, etapas, tamaño de elemento=2): return (bloque_M * bloque_K + bloque_K * bloque_N) * tamaño de elemento * etapas def sección_2(): banner(“2. JERARQUÍA DE MEMORIA – GEMM de núcleo tensor en mosaico”) M = N = K = 2048 bm, bn, bk, st = 128, 128, 32, DEFAULT_STAGES mientras smem_bytes(bm, bn, bk, st) > SMEM_CAP y st > 1: st -= 1 print(f” config: {bm}x{bn}x{bk}, num_stages={st}, ” f”smem={smem_bytes(bm,bn,bk,st)/1024:.0f} KB”) kernel = make_matmul(M, N, K, bm, bn, bk, num_stages=st, threads=128) a = torch.randn(M, K, dispositivo=DEV, dtype=torch.float16) b = torch.randn(K, N, dispositivo=DEV, dtype=torch.float16) c = kernel(a, b) check(c, a @ b, “matmul 2048^3″) flops = 2 * M * N * K ms = banco(lambda: kernel(a, b)) ms_ref = banco(lambda: a @ b) print(f” Tilelang : {ms:7.3f} ms -> {flops/(ms*1e-3)/1e12:6.2f} TFLOP/s”) print(f” cuBLAS : {ms_ref:7.3f} ms -> {flops/(ms_ref*1e-3)/1e12:6.2f} TFLOP/s”) print(f” ratio: {ms_ref/ms*100:5.1f}% de cuBLAS de ~20 líneas de Python”) src = kernel.get_kernel_source() para aguja en (“mma.sync”, “wgmma”, “ldmatrix”, “cp.async”, “tl::gemm”): if aguja en src: print(f” emitido: {aguja}”) return kernel def sección_3(): banner(“3. KNOBS — barriendo el programa a mano”) M = N = K = 2048 a = torch.randn(M, K, dispositivo=DEV, dtype=torch.float16) b = torch.randn(K, N, dispositivo=DEV, dtype=torch.float16) fracasos = 2 * M * N * K candidatos = [
(64, 64, 32, 2, 128, False),
(128, 128, 32, 2, 128, False),
(128, 128, 32, 2, 128, True),
(128, 128, 32, 3, 128, False),
(128, 128, 64, 2, 256, False),
(128, 256, 32, 2, 256, False),
]
print(f” {’tile’:>16} {‘stg’:>4} {‘thr’:>4} {‘swz’:>4} {‘smem’:>7} ” f”{‘ms’:>8} {‘TFLOP/s’:>9}”) resultados = []
para bm, bn, bk, st, thr, swz en candidatos: need = smem_bytes(bm, bn, bk, st) tag = f”{bm}x{bn}x{bk}” si es necesario > SMEM_CAP: print(f” {tag:>16} {st:>4} {thr:>4} {str(swz):>4} {necesidad//1024:>5}KB ” f” SALTADO (sobre el presupuesto de smem)”) continúe intente: k = make_matmul(M, N, K, bm, bn, bk, num_stages=st, threads=thr, use_swizzle=swz) c = k(a, b) ok = (c – (a @ b)).float().norm() / (a @ b).float().norm() < 2e-2 ms = banco(lambda: k(a, b), calentamiento=5, rep=20) resultados.append((ms, tag, st, thr, swz)) print(f" {tag:>16} {st:>4} {thr:>4} {str(swz):>4} {necesidad//1024:>5}KB ” f”{ms:>8.3f} {flops/(ms*1e-3)/1e12:>9.2f}” f”{” if ok else ‘ <- NUMÉRICAMENTE MAL'}") excepto excepción como e: print(f" {tag:>16} {st:>4} {thr:>4} {str(swz):>4} – ” f”falló: {type(e).__name__}”) si resultados: mejor = min(resultados) print(f”\n ganador: {mejor[1]}, etapas={mejor[2]}, hilos={mejor[3]}, ” f”swizzle={mejor[4]} ({mejor[0]:.3f} ms)”) print(” Conclusión: el mejor cronograma depende del arco y la forma, que es”) print(” exactamente por qué existe la sección 7.”)