{"id":103698,"date":"2026-07-12T19:12:01","date_gmt":"2026-07-12T19:12:01","guid":{"rendered":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/"},"modified":"2026-07-12T19:12:01","modified_gmt":"2026-07-12T19:12:01","slug":"a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention","status":"publish","type":"post","link":"https:\/\/youzum.net\/es\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/","title":{"rendered":"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention"},"content":{"rendered":"

In this tutorial, we explore TileGym<\/strong><\/a> GPU programming by building a practical Colab workflow that runs across different hardware conditions. We begin by probing the available CUDA environment, checking whether NVIDIA cuTile runs directly, and falling back to Triton when standard Colab GPUs lack the required cuTile stack. Through this setup, we learn the core tile-programming idea: instead of writing code for one thread at a time, we operate on entire data tiles, load them into the kernel, compute on them efficiently, and store the results back. We use this model to implement vector addition, fused GELU, row-wise softmax, tiled matrix multiplication, and flash attention, while comparing each result against PyTorch for correctness and benchmarking. <\/p>\n

CUDA Environment Probe<\/strong><\/h2>\n
\n
\n
\n
\n
<\/div>\n
<\/div>\n
<\/div>\n<\/div>\n
Copy Code<\/span>Copied<\/span>Use a different Browser<\/span><\/a><\/div>\n<\/div>\n
import os, sys, math, time, textwrap\ndef rule(t=\"\"):\n   print(\"n\" + \"=\" * 78)\n   if t: print(t)\n   print(\"=\" * 78)\nrule(\"0. ENVIRONMENT PROBE\")\ntry:\n   import torch\nexcept ImportError:\n   print(\"Installing torch ...\")\n   os.system(f\"{sys.executable} -m pip install -q torch\")\n   import torch\nHAS_CUDA = torch.cuda.is_available()\nDEV = \"cuda\" if HAS_CUDA else \"cpu\"\ncc = (0, 0)\nif HAS_CUDA:\n   cc = torch.cuda.get_device_capability()\n   print(f\"GPU                 : {torch.cuda.get_device_name(0)}\")\n   print(f\"Compute capability  : {cc[0]}.{cc[1]}\")\n   print(f\"Torch CUDA runtime  : {torch.version.cuda}\")\n   print(f\"Driver \/ torch      : {torch.__version__}\")\nelse:\n   print(\"No CUDA GPU found. In Colab: Runtime -> Change runtime type -> GPU (T4).\")\n   print(\"The tutorial will still run its correctness math on CPU where possible.\")\nCUTILE_HW_OK = HAS_CUDA and cc[0] >= 8\nCUDA_MAJOR = int((torch.version.cuda or \"0\").split(\".\")[0]) if HAS_CUDA else 0\nCUTILE_TOOLKIT_OK = CUDA_MAJOR >= 13\nrule(\"1. ATTEMPTING REAL cuTile (NVIDIA CUDA Tile) BACKEND\")\nct = None\nCUTILE_READY = False\nif CUTILE_HW_OK and CUTILE_TOOLKIT_OK:\n   try:\n       import cuda.tile as ct\n       CUTILE_READY = True\n       print(\"cuda.tile is already importable.\")\n   except Exception:\n       print(\"Installing cuda-tile[tileiras] (this can take a while)...\")\n       os.system(f\"{sys.executable} -m pip install -q 'cuda-tile[tileiras]' cupy-cuda13x\")\n       try:\n           import cuda.tile as ct\n           CUTILE_READY = True\n       except Exception as e:\n           print(\"cuTile import still failed:\", repr(e))\nelse:\n   reasons = []\n   if not HAS_CUDA:            reasons.append(\"no CUDA GPU\")\n   if HAS_CUDA and cc[0] < 8:  reasons.append(f\"compute capability {cc[0]}.{cc[1]} < 8.0 (Turing\/T4 unsupported)\")\n   if not CUTILE_TOOLKIT_OK:   reasons.append(f\"CUDA {torch.version.cuda} < 13.1 required by tileiras\")\n   print(\"Skipping real cuTile install because:\", \"; \".join(reasons) + \".\")\n   print(\"This is expected on a standard Colab T4 \u2014 we fall back to Triton below,\")\n   print(\"which teaches the exact same tile-based programming model.\")\nif CUTILE_READY:\n   BACKEND = \"cutile\"\nelse:\n   try:\n       import triton, triton.language as tl\n       BACKEND = \"triton\" if HAS_CUDA else \"torch\"\n   except ImportError:\n       if HAS_CUDA:\n           print(\"Installing triton ...\")\n           os.system(f\"{sys.executable} -m pip install -q triton\")\n           try:\n               import triton, triton.language as tl\n               BACKEND = \"triton\"\n           except Exception:\n               BACKEND = \"torch\"\n       else:\n           BACKEND = \"torch\"\nrule(f\"ACTIVE EXECUTION BACKEND:  {BACKEND.upper()}\")\nprint({\n   \"cutile\": \"Running NVIDIA cuTile kernels on your Ampere+\/CUDA13 GPU. Nice hardware!\",\n   \"triton\": \"Running Triton tile kernels on your GPU (the standard Colab path).\",\n   \"torch\":  \"No usable GPU kernel backend; showing reference math on CPU only.\",\n}[BACKEND])\nprint(textwrap.dedent(\"\"\"\n   ------------------------------------------------------------------\n   SIMT (classic CUDA)          |   TILE model (cuTile \/ Triton)\n   ------------------------------------------------------------------\n   You write code for ONE       |   You write code for ONE BLOCK that\n   thread. You compute a global |   owns a whole TILE (e.g. 1024 elems\n   index, bounds-check it, and   |   or a 128x128 sub-matrix). You load\n   touch a single element.      |   the tile, do math on the WHOLE tile,\n                                |   store it. The compiler maps the tile\n   C[i] = A[i] + B[i]           |   onto threads \/ tensor cores for you.\n   ------------------------------------------------------------------\n   cuTile primitives:  ct.bid(0), ct.load(...), ct.store(...), a @ b, ct.launch\n   Triton primitives:  tl.program_id, tl.load, tl.store, tl.dot, grid[...]\n   Same idea, two spellings. Below, every kernel is shown in BOTH.\n   \"\"\"))\nCUTILE_SOURCE = {\n\"vector_add\": '''\nimport cuda.tile as ct\n@ct.kernel\ndef vector_add(a, b, c, tile_size: ct.Constant[int]):\n   pid    = ct.bid(0)\n   a_tile = ct.load(a, index=(pid,), shape=(tile_size,))\n   b_tile = ct.load(b, index=(pid,), shape=(tile_size,))\n   ct.store(c, index=(pid,), tile=a_tile + b_tile)\n''',\n\"matmul\": '''\nimport cuda.tile as ct\n@ct.kernel\ndef matmul(A, B, C, K: ct.Constant[int],\n          BM: ct.Constant[int], BN: ct.Constant[int], BK: ct.Constant[int]):\n   m, n = ct.bid(0), ct.bid(1)\n   acc  = ct.zeros((BM, BN), dtype=ct.float32)\n   for k in range(ct.cdiv(K, BK)):\n       a = ct.load(A, index=(m, k), shape=(BM, BK))\n       b = ct.load(B, index=(k, n), shape=(BK, BN))\n       acc = a @ b + acc\n   ct.store(C, index=(m, n), tile=acc)\n'''}\n<\/code><\/pre>\n<\/div>\n<\/div>\n
We begin by setting up the environment, importing the required libraries, and checking whether CUDA is available on the current runtime. We inspect the GPU capabilities, CUDA version, and PyTorch setup to determine whether the real cuTile backend is usable. We then select the active execution backend, explain the tile programming model, and store reference cuTile kernel source strings for comparison.<\/p>\n
Defining Triton Kernels<\/strong><\/h2>\n
\n
\n
\n
\n
<\/div>\n
<\/div>\n
<\/div>\n<\/div>\n
Copy Code<\/span>Copied<\/span>Use a different Browser<\/span><\/a><\/div>\n<\/div>\n
if BACKEND == \"triton\":\n   @triton.jit\n   def _vadd_kernel(a_ptr, b_ptr, c_ptr, n, BLOCK: tl.constexpr):\n       pid  = tl.program_id(0)\n       offs = pid * BLOCK + tl.arange(0, BLOCK)\n       mask = offs < n\n       a = tl.load(a_ptr + offs, mask=mask)\n       b = tl.load(b_ptr + offs, mask=mask)\n       tl.store(c_ptr + offs, a + b, mask=mask)\n   @triton.jit\n   def _fused_gelu_kernel(x_ptr, w_ptr, b_ptr, o_ptr, n, BLOCK: tl.constexpr):\n       pid  = tl.program_id(0)\n       offs = pid * BLOCK + tl.arange(0, BLOCK)\n       mask = offs < n\n       x = tl.load(x_ptr + offs, mask=mask)\n       w = tl.load(w_ptr + offs, mask=mask)\n       b = tl.load(b_ptr + offs, mask=mask)\n       h = x * w + b\n       c = 0.7978845608028654\n       z = c * (h + 0.044715 * h * h * h)\n       e = tl.exp(-2.0 * z)\n       tanh = (1.0 - e) \/ (1.0 + e)\n       g = 0.5 * h * (1.0 + tanh)\n       tl.store(o_ptr + offs, g, mask=mask)\n   @triton.jit\n   def _softmax_kernel(x_ptr, o_ptr, stride, n_cols, BLOCK: tl.constexpr):\n       row  = tl.program_id(0)\n       cols = tl.arange(0, BLOCK)\n       mask = cols < n_cols\n       ptr  = x_ptr + row * stride + cols\n       x    = tl.load(ptr, mask=mask, other=-float(\"inf\"))\n       x    = x - tl.max(x, axis=0)\n       num  = tl.exp(x)\n       den  = tl.sum(num, axis=0)\n       tl.store(o_ptr + row * stride + cols, num \/ den, mask=mask)\n   @triton.jit\n   def _matmul_kernel(A, B, C, M, N, K,\n                      sam, sak, sbk, sbn, scm, scn,\n                      BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr):\n       pid_m = tl.program_id(0)\n       pid_n = tl.program_id(1)\n       offs_m = pid_m * BM + tl.arange(0, BM)\n       offs_n = pid_n * BN + tl.arange(0, BN)\n       offs_k = tl.arange(0, BK)\n       a_ptr = A + offs_m[:, None] * sam + offs_k[None, :] * sak\n       b_ptr = B + offs_k[:, None] * sbk + offs_n[None, :] * sbn\n       acc = tl.zeros((BM, BN), dtype=tl.float32)\n       for k in range(0, K, BK):\n           a = tl.load(a_ptr, mask=offs_k[None, :] < K - k, other=0.0)\n           b = tl.load(b_ptr, mask=offs_k[:, None] < K - k, other=0.0)\n           acc += tl.dot(a, b)\n           a_ptr += BK * sak\n           b_ptr += BK * sbk\n       c_ptr = C + offs_m[:, None] * scm + offs_n[None, :] * scn\n       cmask = (offs_m[:, None] < M) & (offs_n[None, :] < N)\n       tl.store(c_ptr, acc.to(C.dtype.element_ty), mask=cmask)\n   @triton.jit\n   def _flash_kernel(Q, K, V, O, sqz, skz, svz, soz,\n                     L, D, scale,\n                     BL: tl.constexpr, BD: tl.constexpr):\n       pid_l = tl.program_id(0)\n       z     = tl.program_id(1)\n       offs_l = pid_l * BL + tl.arange(0, BL)\n       offs_d = tl.arange(0, BD)\n       q_ptr = Q + z * sqz + offs_l[:, None] * D + offs_d[None, :]\n       q = tl.load(q_ptr, mask=offs_l[:, None] < L, other=0.0)\n       m_i = tl.full((BL,), -float(\"inf\"), dtype=tl.float32)\n       l_i = tl.zeros((BL,), dtype=tl.float32)\n       acc = tl.zeros((BL, BD), dtype=tl.float32)\n       for start in range(0, L, BL):\n           offs_k = start + tl.arange(0, BL)\n           k_ptr = K + z * skz + offs_k[:, None] * D + offs_d[None, :]\n           v_ptr = V + z * svz + offs_k[:, None] * D + offs_d[None, :]\n           k = tl.load(k_ptr, mask=offs_k[:, None] < L, other=0.0)\n           v = tl.load(v_ptr, mask=offs_k[:, None] < L, other=0.0)\n           s = tl.dot(q, tl.trans(k)) * scale\n           s = tl.where(offs_k[None, :] < L, s, -float(\"inf\"))\n           m_ij = tl.maximum(m_i, tl.max(s, axis=1))\n           p    = tl.exp(s - m_ij[:, None])\n           alpha = tl.exp(m_i - m_ij)\n           l_i  = l_i * alpha + tl.sum(p, axis=1)\n           acc  = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v)\n           m_i  = m_ij\n       acc = acc \/ l_i[:, None]\n       o_ptr = O + z * soz + offs_l[:, None] * D + offs_d[None, :]\n       tl.store(o_ptr, acc.to(O.dtype.element_ty), mask=offs_l[:, None] < L)\n   def run_vadd(a, b):\n       c = torch.empty_like(a); n = a.numel()\n       grid = (triton.cdiv(n, 1024),)\n       _vadd_kernel[grid](a, b, c, n, BLOCK=1024)\n       return c\n   def run_fused_gelu(x, w, b):\n       o = torch.empty_like(x); n = x.numel()\n       grid = (triton.cdiv(n, 1024),)\n       _fused_gelu_kernel[grid](x, w, b, o, n, BLOCK=1024)\n       return o\n   def run_softmax(x):\n       m, ncols = x.shape\n       o = torch.empty_like(x)\n       BLOCK = triton.next_power_of_2(ncols)\n       _softmax_kernel[(m,)](x, o, x.stride(0), ncols, BLOCK=BLOCK)\n       return o\n   def run_matmul(a, b):\n       M, K = a.shape; K2, N = b.shape\n       c = torch.empty((M, N), device=a.device, dtype=a.dtype)\n       BM = BN = 64; BK = 32\n       grid = (triton.cdiv(M, BM), triton.cdiv(N, BN))\n       _matmul_kernel[grid](a, b, c, M, N, K,\n                            a.stride(0), a.stride(1), b.stride(0), b.stride(1),\n                            c.stride(0), c.stride(1), BM=BM, BN=BN, BK=BK)\n       return c\n   def run_flash(q, k, v):\n       Z, L, D = q.shape\n       o = torch.empty_like(q)\n       scale = 1.0 \/ math.sqrt(D)\n       BL = 64\n       grid = (triton.cdiv(L, BL), Z)\n       _flash_kernel[grid](q, k, v, o,\n                           q.stride(0), k.stride(0), v.stride(0), o.stride(0),\n                           L, D, scale, BL=BL, BD=D)\n       return o\nelse:\n   def run_vadd(a, b):            return a + b\n   def run_fused_gelu(x, w, b):   return torch.nn.functional.gelu(x * w + b, approximate=\"tanh\")\n   def run_softmax(x):            return torch.softmax(x, dim=-1)\n   def run_matmul(a, b):          return a @ b\n   def run_flash(q, k, v):        return torch.nn.functional.scaled_dot_product_attention(q, k, v)\n<\/code><\/pre>\n<\/div>\n<\/div>\n
We define the Triton implementations for vector addition, fused GELU, row softmax, tiled matrix multiplication, and flash attention. We express each operation using tile-level loads, computations, reductions, dot products, and stores, so that the GPU can handle blocks of data efficiently. We also provide pure PyTorch fallback functions so the tutorial still runs when Triton or a supported GPU backend is unavailable.<\/p>\n
Vector Add and GELU<\/strong><\/h2>\n
\n
\n
\n
\n
<\/div>\n
<\/div>\n
<\/div>\n<\/div>\n
Copy Code<\/span>Copied<\/span>Use a different Browser<\/span><\/a><\/div>\n<\/div>\n
def bench(fn, *a, iters=50, warmup=10):\n   if HAS_CUDA:\n       for _ in range(warmup): fn(*a)\n       torch.cuda.synchronize()\n       t0 = time.perf_counter()\n       for _ in range(iters): fn(*a)\n       torch.cuda.synchronize()\n       return (time.perf_counter() - t0) \/ iters * 1e3\n   else:\n       t0 = time.perf_counter()\n       for _ in range(iters): fn(*a)\n       return (time.perf_counter() - t0) \/ iters * 1e3\ndef check(name, got, ref, atol=1e-2, rtol=1e-2):\n   got = got.float().cpu(); ref = ref.float().cpu()\n   ok = torch.allclose(got, ref, atol=atol, rtol=rtol)\n   md = (got - ref).abs().max().item()\n   print(f\"  correctness [{name:12s}] : {'PASS' if ok else 'FAIL'}  (max abs diff {md:.2e})\")\n   return ok\ndtype = torch.float16 if HAS_CUDA else torch.float32\nrule(\"KERNEL 1 \u2014 VECTOR ADD (load tile -> add -> store tile)\")\nprint(\"cuTile version of this kernel:n\" + CUTILE_SOURCE[\"vector_add\"])\nn = 1 << 20\na = torch.randn(n, device=DEV, dtype=torch.float32)\nb = torch.randn(n, device=DEV, dtype=torch.float32)\ncheck(\"vector_add\", run_vadd(a, b), a + b)\nif BACKEND != \"torch\":\n   print(f\"  {BACKEND} time: {bench(run_vadd, a, b):.4f} ms   torch: {bench(lambda x,y:x+y, a, b):.4f} ms\")\nrule(\"KERNEL 2 \u2014 FUSED GELU(x*w + b)  (three ops fused into one memory pass)\")\nx = torch.randn(n, device=DEV, dtype=torch.float32)\nw = torch.randn(n, device=DEV, dtype=torch.float32)\nbb = torch.randn(n, device=DEV, dtype=torch.float32)\nref = torch.nn.functional.gelu(x * w + bb, approximate=\"tanh\")\ncheck(\"fused_gelu\", run_fused_gelu(x, w, bb), ref)\nif BACKEND != \"torch\":\n   torch_fn = lambda x,w,b: torch.nn.functional.gelu(x*w+b, approximate=\"tanh\")\n   print(f\"  {BACKEND} (1 pass): {bench(run_fused_gelu, x, w, bb):.4f} ms   \"\n         f\"torch (3 passes): {bench(torch_fn, x, w, bb):.4f} ms\")\n<\/code><\/pre>\n<\/div>\n<\/div>\n
We build benchmarking and correctness-checking utilities that compare each custom kernel against a PyTorch reference implementation. We then run the vector-addition kernel and verify that the tile-based output matches the standard PyTorch addition. We also test the fused GELU kernel, demonstrating how multiplication, bias addition, and GELU activation are combined into a single efficient pass.<\/p>\n
Softmax and Tiled Matmul<\/strong><\/h2>\n
\n
\n
\n
\n
<\/div>\n
<\/div>\n
<\/div>\n<\/div>\n
Copy Code<\/span>Copied<\/span>Use a different Browser<\/span><\/a><\/div>\n<\/div>\n
rule(\"KERNEL 3 \u2014 ROW SOFTMAX (tile reductions: max then sum, numerically stable)\")\nrows, cols = 4096, 1024\nx = torch.randn(rows, cols, device=DEV, dtype=torch.float32)\ncheck(\"softmax\", run_softmax(x), torch.softmax(x, dim=-1))\nif BACKEND != \"torch\":\n   print(f\"  {BACKEND} time: {bench(run_softmax, x):.4f} ms   \"\n         f\"torch: {bench(lambda z: torch.softmax(z,-1), x):.4f} ms\")\nrule(\"KERNEL 4 \u2014 TILED MATMUL (K-loop accumulation -> tensor cores)\")\nprint(\"cuTile version of this kernel:n\" + CUTILE_SOURCE[\"matmul\"])\nM = N = K = 1024\na = torch.randn(M, K, device=DEV, dtype=dtype)\nb = torch.randn(K, N, device=DEV, dtype=dtype)\ncheck(\"matmul\", run_matmul(a, b), a @ b, atol=1e-1, rtol=1e-1)\nif BACKEND != \"torch\":\n   t = bench(run_matmul, a, b)\n   flops = 2 * M * N * K\n   print(f\"  {BACKEND}: {t:.4f} ms  ({flops\/ (t*1e-3) \/ 1e12:.2f} TFLOP\/s)   \"\n         f\"torch: {bench(lambda x,y:x@y, a, b):.4f} ms\")\n<\/code><\/pre>\n<\/div>\n<\/div>\n
We run the row-wise softmax kernel and compare it against PyTorch\u2019s softmax to verify numerical correctness. We then perform tiled matrix multiplication, multiplying matrix blocks and accumulating along the K dimension. We benchmark these kernels against PyTorch to observe how tile-based execution performs on the active backend.<\/p>\n
Flash Attention Kernel<\/strong><\/h2>\n
\n
\n
\n
\n
<\/div>\n
<\/div>\n
<\/div>\n<\/div>\n
Copy Code<\/span>Copied<\/span>Use a different Browser<\/span><\/a><\/div>\n<\/div>\n
rule(\"KERNEL 5 \u2014 FLASH ATTENTION (online softmax; the advanced capstone)\")\nZ, L, D = 8, 512, 64\nq = torch.randn(Z, L, D, device=DEV, dtype=dtype)\nk = torch.randn(Z, L, D, device=DEV, dtype=dtype)\nv = torch.randn(Z, L, D, device=DEV, dtype=dtype)\nref = torch.nn.functional.scaled_dot_product_attention(q, k, v)\ncheck(\"flash_attn\", run_flash(q, k, v), ref, atol=2e-2, rtol=2e-2)\nif BACKEND != \"torch\":\n   sdpa = lambda q,k,v: torch.nn.functional.scaled_dot_product_attention(q,k,v)\n   print(f\"  {BACKEND}: {bench(run_flash, q, k, v):.4f} ms   \"\n         f\"torch sdpa: {bench(sdpa, q, k, v):.4f} ms\")\nrule(\"DONE\")\nprint(f\"\"\"\nSummary\n-------\nBackend that ran : {BACKEND}\nWhat you learned : the tile programming model (whole-tile load\/compute\/store),\n                  fusion, tile reductions, K-loop matmul on tensor cores, and\n                  an online-softmax flash-attention kernel.\nTo run the REAL cuTile kernels shown above you need CUDA 13.1+ and an\nAmpere\/Ada\/Blackwell GPU. On such a machine:\n   pip install 'cuda-tile[tileiras]' cupy-cuda13x\n   pip install tilegym[tileiras]\nThen the strings in CUTILE_SOURCE run as-is via ct.launch(...).\n\"\"\")\n<\/code><\/pre>\n<\/div>\n<\/div>\n
We finish with the flash attention kernel, which applies online softmax to compute attention without materializing the full attention matrix. We compare its output to PyTorch\u2019s scaled dot-product attention and benchmark runtime performance when a GPU backend is available. We close the tutorial by summarizing the backend we used and the main tile programming concepts we learned.<\/p>\n
Conclusion<\/strong><\/h2>\n
In conclusion, we understood how tile-based kernels map high-level mathematical operations onto efficient GPU execution patterns. We saw how fusion reduces memory traffic, how tile reductions stabilize and make softmax efficient, how tiled matrix multiplication accumulates over K-blocks, and how flash attention uses online softmax to avoid materializing the full attention matrix. We also gained a path for experimentation: we ran Triton kernels on common Colab GPUs today while still seeing how the same concepts translate to real cuTile kernels on newer CUDA 13.1+ Ampere, Ada, or Blackwell systems.<\/p>\n
\n
\n<\/p>
Check out the\u00a0Full Codes with Notebook<\/a><\/strong>.\u00a0<\/strong>Also,\u00a0feel free to follow us on\u00a0Twitter<\/mark><\/a><\/strong>\u00a0and don\u2019t forget to join our\u00a0150k+ML SubReddit<\/a><\/strong>\u00a0and Subscribe to\u00a0our Newsletter<\/a><\/strong>. Wait! are you on telegram?\u00a0now you can join us on telegram as well.<\/a><\/strong><\/p>\n
Need to partner with us for promoting your GitHub Repo OR Hugging Face Page OR Product Release OR Webinar etc.?\u00a0Connect with us<\/mark><\/a><\/strong><\/p>\n
The post A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention<\/a> appeared first on MarkTechPost<\/a>.<\/p>","protected":false},"excerpt":{"rendered":"
In this tutorial, we explore TileGym GPU programming by building a practical Colab workflow that runs across different hardware conditions. We begin by probing the available CUDA environment, checking whether NVIDIA cuTile runs directly, and falling back to Triton when standard Colab GPUs lack the required cuTile stack. Through this setup, we learn the core tile-programming idea: instead of writing code for one thread at a time, we operate on entire data tiles, load them into the kernel, compute on them efficiently, and store the results back. We use this model to implement vector addition, fused GELU, row-wise softmax, tiled matrix multiplication, and flash attention, while comparing each result against PyTorch for correctness and benchmarking. CUDA Environment Probe Copy CodeCopiedUse a different Browser import os, sys, math, time, textwrap def rule(t=””): print(“n” + “=” * 78) if t: print(t) print(“=” * 78) rule(“0. ENVIRONMENT PROBE”) try: import torch except ImportError: print(“Installing torch …”) os.system(f”{sys.executable} -m pip install -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 capability : {cc[0]}.{cc[1]}”) print(f”Torch CUDA runtime : {torch.version.cuda}”) print(f”Driver \/ torch : {torch.__version__}”) else: print(“No CUDA GPU found. In Colab: Runtime -> Change runtime type -> GPU (T4).”) print(“The tutorial will still run its correctness math on CPU where possible.”) CUTILE_HW_OK = HAS_CUDA and cc[0] >= 8 CUDA_MAJOR = int((torch.version.cuda or “0”).split(“.”)[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: try: import cuda.tile as ct CUTILE_READY = True print(“cuda.tile is already importable.”) except Exception: print(“Installing cuda-tile[tileiras] (this can take a while)…”) os.system(f”{sys.executable} -m pip install -q ‘cuda-tile[tileiras]’ cupy-cuda13x”) try: import cuda.tile as ct CUTILE_READY = True except Exception as e: print(“cuTile import still failed:”, repr(e)) else: reasons = [] if not HAS_CUDA: reasons.append(“no CUDA GPU”) if HAS_CUDA and cc[0] < 8: reasons.append(f”compute capability {cc[0]}.{cc[1]} < 8.0 (Turing\/T4 unsupported)”) if not CUTILE_TOOLKIT_OK: reasons.append(f”CUDA {torch.version.cuda} < 13.1 required by tileiras”) print(“Skipping real cuTile install because:”, “; “.join(reasons) + “.”) print(“This is expected on a standard Colab T4 \u2014 we fall back to Triton below,”) print(“which teaches the exact same tile-based programming model.”) if CUTILE_READY: BACKEND = “cutile” else: try: import triton, triton.language as tl BACKEND = “triton” if HAS_CUDA else “torch” except ImportError: if HAS_CUDA: print(“Installing triton …”) os.system(f”{sys.executable} -m pip install -q triton”) try: import triton, triton.language as tl BACKEND = “triton” except Exception: BACKEND = “torch” else: BACKEND = “torch” rule(f”ACTIVE EXECUTION BACKEND: {BACKEND.upper()}”) print({ “cutile”: “Running NVIDIA cuTile kernels on your Ampere+\/CUDA13 GPU. Nice hardware!”, “triton”: “Running Triton tile kernels on your GPU (the standard Colab path).”, “torch”: “No usable GPU kernel backend; showing reference math on CPU only.”, }[BACKEND]) print(textwrap.dedent(“”” —————————————————————— SIMT (classic CUDA) | TILE model (cuTile \/ Triton) —————————————————————— You write code for ONE | You write code for ONE BLOCK that thread. You compute a global | owns a whole TILE (e.g. 1024 elems index, bounds-check it, and | or a 128×128 sub-matrix). You load touch a single element. | the tile, do math on the WHOLE tile, | store 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.store(…), a @ b, ct.launch Triton primitives: tl.program_id, tl.load, tl.store, tl.dot, grid[…] Same idea, two spellings. Below, every kernel is shown 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,), shape=(tile_size,)) b_tile = ct.load(b, index=(pid,), shape=(tile_size,)) ct.store(c, index=(pid,), tile=a_tile + b_tile) ”’, “matmul”: ”’ import cuda.tile as ct @ct.kernel def matmul(A, B, C, K: 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 k in range(ct.cdiv(K, BK)): a = ct.load(A, index=(m, k), shape=(BM, BK)) b = ct.load(B, index=(k, n), shape=(BK, BN)) acc = a @ b + acc ct.store(C, index=(m, n), tile=acc) ”’} We begin by setting up the environment, importing the required libraries, and checking whether CUDA is available on the current runtime. We inspect the GPU capabilities, CUDA version, and PyTorch setup to determine whether the real cuTile backend is usable. We then select the active execution backend, explain the tile programming model, and store reference cuTile kernel source strings for comparison. Defining Triton Kernels Copy CodeCopiedUse a different Browser 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) mask = offs < n a = tl.load(a_ptr + offs, mask=mask) b = tl.load(b_ptr + offs, mask=mask) tl.store(c_ptr + offs, a + b, mask=mask) @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) mask = offs < n x = tl.load(x_ptr + offs, mask=mask) w = tl.load(w_ptr + offs, mask=mask) b = tl.load(b_ptr + offs, mask=mask) 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.store(o_ptr + offs, g, mask=mask) @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) mask = cols < n_cols ptr = x_ptr + row * stride + cols x = tl.load(ptr, mask=mask, other=-float(“inf”)) x = x – tl.max(x, axis=0) num = tl.exp(x) den = tl.sum(num, axis=0) tl.store(o_ptr + row * stride + cols, num \/ den, mask=mask) @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<\/p>","protected":false},"author":2,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"_acf_changed":false,"pmpro_default_level":"","site-sidebar-layout":"default","site-content-layout":"","ast-site-content-layout":"","site-content-style":"default","site-sidebar-style":"default","ast-global-header-display":"","ast-banner-title-visibility":"","ast-main-header-display":"","ast-hfb-above-header-display":"","ast-hfb-below-header-display":"","ast-hfb-mobile-header-display":"","site-post-title":"","ast-breadcrumbs-content":"","ast-featured-img":"","footer-sml-layout":"","theme-transparent-header-meta":"","adv-header-id-meta":"","stick-header-meta":"","header-above-stick-meta":"","header-main-stick-meta":"","header-below-stick-meta":"","astra-migrate-meta-layouts":"default","ast-page-background-enabled":"default","ast-page-background-meta":{"desktop":{"background-color":"var(--ast-global-color-4)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"tablet":{"background-color":"","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"mobile":{"background-color":"","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""}},"ast-content-background-meta":{"desktop":{"background-color":"var(--ast-global-color-5)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"tablet":{"background-color":"var(--ast-global-color-5)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"mobile":{"background-color":"var(--ast-global-color-5)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""}},"_pvb_checkbox_block_on_post":false,"footnotes":""},"categories":[52,5,7,1],"tags":[],"class_list":["post-103698","post","type-post","status-publish","format-standard","hentry","category-ai-club","category-committee","category-news","category-uncategorized","pmpro-has-access"],"acf":[],"yoast_head":"\nA Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention - YouZum<\/title>\n<meta name=\"description\" content=\"\u0e01\u0e34\u0e08\u0e01\u0e23\u0e23\u0e21\u0e40\u0e01\u0e35\u0e48\u0e22\u0e27\u0e01\u0e31\u0e1a\u0e42\u0e14\u0e23\u0e19\" \/>\n<meta name=\"robots\" content=\"index, follow, max-snippet:-1, max-image-preview:large, max-video-preview:-1\" \/>\n<link rel=\"canonical\" href=\"https:\/\/youzum.net\/es\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\" \/>\n<meta property=\"og:locale\" content=\"es_ES\" \/>\n<meta property=\"og:type\" content=\"article\" \/>\n<meta property=\"og:title\" content=\"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention - YouZum\" \/>\n<meta property=\"og:description\" content=\"\u0e01\u0e34\u0e08\u0e01\u0e23\u0e23\u0e21\u0e40\u0e01\u0e35\u0e48\u0e22\u0e27\u0e01\u0e31\u0e1a\u0e42\u0e14\u0e23\u0e19\" \/>\n<meta property=\"og:url\" content=\"https:\/\/youzum.net\/es\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\" \/>\n<meta property=\"og:site_name\" content=\"YouZum\" \/>\n<meta property=\"article:publisher\" content=\"https:\/\/www.facebook.com\/DroneAssociationTH\/\" \/>\n<meta property=\"article:published_time\" content=\"2026-07-12T19:12:01+00:00\" \/>\n<meta name=\"author\" content=\"admin NU\" \/>\n<meta name=\"twitter:card\" content=\"summary_large_image\" \/>\n<meta name=\"twitter:label1\" content=\"Escrito por\" \/>\n\t<meta name=\"twitter:data1\" content=\"admin NU\" \/>\n\t<meta name=\"twitter:label2\" content=\"Tiempo de lectura\" \/>\n\t<meta name=\"twitter:data2\" content=\"14 minutos\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\/\/schema.org\",\"@graph\":[{\"@type\":\"Article\",\"@id\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#article\",\"isPartOf\":{\"@id\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\"},\"author\":{\"name\":\"admin NU\",\"@id\":\"https:\/\/yousum.gpucore.co\/#\/schema\/person\/97fa48242daf3908e4d9a5f26f4a059c\"},\"headline\":\"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention\",\"datePublished\":\"2026-07-12T19:12:01+00:00\",\"mainEntityOfPage\":{\"@id\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\"},\"wordCount\":647,\"commentCount\":0,\"publisher\":{\"@id\":\"https:\/\/yousum.gpucore.co\/#organization\"},\"articleSection\":[\"AI\",\"Committee\",\"News\",\"Uncategorized\"],\"inLanguage\":\"es\",\"potentialAction\":[{\"@type\":\"CommentAction\",\"name\":\"Comment\",\"target\":[\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#respond\"]}]},{\"@type\":\"WebPage\",\"@id\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\",\"url\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\",\"name\":\"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention - YouZum\",\"isPartOf\":{\"@id\":\"https:\/\/yousum.gpucore.co\/#website\"},\"datePublished\":\"2026-07-12T19:12:01+00:00\",\"description\":\"\u0e01\u0e34\u0e08\u0e01\u0e23\u0e23\u0e21\u0e40\u0e01\u0e35\u0e48\u0e22\u0e27\u0e01\u0e31\u0e1a\u0e42\u0e14\u0e23\u0e19\",\"breadcrumb\":{\"@id\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#breadcrumb\"},\"inLanguage\":\"es\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/\"]}]},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"Home\",\"item\":\"https:\/\/youzum.net\/\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention\"}]},{\"@type\":\"WebSite\",\"@id\":\"https:\/\/yousum.gpucore.co\/#website\",\"url\":\"https:\/\/yousum.gpucore.co\/\",\"name\":\"YouSum\",\"description\":\"\",\"publisher\":{\"@id\":\"https:\/\/yousum.gpucore.co\/#organization\"},\"potentialAction\":[{\"@type\":\"SearchAction\",\"target\":{\"@type\":\"EntryPoint\",\"urlTemplate\":\"https:\/\/yousum.gpucore.co\/?s={search_term_string}\"},\"query-input\":{\"@type\":\"PropertyValueSpecification\",\"valueRequired\":true,\"valueName\":\"search_term_string\"}}],\"inLanguage\":\"es\"},{\"@type\":\"Organization\",\"@id\":\"https:\/\/yousum.gpucore.co\/#organization\",\"name\":\"Drone Association Thailand\",\"url\":\"https:\/\/yousum.gpucore.co\/\",\"logo\":{\"@type\":\"ImageObject\",\"inLanguage\":\"es\",\"@id\":\"https:\/\/yousum.gpucore.co\/#\/schema\/logo\/image\/\",\"url\":\"https:\/\/youzum.net\/wp-content\/uploads\/2024\/11\/tranparent-logo.png\",\"contentUrl\":\"https:\/\/youzum.net\/wp-content\/uploads\/2024\/11\/tranparent-logo.png\",\"width\":300,\"height\":300,\"caption\":\"Drone Association Thailand\"},\"image\":{\"@id\":\"https:\/\/yousum.gpucore.co\/#\/schema\/logo\/image\/\"},\"sameAs\":[\"https:\/\/www.facebook.com\/DroneAssociationTH\/\"]},{\"@type\":\"Person\",\"@id\":\"https:\/\/yousum.gpucore.co\/#\/schema\/person\/97fa48242daf3908e4d9a5f26f4a059c\",\"name\":\"admin NU\",\"image\":{\"@type\":\"ImageObject\",\"inLanguage\":\"es\",\"@id\":\"https:\/\/yousum.gpucore.co\/#\/schema\/person\/image\/\",\"url\":\"https:\/\/youzum.net\/wp-content\/uploads\/avatars\/2\/1746849356-bpfull.png\",\"contentUrl\":\"https:\/\/youzum.net\/wp-content\/uploads\/avatars\/2\/1746849356-bpfull.png\",\"caption\":\"admin NU\"},\"url\":\"https:\/\/youzum.net\/es\/members\/adminnu\/\"}]}<\/script>\n<!-- \/ Yoast SEO plugin. -->","yoast_head_json":{"title":"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention - YouZum","description":"\u0e01\u0e34\u0e08\u0e01\u0e23\u0e23\u0e21\u0e40\u0e01\u0e35\u0e48\u0e22\u0e27\u0e01\u0e31\u0e1a\u0e42\u0e14\u0e23\u0e19","robots":{"index":"index","follow":"follow","max-snippet":"max-snippet:-1","max-image-preview":"max-image-preview:large","max-video-preview":"max-video-preview:-1"},"canonical":"https:\/\/youzum.net\/es\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/","og_locale":"es_ES","og_type":"article","og_title":"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention - YouZum","og_description":"\u0e01\u0e34\u0e08\u0e01\u0e23\u0e23\u0e21\u0e40\u0e01\u0e35\u0e48\u0e22\u0e27\u0e01\u0e31\u0e1a\u0e42\u0e14\u0e23\u0e19","og_url":"https:\/\/youzum.net\/es\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/","og_site_name":"YouZum","article_publisher":"https:\/\/www.facebook.com\/DroneAssociationTH\/","article_published_time":"2026-07-12T19:12:01+00:00","author":"admin NU","twitter_card":"summary_large_image","twitter_misc":{"Escrito por":"admin NU","Tiempo de lectura":"14 minutos"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"Article","@id":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#article","isPartOf":{"@id":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/"},"author":{"name":"admin NU","@id":"https:\/\/yousum.gpucore.co\/#\/schema\/person\/97fa48242daf3908e4d9a5f26f4a059c"},"headline":"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention","datePublished":"2026-07-12T19:12:01+00:00","mainEntityOfPage":{"@id":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/"},"wordCount":647,"commentCount":0,"publisher":{"@id":"https:\/\/yousum.gpucore.co\/#organization"},"articleSection":["AI","Committee","News","Uncategorized"],"inLanguage":"es","potentialAction":[{"@type":"CommentAction","name":"Comment","target":["https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#respond"]}]},{"@type":"WebPage","@id":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/","url":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/","name":"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention - YouZum","isPartOf":{"@id":"https:\/\/yousum.gpucore.co\/#website"},"datePublished":"2026-07-12T19:12:01+00:00","description":"\u0e01\u0e34\u0e08\u0e01\u0e23\u0e23\u0e21\u0e40\u0e01\u0e35\u0e48\u0e22\u0e27\u0e01\u0e31\u0e1a\u0e42\u0e14\u0e23\u0e19","breadcrumb":{"@id":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#breadcrumb"},"inLanguage":"es","potentialAction":[{"@type":"ReadAction","target":["https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/"]}]},{"@type":"BreadcrumbList","@id":"https:\/\/youzum.net\/a-coding-guide-to-nvidias-tile-based-gpu-programming-from-cutile-and-triton-kernels-to-flash-attention\/#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https:\/\/youzum.net\/"},{"@type":"ListItem","position":2,"name":"A Coding Guide to NVIDIA\u2019s Tile-Based GPU Programming: From cuTile and Triton Kernels to Flash Attention"}]},{"@type":"WebSite","@id":"https:\/\/yousum.gpucore.co\/#website","url":"https:\/\/yousum.gpucore.co\/","name":"YouSum","description":"","publisher":{"@id":"https:\/\/yousum.gpucore.co\/#organization"},"potentialAction":[{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https:\/\/yousum.gpucore.co\/?s={search_term_string}"},"query-input":{"@type":"PropertyValueSpecification","valueRequired":true,"valueName":"search_term_string"}}],"inLanguage":"es"},{"@type":"Organization","@id":"https:\/\/yousum.gpucore.co\/#organization","name":"Drone Association Thailand","url":"https:\/\/yousum.gpucore.co\/","logo":{"@type":"ImageObject","inLanguage":"es","@id":"https:\/\/yousum.gpucore.co\/#\/schema\/logo\/image\/","url":"https:\/\/youzum.net\/wp-content\/uploads\/2024\/11\/tranparent-logo.png","contentUrl":"https:\/\/youzum.net\/wp-content\/uploads\/2024\/11\/tranparent-logo.png","width":300,"height":300,"caption":"Drone Association Thailand"},"image":{"@id":"https:\/\/yousum.gpucore.co\/#\/schema\/logo\/image\/"},"sameAs":["https:\/\/www.facebook.com\/DroneAssociationTH\/"]},{"@type":"Person","@id":"https:\/\/yousum.gpucore.co\/#\/schema\/person\/97fa48242daf3908e4d9a5f26f4a059c","name":"admin NU","image":{"@type":"ImageObject","inLanguage":"es","@id":"https:\/\/yousum.gpucore.co\/#\/schema\/person\/image\/","url":"https:\/\/youzum.net\/wp-content\/uploads\/avatars\/2\/1746849356-bpfull.png","contentUrl":"https:\/\/youzum.net\/wp-content\/uploads\/avatars\/2\/1746849356-bpfull.png","caption":"admin NU"},"url":"https:\/\/youzum.net\/es\/members\/adminnu\/"}]}},"rttpg_featured_image_url":null,"rttpg_author":{"display_name":"admin NU","author_link":"https:\/\/youzum.net\/es\/members\/adminnu\/"},"rttpg_comment":0,"rttpg_category":"<a href=\"https:\/\/youzum.net\/es\/category\/ai-club\/\" rel=\"category tag\">AI<\/a> <a href=\"https:\/\/youzum.net\/es\/category\/committee\/\" rel=\"category tag\">Committee<\/a> <a href=\"https:\/\/youzum.net\/es\/category\/news\/\" rel=\"category tag\">News<\/a> <a href=\"https:\/\/youzum.net\/es\/category\/uncategorized\/\" rel=\"category tag\">Uncategorized<\/a>","rttpg_excerpt":"In this tutorial, we explore TileGym GPU programming by building a practical Colab workflow that runs across different hardware conditions. We begin by probing the available CUDA environment, checking whether NVIDIA cuTile runs directly, and falling back to Triton when standard Colab GPUs lack the required cuTile stack. Through this setup, we learn the core…","_links":{"self":[{"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/posts\/103698","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/users\/2"}],"replies":[{"embeddable":true,"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/comments?post=103698"}],"version-history":[{"count":0,"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/posts\/103698\/revisions"}],"wp:attachment":[{"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/media?parent=103698"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/categories?post=103698"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/youzum.net\/es\/wp-json\/wp\/v2\/tags?post=103698"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}