[Doc] Fix some typo of reduce optimize

docs/09_optimize_reduce/02_bank_conflict/README.md

@@ -4,27 +4,27 @@
 
 ## 1. Bank Conflict
 
-Bank conflict 是指当多个线程同时访问同一个 bank 的时候，会出现 bank conflict。bank 是共享内存的最小单元，每个 bank 可以同时为一个线程提供数据。如果多个线程同时访问同一个 bank，那么它们的访问就会被串行化，从而降低性能。为了避免或减少访存冲突，可以使用一些技巧，如改变数据布局、使用 padding、使用 shuffle 指令等。
+> 为了获得高的内存带宽，共享内存在物理上被分为 32 个同样宽度的、能被同时访问的内存 bank。共享内存中每连续的 128 字节的内容分摊到 32 个 bank 的同一层中。bank 是共享内存的最小单元。
 
-同一个 Block 的线程会共享一块共享内存，共享内存在逻辑上被分为 32 个 bank。当同一个 warp 中的线程访问同一个 bank 的时候，会出现 bank conflict。在最坏的情况下，即一个 warp 中的所有线程访问了相同 bank 的 32 个不同地址的话，那么这 32 个访问操作将会全部被序列化，大大降低了内存带宽。在一个 warp 内对同一个 bank 中的 n 个地址同时访问将导致 n 次内存事务，称为发生了 n 路 bank conflict。需要注意的是，不同 warp 中的线程之间不存在 bank conflict。
+同一个 Block 的线程会共享一块共享内存，Bank conflict 是指一个 warp 内的多个线程同时访问同一个 bank 的不同地址，那么它们的访问就会被串行化，从而降低性能。在最坏的情况下，即一个 warp 中的所有线程访问了相同 bank 的 32 个不同地址的话，那么这 32 个访问操作将会全部被序列化，大大降低了内存带宽。在一个 warp 内对同一个 bank 中的 n 个地址同时访问将导致 n 次内存事务，称为发生了 n 路 bank conflict。需要注意的是，不同 warp 中的线程之间不存在 bank conflict。为了避免或减少访存冲突，可以使用一些技巧，如改变数据布局、使用 padding、使用 shuffle 指令等。
 
 接下来让我们来分析上一个 Kernel 中的 bank conflict。上一个 Kernel 中，我们使用了交叉寻址的方式，使得连续的线程访问的地址不再连续。这样做的目的是为了避免 warp divergence。但是这样做也会导致 bank conflict。
 
 我们以 0 号 warp 为例。在第一次迭代中，0 号线程需要加载 shared memory 的 0 号和 1 号地址，然后写回 0 号地址。同时，0 号 warp 的 16 号线程需要加载 shared memory 的 32 和 33 号地址，并写回 32 号地址。因此，在一个 warp 内同时访问了一个 bank 的不同内存地址，导致发生了 2 路的 Bank Conflict，如下图所示：
 
-![picture 1](images/ef322be7c3e5b6b9be69d2b90e88083f50569a58a97129f348e483b946ab4edf.png)  
+![picture 1](images/ef322be7c3e5b6b9be69d2b90e88083f50569a58a97129f348e483b946ab4edf.png)
 
-类似地，在第二次迭代过程中，0 号 warp 的 0 号线程会加载 0 号和 2 号地址并写回 0 号地址。然后，0 号 warp 的 8 号线程需要加载 shared memory 的 32 号和 34 号地址（228=32，32+2=34），并写回 32 号线程。此时，16 号线程会加载 64 号和 68 号地址，24 号线程会加载 96 号和 100 号地址。由于 0 号、32 号、64 号、96 号地址都在一个 bank 中，产生了 4 路的 Bank Conflict。这样以此类推，下一次迭代会产生 8 路的 Bank Conflict，使得整个 Kernel 一直受到 Bank Conflict 的影响。
+类似地，在第二次迭代过程中，0 号 warp 的 0 号线程会加载 0 号和 2 号地址并写回 0 号地址。然后，0 号 warp 的 8 号线程需要加载 shared memory 的 32 号和 34 号地址（4*8=32，32+2=34），并写回 32 号线程。此时，16 号线程会加载 64 号和 66 号地址，24 号线程会加载 96 号和 98 号地址。由于 0 号、32 号、64 号、96 号地址都在一个 bank 中，产生了 4 路的 Bank Conflict。这样以此类推，下一次迭代会产生 8 路的 Bank Conflict，使得整个 Kernel 一直受到 Bank Conflict 的影响。
 
 我们可以使用 nvprof 来查看 bank conflict 的情况。
 
 ```bash
 nvprof --events shared_st_bank_conflict ./reduce_interleaved_addressing
 ```
 
-| Invocations | Event Name               | Min      | Max      | Avg      | Total    |
-|-------------|--------------------------|----------|----------|----------|----------|
-| 1           | shared_st_bank_conflict | 4587520  | 4587520  | 4587520  | 4587520  |
+| Invocations | Event Name              | Min     | Max     | Avg     | Total   |
+| ----------- | ----------------------- | ------- | ------- | ------- | ------- |
+| 1           | shared_st_bank_conflict | 4587520 | 4587520 | 4587520 | 4587520 |
 
 
 如果你的设备不支持 nvprof，你可以使用 nsight-compute 的命令行工具 ncu 来查看 bank conflict 的情况。
@@ -44,13 +44,13 @@ sudo ncu --metrics      l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum
 
 每个 Block 里面有一半的线程是需要加载数据的，这些加载数据的线程就可能会发生 bank conflict。我们可以让这些线程访问不同的 bank。前面我们已经说过了为了防止线程束分化，所以这一半加载数据的线程就是 0-127 号线程。这 128 个线程可以分成 4 个 warp，每个 warp 有 32 个线程。我们可以让这 4 个 warp 分别访问 4 个不同的 bank。每个 warp 访问一个 bank。仔细看看上面的布局图，我们一共有 8 行，正好可以分成 4 个 warp，每个 warp 有 2 行。结合交叉寻址的方式，我们可以让每个 warp 访问一个 bank。这样就可以避免 bank conflict 了。我用不同颜色的线表示了不同的 warp，如下图所示：
 
-![picture 4](images/e69b477993846936b270e82a37615c00424010cd8003f429354aa27325c96f57.png)  
+![picture 4](images/e69b477993846936b270e82a37615c00424010cd8003f429354aa27325c96f57.png)
 
 0 号 warp 的 0 号线程访问 0 号 bank 的 0 号地址和 128 号地址，1 号 warp 的 0 号线程访问 1 号 bank 的 32 号地址和 128 号地址，2 号 warp 的 0 号线程访问 2 号 bank 的 64 号地址和 160 号地址，以此类推。
 
 整体过程如下图所示：
 
-![picture 6](images/0f65c7d9e911014e31ddd84c583dea859ba24ebd48715c2680eb604e7ebb9a2b.png)  
+![picture 6](images/0f65c7d9e911014e31ddd84c583dea859ba24ebd48715c2680eb604e7ebb9a2b.png)
 
 > [!NOTE]
 > 图里面的 block size 是 16 而不是 256，这是为了方便说明。实际上，我们的 block size 是 256。
@@ -79,17 +79,17 @@ for(int s=blockDim.x/2; s>0; s >>= 1) {
 
 修改后我们可以再次运行 nvprof 来查看 bank conflict 的情况，输出如下所示：
 
-| Invocations | Event Name               | Min      | Max      | Avg      | Total    |
-|-------------|--------------------------|----------|----------|----------|----------|
-| 1           | shared_st_bank_conflict | 0  | 0  | 0  | 0  |
+| Invocations | Event Name              | Min | Max | Avg | Total |
+| ----------- | ----------------------- | --- | --- | --- | ----- |
+| 1           | shared_st_bank_conflict | 0   | 0   | 0   | 0     |
 
 性能和带宽的测试情况如下：
 
-| 优化手段 | 运行时间(us) | 带宽 | 加速比 |
-| --- | --- | --- | --- |
-| Baseline | 3118.4 | 42.503GB/s | ~ |
-| 交错寻址 | 1904.4 | 73.522GB/s | 1.64 |
-| 解决 bank conflict | 1475.2 | 97.536GB/s | 2.29 |
+| 优化手段           | 运行时间(us) | 带宽       | 加速比 |
+| ------------------ | ------------ | ---------- | ------ |
+| Baseline           | 3118.4       | 42.503GB/s | ~      |
+| 交错寻址           | 1904.4       | 73.522GB/s | 1.64   |
+| 解决 bank conflict | 1475.2       | 97.536GB/s | 2.29   |
 
 可以看到，解决 bank conflict 之后，性能和带宽都有了很大的提升。
 
@@ -99,7 +99,7 @@ for(int s=blockDim.x/2; s>0; s >>= 1) {
 nvcc -o reduce_bank_conflict_free reduce_bank_conflict_free.cu
 ```
 
-## Reference 
+## Reference
 
 - [https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/](https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/)
 - [http://giantpandacv.com/project/OneFlow/%E3%80%90BBuf%E7%9A%84CUDA%E7%AC%94%E8%AE%B0%E3%80%91%E4%B8%89%EF%BC%8Creduce%E4%BC%98%E5%8C%96%E5%85%A5%E9%97%A8%E5%AD%A6%E4%B9%A0%E7%AC%94%E8%AE%B0/#2-bank-conflict](http://giantpandacv.com/project/OneFlow/%E3%80%90BBuf%E7%9A%84CUDA%E7%AC%94%E8%AE%B0%E3%80%91%E4%B8%89%EF%BC%8Creduce%E4%BC%98%E5%8C%96%E5%85%A5%E9%97%A8%E5%AD%A6%E4%B9%A0%E7%AC%94%E8%AE%B0/#2-bank-conflict)

docs/09_optimize_reduce/04_unroll/README.md

@@ -1,6 +1,6 @@
 # 展开 Warp
 
-现在我们已经使用了 3 种方法对 Reduce Kernel 进行优化 （交错寻址、去除 Bank Confilt、减少空闲线程）。
+现在我们已经使用了 3 种方法对 Reduce Kernel 进行优化 （交错寻址、去除 Bank Conflilt、减少空闲线程）。
 当下实现的 Kernel 距离理论带宽还有一定距离，我们可以继续优化。Reduce 并不是一个算术密集型的 Kernel。
 对于这种 Kernel，一个可能的瓶颈就是地址算术指令和循环的开销。
 
@@ -28,8 +28,7 @@ for (int s = blockDim.x / 2; s > 0; s >>= 1)
 }
 ```
 
-每一次循环都会进行一个 BLOCK 中线程的同步。但是实际上当 `s <= 32` 的时候我们只用到了一个 Warp 的线程。
-在最后一个 Warp 中，由于一个 Warp 中的线程都是在同一个 simd 单元上的这些线程本来就是同步的，所以这个同步就是没有必要的了。
+每一次循环都会进行一个 BLOCK 中线程的同步。但是实际上当 `s <= 32` 的时候，由于 `tid <= s` 所以我们只用到了一个 Warp 的线程。由于 cuda 是单指令多线程的设计，所以同一个 Warp 中的线程都是并行执行的。所以最后一个 Warp 在同一个 simd 单元上的这些线程本来就是同步的，所以这个 `__syncthreads()` 同步就是没有必要的了。
 
 
 ## 2. 优化方案
@@ -41,12 +40,12 @@ for (int s = blockDim.x / 2; s > 0; s >>= 1)
 ```cpp
 __device__ void warp_reduce(volatile int *sdata, int tid)
 {
-    sdada[tid] += sdata[tid + 32];
-    sdada[tid] += sdata[tid + 16];
-    sdada[tid] += sdata[tid + 8];
-    sdada[tid] += sdata[tid + 4];
-    sdada[tid] += sdata[tid + 2];
-    sdada[tid] += sdata[tid + 1];
+    sdata[tid] += sdata[tid + 32];
+    sdata[tid] += sdata[tid + 16];
+    sdata[tid] += sdata[tid + 8];
+    sdata[tid] += sdata[tid + 4];
+    sdata[tid] += sdata[tid + 2];
+    sdata[tid] += sdata[tid + 1];
 }
 ```
 
@@ -87,13 +86,13 @@ nvcc -o reduce_unroll_last_warp reduce_unroll_last_warp.cu
 对上面的 Kernel 进行性能分析结果如下：
 
 
-| 优化手段 | 运行时间(us) | 带宽 | 加速比 |
-| --- | --- | --- | --- |
-| Baseline | 3118.4 | 42.503GB/s | ~ |
-| 交错寻址 | 1904.4 | 73.522GB/s | 1.64 |
-| 解决 bank conflict | 1475.2 | 97.536GB/s | 2.29 |
-| 去除 idle 线程 | 758.38 | 189.78GB/s | 4.11 |
-| 展开最后一个 Warp | 484.01 | 287.25GB/s | 6.44 |
+| 优化手段           | 运行时间(us) | 带宽       | 加速比 |
+| ------------------ | ------------ | ---------- | ------ |
+| Baseline           | 3118.4       | 42.503GB/s | ~      |
+| 交错寻址           | 1904.4       | 73.522GB/s | 1.64   |
+| 解决 bank conflict | 1475.2       | 97.536GB/s | 2.29   |
+| 去除 idle 线程     | 758.38       | 189.78GB/s | 4.11   |
+| 展开最后一个 Warp  | 484.01       | 287.25GB/s | 6.44   |
 
 ### 2.2. 完全展开
 
@@ -173,14 +172,14 @@ nvcc -o reduce_unroll_all reduce_unroll_all.cu
 
 对上面的 Kernel 进行性能分析结果如下：
 
-| 优化手段 | 运行时间(us) | 带宽(GB/s) | 加速比 |
-| --- | --- | --- | --- |
-| Baseline | 3118.4 | 42.503 | ~ |
-| 交错寻址 | 1904.4 | 73.522 | 1.64 |
-| 解决 bank conflict | 1475.2 | 97.536 | 2.29 |
-| 去除 idle 线程 | 758.38 | 189.78 | 4.11 |
-| 展开最后一个 Warp | 484.01 | 287.25 | 6.44 |
-| 完全展开 | 477.23 |  291.77 | 6.53 |
+| 优化手段           | 运行时间(us) | 带宽(GB/s) | 加速比 |
+| ------------------ | ------------ | ---------- | ------ |
+| Baseline           | 3118.4       | 42.503     | ~      |
+| 交错寻址           | 1904.4       | 73.522     | 1.64   |
+| 解决 bank conflict | 1475.2       | 97.536     | 2.29   |
+| 去除 idle 线程     | 758.38       | 189.78     | 4.11   |
+| 展开最后一个 Warp  | 484.01       | 287.25     | 6.44   |
+| 完全展开           | 477.23       | 291.77     | 6.53   |
 
 
 ## 3. 总结