A Coding Guide to NVIDIA’s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention
In this tutorial, we discover TileGym GPU programming by constructing a sensible Colab workflow that runs throughout completely different {hardware} situations. We start by probing the out there CUDA surroundings, checking whether or not NVIDIA cuTile runs instantly, and falling again to Triton when customary Colab GPUs lack the required cuTile stack. Through this setup, we study the core tile-programming thought: as a substitute of writing code for one thread at a time, we function on whole information tiles, load them into the kernel, compute on them effectively, and retailer the outcomes again. We use this mannequin to implement vector addition, fused GELU, row-wise softmax, tiled matrix multiplication, and flash consideration, whereas evaluating every end result in opposition to PyTorch for correctness and benchmarking.
CUDA Environment Probe
import os, sys, math, time, textwrap
def rule(t=""):
print("n" + "=" * 78)
if t: print(t)
print("=" * 78)
rule("0. ENVIRONMENT PROBE")
strive:
import torch
besides ImportError:
print("Installing torch ...")
os.system(f"{sys.executable} -m pip set up -q torch")
import torch
HAS_CUDA = torch.cuda.is_available()
DEV = "cuda" if HAS_CUDA else "cpu"
cc = (0, 0)
if HAS_CUDA:
cc = torch.cuda.get_device_capability()
print(f"GPU : {torch.cuda.get_device_name(0)}")
print(f"Compute functionality : {cc[0]}.{cc[1]}")
print(f"Torch CUDA runtime : {torch.model.cuda}")
print(f"Driver / torch : {torch.__version__}")
else:
print("No CUDA GPU discovered. In Colab: Runtime -> Change runtime kind -> GPU (T4).")
print("The tutorial will nonetheless run its correctness math on CPU the place attainable.")
CUTILE_HW_OK = HAS_CUDA and cc[0] >= 8
CUDA_MAJOR = int((torch.model.cuda or "0").cut up(".")[0]) if HAS_CUDA else 0
CUTILE_TOOLKIT_OK = CUDA_MAJOR >= 13
rule("1. ATTEMPTING REAL cuTile (NVIDIA CUDA Tile) BACKEND")
ct = None
CUTILE_READY = False
if CUTILE_HW_OK and CUTILE_TOOLKIT_OK:
strive:
import cuda.tile as ct
CUTILE_READY = True
print("cuda.tile is already importable.")
besides Exception:
print("Installing cuda-tile[tileiras] (this may take some time)...")
os.system(f"{sys.executable} -m pip set up -q 'cuda-tile[tileiras]' cupy-cuda13x")
strive:
import cuda.tile as ct
CUTILE_READY = True
besides Exception as e:
print("cuTile import nonetheless failed:", repr(e))
else:
causes = []
if not HAS_CUDA: causes.append("no CUDA GPU")
if HAS_CUDA and cc[0] < 8: causes.append(f"compute functionality {cc[0]}.{cc[1]} < 8.0 (Turing/T4 unsupported)")
if not CUTILE_TOOLKIT_OK: causes.append(f"CUDA {torch.model.cuda} < 13.1 required by tileiras")
print("Skipping actual cuTile set up as a result of:", "; ".be a part of(causes) + ".")
print("This is predicted on a normal Colab T4 — we fall again to Triton beneath,")
print("which teaches the very same tile-based programming mannequin.")
if CUTILE_READY:
BACKEND = "cutile"
else:
strive:
import triton, triton.language as tl
BACKEND = "triton" if HAS_CUDA else "torch"
besides ImportError:
if HAS_CUDA:
print("Installing triton ...")
os.system(f"{sys.executable} -m pip set up -q triton")
strive:
import triton, triton.language as tl
BACKEND = "triton"
besides Exception:
BACKEND = "torch"
else:
BACKEND = "torch"
rule(f"ACTIVE EXECUTION BACKEND: {BACKEND.higher()}")
print({
"cutile": "Running NVIDIA cuTile kernels in your Ampere+/CUDA13 GPU. Nice {hardware}!",
"triton": "Running Triton tile kernels in your GPU (the usual Colab path).",
"torch": "No usable GPU kernel backend; displaying reference math on CPU solely.",
}[BACKEND])
print(textwrap.dedent("""
------------------------------------------------------------------
SIMT (basic CUDA) | TILE mannequin (cuTile / Triton)
------------------------------------------------------------------
You write code for ONE | You write code for ONE BLOCK that
thread. You compute a world | owns a complete TILE (e.g. 1024 elems
index, bounds-check it, and | or a 128x128 sub-matrix). You load
contact a single aspect. | the tile, do math on the WHOLE tile,
| retailer it. The compiler maps the tile
C[i] = A[i] + B[i] | onto threads / tensor cores for you.
------------------------------------------------------------------
cuTile primitives: ct.bid(0), ct.load(...), ct.retailer(...), a @ b, ct.launch
Triton primitives: tl.program_id, tl.load, tl.retailer, tl.dot, grid[...]
Same thought, two spellings. Below, each kernel is proven in BOTH.
"""))
CUTILE_SOURCE = {
"vector_add": '''
import cuda.tile as ct
@ct.kernel
def vector_add(a, b, c, tile_size: ct.Constant[int]):
pid = ct.bid(0)
a_tile = ct.load(a, index=(pid,), form=(tile_size,))
b_tile = ct.load(b, index=(pid,), form=(tile_size,))
ct.retailer(c, index=(pid,), tile=a_tile + b_tile)
''',
"matmul": '''
import cuda.tile as ct
@ct.kernel
def matmul(A, B, C, Okay: ct.Constant[int],
BM: ct.Constant[int], BN: ct.Constant[int], BK: ct.Constant[int]):
m, n = ct.bid(0), ct.bid(1)
acc = ct.zeros((BM, BN), dtype=ct.float32)
for okay in vary(ct.cdiv(Okay, BK)):
a = ct.load(A, index=(m, okay), form=(BM, BK))
b = ct.load(B, index=(okay, n), form=(BK, BN))
acc = a @ b + acc
ct.retailer(C, index=(m, n), tile=acc)
'''}
We start by organising the surroundings, importing the required libraries, and checking whether or not CUDA is on the market on the present runtime. We examine the GPU capabilities, CUDA model, and PyTorch setup to decide whether or not the actual cuTile backend is usable. We then choose the lively execution backend, clarify the tile programming mannequin, and retailer reference cuTile kernel supply strings for comparability.
Defining Triton Kernels
if BACKEND == "triton":
@triton.jit
def _vadd_kernel(a_ptr, b_ptr, c_ptr, n, BLOCK: tl.constexpr):
pid = tl.program_id(0)
offs = pid * BLOCK + tl.arange(0, BLOCK)
masks = offs < n
a = tl.load(a_ptr + offs, masks=masks)
b = tl.load(b_ptr + offs, masks=masks)
tl.retailer(c_ptr + offs, a + b, masks=masks)
@triton.jit
def _fused_gelu_kernel(x_ptr, w_ptr, b_ptr, o_ptr, n, BLOCK: tl.constexpr):
pid = tl.program_id(0)
offs = pid * BLOCK + tl.arange(0, BLOCK)
masks = offs < n
x = tl.load(x_ptr + offs, masks=masks)
w = tl.load(w_ptr + offs, masks=masks)
b = tl.load(b_ptr + offs, masks=masks)
h = x * w + b
c = 0.7978845608028654
z = c * (h + 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.retailer(o_ptr + offs, g, masks=masks)
@triton.jit
def _softmax_kernel(x_ptr, o_ptr, stride, n_cols, BLOCK: tl.constexpr):
row = tl.program_id(0)
cols = tl.arange(0, BLOCK)
masks = cols < n_cols
ptr = x_ptr + row * stride + cols
x = tl.load(ptr, masks=masks, different=-float("inf"))
x = x - tl.max(x, axis=0)
num = tl.exp(x)
den = tl.sum(num, axis=0)
tl.retailer(o_ptr + row * stride + cols, num / den, masks=masks)
@triton.jit
def _matmul_kernel(A, B, C, M, N, Okay,
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)
for okay in vary(0, Okay, BK):
a = tl.load(a_ptr, masks=offs_k[None, :] < Okay - okay, different=0.0)
b = tl.load(b_ptr, masks=offs_k[:, None] < Okay - okay, different=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] < M) & (offs_n[None, :] < N)
tl.retailer(c_ptr, acc.to(C.dtype.element_ty), masks=cmask)
@triton.jit
def _flash_kernel(Q, Okay, V, O, sqz, skz, svz, soz,
L, D, scale,
BL: tl.constexpr, BD: tl.constexpr):
pid_l = tl.program_id(0)
z = tl.program_id(1)
offs_l = pid_l * BL + tl.arange(0, BL)
offs_d = tl.arange(0, BD)
q_ptr = Q + z * sqz + offs_l[:, None] * D + offs_d[None, :]
q = tl.load(q_ptr, masks=offs_l[:, None] < L, different=0.0)
m_i = tl.full((BL,), -float("inf"), dtype=tl.float32)
l_i = tl.zeros((BL,), dtype=tl.float32)
acc = tl.zeros((BL, BD), dtype=tl.float32)
for begin in vary(0, L, BL):
offs_k = begin + tl.arange(0, BL)
k_ptr = Okay + z * skz + offs_k[:, None] * D + offs_d[None, :]
v_ptr = V + z * svz + offs_k[:, None] * D + offs_d[None, :]
okay = tl.load(k_ptr, masks=offs_k[:, None] < L, different=0.0)
v = tl.load(v_ptr, masks=offs_k[:, None] < L, different=0.0)
s = tl.dot(q, tl.trans(okay)) * scale
s = tl.the place(offs_k[None, :] < L, s, -float("inf"))
m_ij = tl.most(m_i, tl.max(s, axis=1))
p = tl.exp(s - m_ij[:, None])
alpha = tl.exp(m_i - m_ij)
l_i = l_i * alpha + tl.sum(p, axis=1)
acc = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v)
m_i = m_ij
acc = acc / l_i[:, None]
o_ptr = O + z * soz + offs_l[:, None] * D + offs_d[None, :]
tl.retailer(o_ptr, acc.to(O.dtype.element_ty), masks=offs_l[:, None] < L)
def run_vadd(a, b):
c = torch.empty_like(a); n = a.numel()
grid = (triton.cdiv(n, 1024),)
_vadd_kernel[grid](a, b, c, n, BLOCK=1024)
return c
def run_fused_gelu(x, w, b):
o = torch.empty_like(x); n = x.numel()
grid = (triton.cdiv(n, 1024),)
_fused_gelu_kernel[grid](x, w, b, o, n, BLOCK=1024)
return o
def run_softmax(x):
m, ncols = x.form
o = torch.empty_like(x)
BLOCK = triton.next_power_of_2(ncols)
_softmax_kernel[(m,)](x, o, x.stride(0), ncols, BLOCK=BLOCK)
return o
def run_matmul(a, b):
M, Okay = a.form; K2, N = b.form
c = torch.empty((M, N), machine=a.machine, dtype=a.dtype)
BM = BN = 64; BK = 32
grid = (triton.cdiv(M, BM), triton.cdiv(N, BN))
_matmul_kernel[grid](a, b, c, M, N, Okay,
a.stride(0), a.stride(1), b.stride(0), b.stride(1),
c.stride(0), c.stride(1), BM=BM, BN=BN, BK=BK)
return c
def run_flash(q, okay, v):
Z, L, D = q.form
o = torch.empty_like(q)
scale = 1.0 / math.sqrt(D)
BL = 64
grid = (triton.cdiv(L, BL), Z)
_flash_kernel[grid](q, okay, v, o,
q.stride(0), okay.stride(0), v.stride(0), o.stride(0),
L, D, scale, BL=BL, BD=D)
return o
else:
def run_vadd(a, b): return a + b
def run_fused_gelu(x, w, b): return torch.nn.purposeful.gelu(x * w + b, approximate="tanh")
def run_softmax(x): return torch.softmax(x, dim=-1)
def run_matmul(a, b): return a @ b
def run_flash(q, okay, v): return torch.nn.purposeful.scaled_dot_product_attention(q, okay, v)
We outline the Triton implementations for vector addition, fused GELU, row softmax, tiled matrix multiplication, and flash consideration. We specific every operation utilizing tile-level hundreds, computations, reductions, dot merchandise, and shops, in order that the GPU can deal with blocks of information effectively. We additionally present pure PyTorch fallback capabilities so the tutorial nonetheless runs when Triton or a supported GPU backend is unavailable.
Vector Add and GELU
def bench(fn, *a, iters=50, warmup=10):
if HAS_CUDA:
for _ in vary(warmup): fn(*a)
torch.cuda.synchronize()
t0 = time.perf_counter()
for _ in vary(iters): fn(*a)
torch.cuda.synchronize()
return (time.perf_counter() - t0) / iters * 1e3
else:
t0 = time.perf_counter()
for _ in vary(iters): fn(*a)
return (time.perf_counter() - t0) / iters * 1e3
def test(title, received, ref, atol=1e-2, rtol=1e-2):
received = received.float().cpu(); ref = ref.float().cpu()
okay = torch.allclose(received, ref, atol=atol, rtol=rtol)
md = (received - ref).abs().max().merchandise()
print(f" correctness [{name:12s}] : {'PASS' if okay else 'FAIL'} (max abs diff {md:.2e})")
return okay
dtype = torch.float16 if HAS_CUDA else torch.float32
rule("KERNEL 1 — VECTOR ADD (load tile -> add -> retailer tile)")
print("cuTile model of this kernel:n" + CUTILE_SOURCE["vector_add"])
n = 1 << 20
a = torch.randn(n, machine=DEV, dtype=torch.float32)
b = torch.randn(n, machine=DEV, dtype=torch.float32)
test("vector_add", run_vadd(a, b), a + b)
if BACKEND != "torch":
print(f" {BACKEND} time: {bench(run_vadd, a, b):.4f} ms torch: {bench(lambda x,y:x+y, a, b):.4f} ms")
rule("KERNEL 2 — FUSED GELU(x*w + b) (three ops fused into one reminiscence cross)")
x = torch.randn(n, machine=DEV, dtype=torch.float32)
w = torch.randn(n, machine=DEV, dtype=torch.float32)
bb = torch.randn(n, machine=DEV, dtype=torch.float32)
ref = torch.nn.purposeful.gelu(x * w + bb, approximate="tanh")
test("fused_gelu", run_fused_gelu(x, w, bb), ref)
if BACKEND != "torch":
torch_fn = lambda x,w,b: torch.nn.purposeful.gelu(x*w+b, approximate="tanh")
print(f" {BACKEND} (1 cross): {bench(run_fused_gelu, x, w, bb):.4f} ms "
f"torch (3 passes): {bench(torch_fn, x, w, bb):.4f} ms")
We construct benchmarking and correctness-checking utilities that examine every customized kernel in opposition to a PyTorch reference implementation. We then run the vector-addition kernel and confirm that the tile-based output matches the usual PyTorch addition. We additionally take a look at the fused GELU kernel, demonstrating how multiplication, bias addition, and GELU activation are mixed right into a single environment friendly cross.
Softmax and Tiled Matmul
rule("KERNEL 3 — ROW SOFTMAX (tile reductions: max then sum, numerically steady)")
rows, cols = 4096, 1024
x = torch.randn(rows, cols, machine=DEV, dtype=torch.float32)
test("softmax", run_softmax(x), torch.softmax(x, dim=-1))
if BACKEND != "torch":
print(f" {BACKEND} time: {bench(run_softmax, x):.4f} ms "
f"torch: {bench(lambda z: torch.softmax(z,-1), x):.4f} ms")
rule("KERNEL 4 — TILED MATMUL (Okay-loop accumulation -> tensor cores)")
print("cuTile model of this kernel:n" + CUTILE_SOURCE["matmul"])
M = N = Okay = 1024
a = torch.randn(M, Okay, machine=DEV, dtype=dtype)
b = torch.randn(Okay, N, machine=DEV, dtype=dtype)
test("matmul", run_matmul(a, b), a @ b, atol=1e-1, rtol=1e-1)
if BACKEND != "torch":
t = bench(run_matmul, a, b)
flops = 2 * M * N * Okay
print(f" {BACKEND}: {t:.4f} ms ({flops/ (t*1e-3) / 1e12:.2f} TFLOP/s) "
f"torch: {bench(lambda x,y:x@y, a, b):.4f} ms")
We run the row-wise softmax kernel and examine it in opposition to PyTorch’s softmax to confirm numerical correctness. We then carry out tiled matrix multiplication, multiplying matrix blocks and accumulating alongside the Okay dimension. We benchmark these kernels in opposition to PyTorch to observe how tile-based execution performs on the lively backend.
Flash Attention Kernel
rule("KERNEL 5 — FLASH ATTENTION (on-line softmax; the superior capstone)")
Z, L, D = 8, 512, 64
q = torch.randn(Z, L, D, machine=DEV, dtype=dtype)
okay = torch.randn(Z, L, D, machine=DEV, dtype=dtype)
v = torch.randn(Z, L, D, machine=DEV, dtype=dtype)
ref = torch.nn.purposeful.scaled_dot_product_attention(q, okay, v)
test("flash_attn", run_flash(q, okay, v), ref, atol=2e-2, rtol=2e-2)
if BACKEND != "torch":
sdpa = lambda q,okay,v: torch.nn.purposeful.scaled_dot_product_attention(q,okay,v)
print(f" {BACKEND}: {bench(run_flash, q, okay, v):.4f} ms "
f"torch sdpa: {bench(sdpa, q, okay, v):.4f} ms")
rule("DONE")
print(f"""
Summary
-------
Backend that ran : {BACKEND}
What you realized : the tile programming mannequin (whole-tile load/compute/retailer),
fusion, tile reductions, Okay-loop matmul on tensor cores, and
an online-softmax flash-attention kernel.
To run the REAL cuTile kernels proven above you want CUDA 13.1+ and an
Ampere/Ada/Blackwell GPU. On such a machine:
pip set up 'cuda-tile[tileiras]' cupy-cuda13x
pip set up tilegym[tileiras]
Then the strings in CUTILE_SOURCE run as-is through ct.launch(...).
""")
We end with the flash consideration kernel, which applies on-line softmax to compute consideration with out materializing the complete consideration matrix. We examine its output to PyTorch’s scaled dot-product consideration and benchmark runtime efficiency when a GPU backend is on the market. We shut the tutorial by summarizing the backend we used and the principle tile programming ideas we realized.
Conclusion
In conclusion, we understood how tile-based kernels map high-level mathematical operations onto environment friendly GPU execution patterns. We noticed how fusion reduces reminiscence site visitors, how tile reductions stabilize and make softmax environment friendly, how tiled matrix multiplication accumulates over Okay-blocks, and how flash consideration makes use of on-line softmax to keep away from materializing the complete consideration matrix. We additionally gained a path for experimentation: we ran Triton kernels on frequent Colab GPUs at the moment whereas nonetheless seeing how the identical ideas translate to actual cuTile kernels on newer CUDA 13.1+ Ampere, Ada, or Blackwell methods.
Check out the Full Codes with Notebook. Also, be at liberty to observe us on Twitter and don’t overlook to be a part of our 150k+ML SubReddit and Subscribe to our Newsletter. Wait! are you on telegram? now you can join us on telegram as well.
Need to companion with us for selling your GitHub Repo OR Hugging Face Page OR Product Release OR Webinar and so on.? Connect with us
The submit A Coding Guide to NVIDIA’s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention appeared first on MarkTechPost.
