update

Author: AndSonder

.gitignore

@@ -31,3 +31,4 @@
 *.out
 *.app
 .DS_Store
+.vscode

README.md

@@ -54,9 +54,9 @@
 + 使用多个 GPU 计算
 + ...(补充中)
 
-### 大师系列 💡
+### Triton 系列 💡
 
-我现在还不知道写啥，毕竟我现在还是菜鸡~~
++ [Triton 编程范式入门](./docs/18_triton/01_triton_programming_paradigms/README.md)
 
 ### LLM 推理技术 🤖
 

docs/18_triton/01_triton_programming_paradigms/README.md

@@ -198,38 +198,11 @@ Triton 的 `mask` 机制则完全不同。`mask = offsets < n_elements` 是一
 
 #### CUDA 的执行方式
 
-```
-配置：blockSize = 256, numBlocks = 40
-
-Grid
-├── Block[0]           Block[1]           ...  Block[39]
-│   ├── Thread[0]      ├── Thread[0]           ├── Thread[0]
-│   │   处理 idx=0     │   处理 idx=256        │   处理 idx=9984
-│   ├── Thread[1]      ├── Thread[1]           ├── Thread[1]
-│   │   处理 idx=1     │   处理 idx=257        │   处理 idx=9985
-│   ├── ...            ├── ...                 ├── ...
-│   └── Thread[255]    └── Thread[255]         └── Thread[255]
-│       处理 idx=255       处理 idx=511             处理 idx=10239 (越界！)
-
-总共启动: 40 × 256 = 10240 个线程
-其中 240 个线程因为 if (idx < 10000) 被过滤掉
-```
+![图 0](images/a476e6f0adb4c9de8c67f9451247ff520d044a1cdc65708a2760b94fa4e803f3.png)
 
 #### Triton 的执行方式
 
-```
-配置：BLOCK_SIZE = 1024, numPrograms = 10
-
-Grid
-├── Program[0]           Program[1]           ...  Program[9]
-│   处理元素             处理元素                   处理元素
-│   [0~1023]            [1024~2047]                [9216~10239]
-│                                                   (其中 10000~10239 被 mask 过滤)
-
-总共启动: 10 个 Program Instance
-每个 Program 处理 1024 个元素（向量化）
-Triton 内部会自动映射到合适的线程配置
-```
+![图 1](images/0d3e9e7c3877312abfbde75dea15acc1e2b671548a26a4e86a114ce6a59a22cd.png)  
 
 从这个例子可以看出，CUDA 启动了 10240 个线程，你需要思考"我是第几号线程"。而 Triton 只启动了 10 个 Program Instance，你要思考的是"我处理哪批数据"。这种抽象层次的提升，让代码更简洁，也更容易理解。
 
@@ -263,26 +236,21 @@ Triton 的 Grid 配置则简单得多，你只需要指定 `BLOCK_SIZE`（每个
 
 ## 四、课后练习
 
-请打开 `01_exercises.py` 完成以下三个练习：练习 1 实现 AXPY 操作（$Z = \alpha \cdot X + Y$），巩固基本的向量化加载和存储；练习 2 测试不同 `BLOCK_SIZE` 的性能影响，理解为什么 Triton 的最优 `BLOCK_SIZE` 比 CUDA 的 `blockDim` 要大；练习 3 实现 1D 卷积，体会如何用向量化方式处理滑动窗口操作。每个练习都包含了测试函数和思考题。
+请打开 [homework.ipynb](https://github.com/PaddleJitLab/CUDATutorial/tree/develop/docs/18_triton/01_triton_programming_paradigms/homework.ipynb) 完成以下练习：练习 1 实现 AXPY 操作（$Z = \alpha \cdot X + Y$），巩固基本的向量化加载和存储；练习 2 实现 1D 卷积，体会如何用向量化方式处理滑动窗口操作。每个练习都包含了测试函数和思考题。
 
 ## 五、常见问题 FAQ
 
-### Q1: Triton 的 BLOCK_SIZE 应该设置多大？
-
-**A**: 如果你有 CUDA 经验，需要注意 CUDA 的经验值在这里不适用。CUDA 的 `blockDim.x` 通常设置为 128/256/512，而 Triton 的 `BLOCK_SIZE` 通常要大得多，一般是 1024/2048/4096。这是因为 Triton 的 `BLOCK_SIZE` 表示的是元素数，而不是线程数。建议从 1024 开始尝试，然后根据性能 profiling 的结果进行调整。影响最优 `BLOCK_SIZE` 的因素包括：寄存器使用量、Shared Memory 大小、以及数据复用程度。
-
-
-### Q2: Triton 内部到底有没有线程？性能会比 CUDA 差吗？
-
-**A**: Triton 内部是有线程的，只是抽象层次更高，不暴露给程序员。Triton 编译器会将你写的向量化代码编译成高效的 PTX（GPU 汇编），最终还是在 GPU 的线程上执行。在性能方面，对于简单算子（如 element-wise 操作），Triton 的性能可以接近手写的优化 CUDA 代码；对于复杂算子（如 Flash Attention），Triton 可以达到优化后 CUDA 的 95% 以上的性能。但在开发效率方面，Triton 远远领先于 CUDA。
+### Q1: Triton 内部到底有没有线程？性能会比 CUDA 差吗？
 
+**A**: 从硬件执行层面看，Triton 代码最终仍然运行在 GPU 的线程和 warp 上，只是 Triton 提供了更高层次的编程抽象，不直接暴露线程和 block 的概念。Triton 编译器会将向量化的程序描述转换为高效的 PTX / SASS，并映射到底层 GPU 执行模型。在性能方面，对于简单算子（如 element-wise 或带宽受限算子），Triton 通常可以达到接近手写 CUDA 的性能；对于高度优化的复杂算子（如 Flash Attention），Triton 在实践中也能达到与优化 CUDA 实现相当、或略低的性能水平。相比之下，Triton 在开发效率和可维护性方面通常具有明显优势。
 
-### Q3: mask 操作会导致性能下降吗？（类似 Warp Divergence）
 
-**A**: Triton 的 `mask` 是向量化的，编译器会生成 predicated instructions（带谓词的指令），不会像 CUDA 的标量 `if` 那样导致严重的 Warp Divergence。性能损失通常可以忽略。从技术细节来看，现代 GPU 支持 predicated execution，每个线程都有独立的 predicate 寄存器。Triton 编译器会自动将 `mask` 映射到这些硬件特性，因此可以在不引入分支的情况下实现条件执行。
+### Q2: mask 操作会导致性能下降吗？（类似 Warp Divergence）
 
+**A**: Triton 的 mask 是向量化语义，编译器通常会将其生成 predicated instructions（带谓词的指令），而不是显式的分支跳转，因此不会像 CUDA 中不当使用 if 那样引入严重的 warp divergence。
+在大多数连续访问、边界检查类场景中，mask 带来的性能开销较小；但如果 mask 覆盖比例很大或访问模式高度稀疏，仍然可能造成一定的算力浪费。总体而言，mask 是 Triton 中推荐且高效的边界处理方式。
 
-### Q4: 什么时候不能用 Triton？
+### Q3: 什么时候不能用 Triton？
 
 **A**: 以下场景建议使用 CUDA：
 1. 需要显式管理 Shared Memory 布局（如手动消除 Bank Conflicts）
@@ -291,7 +259,7 @@ Triton 的 Grid 配置则简单得多，你只需要指定 `BLOCK_SIZE`（每个
 4. 算法严重依赖线程间细粒度通信
 5. 需要与现有 CUDA 代码库深度集成
 
-### Q5: 如何从 CUDA 代码迁移到 Triton？
+### Q4: 如何从 CUDA 代码迁移到 Triton？
 
 **A**: 五步迁移法：
 
@@ -355,4 +323,4 @@ Triton 的 Grid 配置则简单得多，你只需要指定 `BLOCK_SIZE`（每个
 - [OpenAI Triton GitHub](https://github.com/openai/triton)
 - [CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html)
 
-**下一步**：完成所有练习后，进入 **Module 02: 内存与数据搬运**，学习更复杂的内存访问模式！
+**下一步**：完成所有练习后，进入 **02: 内存与数据搬运**，学习更复杂的内存访问模式！

docs/18_triton/01_triton_programming_paradigms/homework.ipynb

@@ -132,145 +132,7 @@
    "source": [
     "---\n",
     "\n",
-    "## 练习 2: 性能测试 - BLOCK_SIZE 的影响\n",
-    "\n",
-    "**目标**：探索不同 `BLOCK_SIZE` 对性能的影响，找出最优配置\n",
-    "\n",
-    "这个练习帮助你理解为什么 Triton 的 `BLOCK_SIZE` 通常比 CUDA 的 `blockDim` 大得多。\n",
-    "\n",
-    "**测试方案**：\n",
-    "- 使用向量加法作为基准测试\n",
-    "- 测试不同的 `BLOCK_SIZE`: [128, 256, 512, 1024, 2048, 4096]\n",
-    "- 测量执行时间和内存带宽"
-   ]
-  },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "metadata": {},
-   "outputs": [],
-   "source": [
-    "# 向量加法 Kernel（用于性能测试）\n",
-    "@triton.jit\n",
-    "def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):\n",
-    "    pid = tl.program_id(axis=0)\n",
-    "    block_start = pid * BLOCK_SIZE\n",
-    "    offsets = block_start + tl.arange(0, BLOCK_SIZE)\n",
-    "    mask = offsets < n_elements\n",
-    "    x = tl.load(x_ptr + offsets, mask=mask)\n",
-    "    y = tl.load(y_ptr + offsets, mask=mask)\n",
-    "    output = x + y\n",
-    "    tl.store(output_ptr + offsets, output, mask=mask)"
-   ]
-  },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "metadata": {},
-   "outputs": [],
-   "source": [
-    "def benchmark_block_size(block_size, x, y, output, warmup=10, repeat=100):\n",
-    "    \"\"\"基准测试单个 BLOCK_SIZE\"\"\"\n",
-    "    n_elements = x.numel()\n",
-    "    grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)\n",
-    "    \n",
-    "    # Warmup\n",
-    "    for _ in range(warmup):\n",
-    "        add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=block_size)\n",
-    "    \n",
-    "    # Timing\n",
-    "    torch.cuda.synchronize()\n",
-    "    start_event = torch.cuda.Event(enable_timing=True)\n",
-    "    end_event = torch.cuda.Event(enable_timing=True)\n",
-    "    \n",
-    "    start_event.record()\n",
-    "    for _ in range(repeat):\n",
-    "        add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=block_size)\n",
-    "    end_event.record()\n",
-    "    \n",
-    "    torch.cuda.synchronize()\n",
-    "    time_ms = start_event.elapsed_time(end_event) / repeat\n",
-    "    \n",
-    "    # 计算带宽 (读 x, 读 y, 写 output)\n",
-    "    total_bytes = 3 * n_elements * 4  # float32 = 4 bytes\n",
-    "    bandwidth_gb_s = total_bytes / (time_ms * 1e-3) / 1e9\n",
-    "    \n",
-    "    return time_ms, bandwidth_gb_s"
-   ]
-  },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "metadata": {},
-   "outputs": [],
-   "source": [
-    "# 运行基准测试\n",
-    "size = 1024 * 1024 * 10  # 10M elements\n",
-    "x = torch.randn(size, device='cuda', dtype=torch.float32)\n",
-    "y = torch.randn(size, device='cuda', dtype=torch.float32)\n",
-    "output = torch.empty_like(x)\n",
-    "\n",
-    "block_sizes = [128, 256, 512, 1024, 2048, 4096]\n",
-    "results = []\n",
-    "\n",
-    "print(f\"{'BLOCK_SIZE':<15} {'Time (ms)':<15} {'Bandwidth (GB/s)':<20}\")\n",
-    "print(\"-\" * 50)\n",
-    "\n",
-    "for bs in block_sizes:\n",
-    "    time_ms, bandwidth = benchmark_block_size(bs, x, y, output)\n",
-    "    results.append((bs, time_ms, bandwidth))\n",
-    "    print(f\"{bs:<15} {time_ms:<15.3f} {bandwidth:<20.2f}\")"
-   ]
-  },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "metadata": {},
-   "outputs": [],
-   "source": [
-    "# 可视化结果\n",
-    "block_sizes_list = [r[0] for r in results]\n",
-    "bandwidths = [r[2] for r in results]\n",
-    "\n",
-    "plt.figure(figsize=(10, 5))\n",
-    "plt.plot(block_sizes_list, bandwidths, marker='o', linewidth=2, markersize=8)\n",
-    "plt.xlabel('BLOCK_SIZE', fontsize=12)\n",
-    "plt.ylabel('Bandwidth (GB/s)', fontsize=12)\n",
-    "plt.title('Triton BLOCK_SIZE vs Memory Bandwidth', fontsize=14)\n",
-    "plt.grid(True, alpha=0.3)\n",
-    "plt.xscale('log', base=2)\n",
-    "plt.xticks(block_sizes_list, block_sizes_list)\n",
-    "\n",
-    "# 标注最佳 BLOCK_SIZE\n",
-    "best_idx = bandwidths.index(max(bandwidths))\n",
-    "plt.axvline(x=block_sizes_list[best_idx], color='r', linestyle='--', alpha=0.5)\n",
-    "plt.text(block_sizes_list[best_idx], max(bandwidths) * 0.95, \n",
-    "         f'Best: {block_sizes_list[best_idx]}', ha='center', fontsize=10, color='r')\n",
-    "\n",
-    "plt.tight_layout()\n",
-    "plt.show()\n",
-    "\n",
-    "print(f\"\\n🏆 最优 BLOCK_SIZE: {block_sizes_list[best_idx]}\")\n",
-    "print(f\"🏆 最高带宽: {max(bandwidths):.2f} GB/s\")"
-   ]
-  },
-  {
-   "cell_type": "markdown",
-   "metadata": {},
-   "source": [
-    "**思考题**：\n",
-    "1. 为什么 `BLOCK_SIZE=128` 性能较差？（提示：GPU 利用率）\n",
-    "2. 为什么 `BLOCK_SIZE=4096` 可能也不理想？（提示：寄存器压力）\n",
-    "3. 对比 CUDA 的 `blockDim.x` 常用值（256），Triton 的最优 `BLOCK_SIZE` 为什么更大？"
-   ]
-  },
-  {
-   "cell_type": "markdown",
-   "metadata": {},
-   "source": [
-    "---\n",
-    "\n",
-    "## 练习 3: 1D 卷积（挑战）\n",
+    "## 练习 2: 1D 卷积（挑战）\n",
     "\n",
     "**目标**：实现简单的 1D 卷积（3-tap box filter）：$Y[i] = X[i-1] + X[i] + X[i+1]$\n",
     "\n",
@@ -304,7 +166,6 @@
     "    # ==================== 在下方编写代码 ====================\n",
     "    \n",
     "    \n",
-    "    \n",
     "    # ========================================================\n",
     "    pass\n",
     "\n",
@@ -351,36 +212,13 @@
     "    print(f\"Torch:  {y_torch[:5].cpu().numpy()}\")"
    ]
   },
-  {
-   "cell_type": "code",
-   "execution_count": null,
-   "metadata": {},
-   "outputs": [],
-   "source": [
-    "# 可视化卷积效果（可选）\n",
-    "size = 100\n",
-    "x = torch.randn(size, device='cuda', dtype=torch.float32)\n",
-    "y = run_conv1d(x)\n",
-    "\n",
-    "plt.figure(figsize=(12, 5))\n",
-    "plt.plot(x.cpu().numpy(), label='Input', alpha=0.7)\n",
-    "plt.plot(y.cpu().numpy(), label='Output (Smoothed)', alpha=0.7, linewidth=2)\n",
-    "plt.xlabel('Index')\n",
-    "plt.ylabel('Value')\n",
-    "plt.title('1D Convolution: Box Filter (3-tap)')\n",
-    "plt.legend()\n",
-    "plt.grid(True, alpha=0.3)\n",
-    "plt.tight_layout()\n",
-    "plt.show()"
-   ]
-  },
   {
    "cell_type": "markdown",
    "metadata": {},
    "source": [
     "**思考题**（高级）：\n",
     "1. 为什么这种方法效率不高？（提示：重复加载）\n",
-    "2. 如何优化？（提示：Shared Memory 或加载更大的块然后切片）"
+    "2. 如何优化？（提示：加载更大的块然后切片）"
    ]
   },
   {
@@ -391,12 +229,61 @@
     "\n",
     "## 总结\n",
     "\n",
-    "完成这三个练习后，你应该：\n",
-    "- 掌握了 Triton kernel 的基本写法\n",
-    "- 理解了 `BLOCK_SIZE` 对性能的重要影响\n",
-    "- 学会了如何处理复杂的内存访问模式\n",
+    "完成这三个练习后，你应该掌握了 Triton kernel 的基本写法\n",
+    "\n",
+    "**下一步**：学习 Triton 的 Shared Memory 和 Block Reduction 操作！\n",
+    "\n",
+    "## 课后答案\n",
+    "\n",
+    "```python\n",
+    "@triton.jit\n",
+    "def axpy_kernel(\n",
+    "    x_ptr, y_ptr, z_ptr,\n",
+    "    n_elements,\n",
+    "    alpha,  # 标量参数\n",
+    "    BLOCK_SIZE: tl.constexpr\n",
+    "):\n",
+    "    \"\"\"\n",
+    "    TODO: 实现 AXPY 操作\n",
+    "    1. 计算 pid 和 offsets\n",
+    "    2. 创建 mask\n",
+    "    3. 加载 x 和 y\n",
+    "    4. 计算 z = alpha * x + y\n",
+    "    5. 存储 z\n",
+    "    \"\"\"\n",
+    "    # ==================== 在下方编写代码 ====================\n",
+    "    pid = tl.program_id(0)\n",
+    "    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n",
+    "    mask = offsets < n_elements\n",
+    "    x = tl.load(x_ptr + offsets, mask=mask, other=0.0)\n",
+    "    y = tl.load(y_ptr + offsets, mask=mask, other=0.0)\n",
+    "    z = alpha * x + y\n",
+    "    tl.store(z_ptr + offsets, z, mask=mask)\n",
+    "    # ========================================================\n",
     "\n",
-    "**下一步**：学习 Triton 的 Shared Memory 和 Block Reduction 操作！"
+    "@triton.jit\n",
+    "def conv1d_kernel(\n",
+    "    x_ptr, y_ptr,\n",
+    "    n_elements,\n",
+    "    BLOCK_SIZE: tl.constexpr\n",
+    "):\n",
+    "    \"\"\"\n",
+    "    TODO: 实现 3-tap 1D 卷积\n",
+    "    Y[i] = X[i-1] + X[i] + X[i+1]\n",
+    "    \"\"\"\n",
+    "    # ==================== 在下方编写代码 ====================\n",
+    "    pid = tl.program_id(0)\n",
+    "    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n",
+    "    mask = offsets < n_elements\n",
+    "    \n",
+    "    x_center = tl.load(x_ptr + offsets, mask=mask, other=0.0)\n",
+    "    x_left = tl.load(x_ptr + offsets - 1, mask=offsets > 0, other=0.0)\n",
+    "    x_right = tl.load(x_ptr + offsets + 1, mask=offsets < n_elements - 1, other=0.0)\n",
+    "    \n",
+    "    y = x_left + x_center + x_right\n",
+    "    tl.store(y_ptr + offsets, y, mask=mask)\n",
+    "    # =========================================================\n",
+    "```"
    ]
   }
  ],

docs/18_triton/01_triton_programming_paradigms/images/0d3e9e7c3877312abfbde75dea15acc1e2b671548a26a4e86a114ce6a59a22cd.png

@@ -49,7 +49,9 @@
 + 使用多个 GPU 计算
 + ...(补充中)
 
-### 大师系列 💡
+### Triton 系列 💡
+
++ [Triton 编程范式入门](/triton/triton_programming_paradigms)
 
 我现在还不知道写啥，毕竟我现在还是菜鸡~~
 

docs/18_triton/01_triton_programming_paradigms/images/a476e6f0adb4c9de8c67f9451247ff520d044a1cdc65708a2760b94fa4e803f3.png

@@ -168,6 +168,16 @@ const sidebars = {
                 }
             ],
         },
+        {
+            type: 'category',
+            label: 'Triton 系列',
+            items: [
+                {
+                    type: 'autogenerated',
+                    dirName: '18_triton/01_triton_programming_paradigms'
+                },
+            ]
+        },
         {
             type: 'category',
             label: 'LLM 推理技术',