2626
2727这是一个经典的 2D 操作,能很好地展示多维地址计算。在内存中,矩阵是按** 行优先** 顺序存储的。例如,一个 3×4 的矩阵:
2828
29- ```
29+ ``` plain
3030A = [[1, 2, 3, 4],
3131 [5, 6, 7, 8],
3232 [9, 10, 11, 12]]
@@ -35,7 +35,7 @@ A = [[1, 2, 3, 4],
3535在内存中的实际布局是:` [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] `
3636
3737如果要访问 ` A[i][j] ` (第 i 行,第 j 列),内存地址是:
38- ```
38+ ``` plain
3939物理地址 = base_ptr + i * N + j
4040```
4141
@@ -156,7 +156,7 @@ for (int i = 0; i < 64; i++) {
156156
157157理解了 2D 地址计算后,咱们来看看一个对性能影响巨大的因素:** 内存连续性** 。
158158
159- ### 2.1 什么是内存连续性?
159+ ### 2.1 什么是内存连续性
160160
161161在 GPU 编程中,有一个非常重要的性能优化原则:** Memory Coalescing** 。
162162
@@ -210,7 +210,7 @@ def vector_add_strided(x_ptr, y_ptr, out_ptr, n_elements, stride, BLOCK_SIZE: tl
210210
211211** 性能测试结果** :
212212
213- ```
213+ ``` plain
214214Size | Stride | Time (ms) | Slowdown
215215--------|--------|-----------|----------
2162165242880 | 1 | 0.02 | 1.00x
@@ -228,7 +228,7 @@ Size | Stride | Time (ms) | Slowdown
228228GPU 的内存访问是以** 事务为单位** 的。一个 warp(32 个线程)访问连续的 32 个 float(128 字节)时,可以合并为** 单个内存事务** 。
229229
230230但当 stride=8 时,这 32 个线程访问的是间隔 8 个元素的数据:
231- ```
231+ ``` plain
232232线程 0: addr[0]
233233线程 1: addr[8]
234234线程 2: addr[16]
@@ -473,7 +473,7 @@ cache hints 的效果因 GPU 架构和数据访问模式而异,需要实际测
473473
474474前面我们学习了如何手动计算 2D/多维张量的地址。但实际上 Triton 提供了一个更强大的工具:` tl.make_block_ptr ` 。
475475
476- ### 4.1 什么是 Block Pointer?
476+ ### 4.1 什么是 Block Pointer
477477
478478` tl.make_block_ptr ` 是一个专门用于处理多维张量访问的高级工具。它可以:
479479
@@ -613,7 +613,7 @@ tl.store(c_block_ptr, result, boundary_check=(0, 1)) # 正确
613613
614614:::
615615
616- ### 4.6 什么时候使用 Block Pointer?
616+ ### 4.6 什么时候使用 Block Pointer
617617
618618| 场景 | 推荐方式 |
619619| ------| ---------|
0 commit comments