Una guía de codificación para la programación de GPU basada en Tile de NVIDIA: desde los núcleos cuTile y Triton hasta Flash Attention

if BACKEND == “triton”: @triton.jit def _vadd_kernel(a_ptr, b_ptr, c_ptr, n, BLOQUE: tl.constexpr): pid = tl.program_id(0) offs = pid * BLOQUE + tl.arange(0, BLOQUE) máscara = apagados < n a = tl.load(a_ptr + apagados, máscara=máscara) b = tl.load(b_ptr + offs, máscara=máscara) tl.store(c_ptr + offs, a + b, máscara=máscara) @triton.jit def _fused_gelu_kernel(x_ptr, w_ptr, b_ptr, o_ptr, n, BLOQUE: tl.constexpr): pid = tl.program_id(0) offs = pid * BLOQUE + + 0.044715 * h * h * h) e = tl.exp(-2.0 * z) tanh = (1.0 - e) / (1.0 + e) g = 0.5 * h * (1.0 + tanh) tl.store(o_ptr + offs, g, máscara=máscara) @triton.jit def _softmax_kernel(x_ptr, o_ptr, zancada, n_cols, BLOQUE: tl.constexpr): fila = tl.program_id(0) cols = tl.arange(0, BLOQUE) máscara = cols < n_cols ptr = x_ptr + fila * zancada + cols x = tl.load(ptr, máscara=máscara, otro=-float("inf")) x = x - tl.max(x, eje=0) num = tl.exp(x) guarida = tl.sum(num, eje=0) tl.store(o_ptr + fila * zancada + columnas, num / guarida, máscara=máscara) @triton.jit def _matmul_kernel(A, B, C, M, N, K, sam, sak, sbk, sbn, scm, scn, BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr): pid_m = tl.program_id(0) pid_n = tl.program_id(1) offs_m = pid_m * BM + tl.arange(0, BM) offs_n = pid_n * BN + tl.arange(0, BN) offs_k = tl.arange(0, BK) a_ptr = A + offs_m[:, None] * sam + offs_k[None, :] * sak b_ptr = B + offs_k[:, None] * sbk + offs_n[None, :] * sbn acc = tl.zeros((BM, BN), dtype=tl.float32) para k en el rango(0, K, BK): a = tl.load(a_ptr, mask=offs_k[None, :] < K - k, otro=0.0) b = tl.load(b_ptr, máscara=offs_k[:, None] < K - k, otro=0.0) acc += tl.dot(a, b) a_ptr += BK * sak b_ptr += BK * sbk c_ptr = C + offs_m[:, None] * scm + offs_n[None, :] * scn cmask = (offs_m[:, None]