Skip to content

Commit 9037af5

Browse files
muyuuuumuyuuuu
authored
[Doc] Fix some typo of reduce optimize (#52)
1 parent c0c6b4b commit 9037af5

File tree

2 files changed

+41
-42
lines changed

2 files changed

+41
-42
lines changed

‎docs/09_optimize_reduce/02_bank_conflict/README.md

+18-18
Original file line numberDiff line numberDiff line change
@@ -4,27 +4,27 @@
44

55
## 1. Bank Conflict
66

7-
Bank conflict 是指当多个线程同时访问同一个 bank 的时候，会出现 bank conflict。bank 是共享内存的最小单元，每个 bank 可以同时为一个线程提供数据。如果多个线程同时访问同一个 bank，那么它们的访问就会被串行化，从而降低性能。为了避免或减少访存冲突，可以使用一些技巧，如改变数据布局、使用 padding、使用 shuffle 指令等。
7+
> 为了获得高的内存带宽，共享内存在物理上被分为 32 个同样宽度的、能被同时访问的内存 bank。共享内存中每连续的 128 字节的内容分摊到 32 个 bank 的同一层中。bank 是共享内存的最小单元。
88
9-
同一个 Block 的线程会共享一块共享内存，共享内存在逻辑上被分为 32 个 bank。当同一个 warp 中的线程访问同一个 bank 的时候，会出现 bank conflict。在最坏的情况下，即一个 warp 中的所有线程访问了相同 bank 的 32 个不同地址的话，那么这 32 个访问操作将会全部被序列化，大大降低了内存带宽。在一个 warp 内对同一个 bank 中的 n 个地址同时访问将导致 n 次内存事务，称为发生了 n 路 bank conflict。需要注意的是，不同 warp 中的线程之间不存在 bank conflict。
9+
同一个 Block 的线程会共享一块共享内存，Bank conflict 是指一个 warp 内的多个线程同时访问同一个 bank 的不同地址，那么它们的访问就会被串行化，从而降低性能。在最坏的情况下，即一个 warp 中的所有线程访问了相同 bank 的 32 个不同地址的话，那么这 32 个访问操作将会全部被序列化，大大降低了内存带宽。在一个 warp 内对同一个 bank 中的 n 个地址同时访问将导致 n 次内存事务，称为发生了 n 路 bank conflict。需要注意的是，不同 warp 中的线程之间不存在 bank conflict。为了避免或减少访存冲突，可以使用一些技巧，如改变数据布局、使用 padding、使用 shuffle 指令等。
1010

1111
接下来让我们来分析上一个 Kernel 中的 bank conflict。上一个 Kernel 中，我们使用了交叉寻址的方式，使得连续的线程访问的地址不再连续。这样做的目的是为了避免 warp divergence。但是这样做也会导致 bank conflict。
1212

1313
我们以 0 号 warp 为例。在第一次迭代中，0 号线程需要加载 shared memory 的 0 号和 1 号地址，然后写回 0 号地址。同时，0 号 warp 的 16 号线程需要加载 shared memory 的 32 和 33 号地址，并写回 32 号地址。因此，在一个 warp 内同时访问了一个 bank 的不同内存地址，导致发生了 2 路的 Bank Conflict，如下图所示：
1414

15-
![picture 1](images/ef322be7c3e5b6b9be69d2b90e88083f50569a58a97129f348e483b946ab4edf.png)
15+
![picture 1](images/ef322be7c3e5b6b9be69d2b90e88083f50569a58a97129f348e483b946ab4edf.png)
1616

17-
类似地，在第二次迭代过程中，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 的影响。
17+
类似地，在第二次迭代过程中，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 的影响。
1818

1919
我们可以使用 nvprof 来查看 bank conflict 的情况。
2020

2121
```bash
2222
nvprof --events shared_st_bank_conflict ./reduce_interleaved_addressing
2323
```
2424

25-
| Invocations | Event Name | Min | Max | Avg | Total |
26-
|-------------|--------------------------|----------|----------|----------|----------|
27-
| 1 | shared_st_bank_conflict | 4587520 | 4587520 | 4587520 | 4587520 |
25+
| Invocations | Event Name | Min | Max | Avg | Total |
26+
| ----------- | ----------------------- | ------- | ------- | ------- | ------- |
27+
| 1 | shared_st_bank_conflict | 4587520 | 4587520 | 4587520 | 4587520 |
2828

2929

3030
如果你的设备不支持 nvprof，你可以使用 nsight-compute 的命令行工具 ncu 来查看 bank conflict 的情况。
@@ -44,13 +44,13 @@ sudo ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum
4444

4545
每个 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，如下图所示：
4646

47-
![picture 4](images/e69b477993846936b270e82a37615c00424010cd8003f429354aa27325c96f57.png)
47+
![picture 4](images/e69b477993846936b270e82a37615c00424010cd8003f429354aa27325c96f57.png)
4848

4949
0 号 warp 的 0 号线程访问 0 号 bank 的 0 号地址和 128 号地址，1 号 warp 的 0 号线程访问 1 号 bank 的 32 号地址和 128 号地址，2 号 warp 的 0 号线程访问 2 号 bank 的 64 号地址和 160 号地址，以此类推。
5050

5151
整体过程如下图所示：
5252

53-
![picture 6](images/0f65c7d9e911014e31ddd84c583dea859ba24ebd48715c2680eb604e7ebb9a2b.png)
53+
![picture 6](images/0f65c7d9e911014e31ddd84c583dea859ba24ebd48715c2680eb604e7ebb9a2b.png)
5454

5555
> [!NOTE]
5656
> 图里面的 block size 是 16 而不是 256，这是为了方便说明。实际上，我们的 block size 是 256。
@@ -79,17 +79,17 @@ for(int s=blockDim.x/2; s>0; s >>= 1) {
7979

8080
修改后我们可以再次运行 nvprof 来查看 bank conflict 的情况，输出如下所示：
8181

82-
| Invocations | Event Name | Min | Max | Avg | Total |
83-
|-------------|--------------------------|----------|----------|----------|----------|
84-
| 1 | shared_st_bank_conflict | 0 | 0 | 0 | 0 |
82+
| Invocations | Event Name | Min | Max | Avg | Total |
83+
| ----------- | ----------------------- | --- | --- | --- | ----- |
84+
| 1 | shared_st_bank_conflict | 0 | 0 | 0 | 0 |
8585

8686
性能和带宽的测试情况如下：
8787

88-
| 优化手段 | 运行时间(us) | 带宽 | 加速比 |
89-
| --- | --- | --- | --- |
90-
| Baseline | 3118.4 | 42.503GB/s | ~ |
91-
| 交错寻址 | 1904.4 | 73.522GB/s | 1.64 |
92-
| 解决 bank conflict | 1475.2 | 97.536GB/s | 2.29 |
88+
| 优化手段 | 运行时间(us) | 带宽 | 加速比 |
89+
| ------------------ | ------------ | ---------- | ------ |
90+
| Baseline | 3118.4 | 42.503GB/s | ~ |
91+
| 交错寻址 | 1904.4 | 73.522GB/s | 1.64 |
92+
| 解决 bank conflict | 1475.2 | 97.536GB/s | 2.29 |
9393

9494
可以看到，解决 bank conflict 之后，性能和带宽都有了很大的提升。
9595

@@ -99,7 +99,7 @@ for(int s=blockDim.x/2; s>0; s >>= 1) {
9999
nvcc -o reduce_bank_conflict_free reduce_bank_conflict_free.cu
100100
```
101101

102-
## Reference
102+
## Reference
103103

104104
- [https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/](https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/)
105105
- [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

+23-24
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
# 展开 Warp
22

3-
现在我们已经使用了 3 种方法对 Reduce Kernel 进行优化 （交错寻址、去除 Bank Confilt、减少空闲线程）。
3+
现在我们已经使用了 3 种方法对 Reduce Kernel 进行优化 （交错寻址、去除 Bank Conflilt、减少空闲线程）。
44
当下实现的 Kernel 距离理论带宽还有一定距离，我们可以继续优化。Reduce 并不是一个算术密集型的 Kernel。
55
对于这种 Kernel，一个可能的瓶颈就是地址算术指令和循环的开销。
66

@@ -28,8 +28,7 @@ for (int s = blockDim.x / 2; s > 0; s >>= 1)
2828
}
2929
```
3030

31-
每一次循环都会进行一个 BLOCK 中线程的同步。但是实际上当 `s <= 32` 的时候我们只用到了一个 Warp 的线程。
32-
在最后一个 Warp 中，由于一个 Warp 中的线程都是在同一个 simd 单元上的这些线程本来就是同步的，所以这个同步就是没有必要的了。
31+
每一次循环都会进行一个 BLOCK 中线程的同步。但是实际上当 `s <= 32` 的时候，由于 `tid <= s` 所以我们只用到了一个 Warp 的线程。由于 cuda 是单指令多线程的设计，所以同一个 Warp 中的线程都是并行执行的。所以最后一个 Warp 在同一个 simd 单元上的这些线程本来就是同步的，所以这个 `__syncthreads()` 同步就是没有必要的了。
3332

3433

3534
## 2. 优化方案
@@ -41,12 +40,12 @@ for (int s = blockDim.x / 2; s > 0; s >>= 1)
4140
```cpp
4241
__device__ void warp_reduce(volatile int *sdata, int tid)
4342
{
44-
sdada[tid] += sdata[tid + 32];
45-
sdada[tid] += sdata[tid + 16];
46-
sdada[tid] += sdata[tid + 8];
47-
sdada[tid] += sdata[tid + 4];
48-
sdada[tid] += sdata[tid + 2];
49-
sdada[tid] += sdata[tid + 1];
43+
sdata[tid] += sdata[tid + 32];
44+
sdata[tid] += sdata[tid + 16];
45+
sdata[tid] += sdata[tid + 8];
46+
sdata[tid] += sdata[tid + 4];
47+
sdata[tid] += sdata[tid + 2];
48+
sdata[tid] += sdata[tid + 1];
5049
}
5150
```
5251
@@ -87,13 +86,13 @@ nvcc -o reduce_unroll_last_warp reduce_unroll_last_warp.cu
8786
对上面的 Kernel 进行性能分析结果如下：
8887

8988

90-
| 优化手段 | 运行时间(us) | 带宽 | 加速比 |
91-
| --- | --- | --- | --- |
92-
| Baseline | 3118.4 | 42.503GB/s | ~ |
93-
| 交错寻址 | 1904.4 | 73.522GB/s | 1.64 |
94-
| 解决 bank conflict | 1475.2 | 97.536GB/s | 2.29 |
95-
| 去除 idle 线程 | 758.38 | 189.78GB/s | 4.11 |
96-
| 展开最后一个 Warp | 484.01 | 287.25GB/s | 6.44 |
89+
| 优化手段 | 运行时间(us) | 带宽 | 加速比 |
90+
| ------------------ | ------------ | ---------- | ------ |
91+
| Baseline | 3118.4 | 42.503GB/s | ~ |
92+
| 交错寻址 | 1904.4 | 73.522GB/s | 1.64 |
93+
| 解决 bank conflict | 1475.2 | 97.536GB/s | 2.29 |
94+
| 去除 idle 线程 | 758.38 | 189.78GB/s | 4.11 |
95+
| 展开最后一个 Warp | 484.01 | 287.25GB/s | 6.44 |
9796

9897
### 2.2. 完全展开
9998

@@ -173,14 +172,14 @@ nvcc -o reduce_unroll_all reduce_unroll_all.cu
173172

174173
对上面的 Kernel 进行性能分析结果如下：
175174

176-
| 优化手段 | 运行时间(us) | 带宽(GB/s) | 加速比 |
177-
| --- | --- | --- | --- |
178-
| Baseline | 3118.4 | 42.503 | ~ |
179-
| 交错寻址 | 1904.4 | 73.522 | 1.64 |
180-
| 解决 bank conflict | 1475.2 | 97.536 | 2.29 |
181-
| 去除 idle 线程 | 758.38 | 189.78 | 4.11 |
182-
| 展开最后一个 Warp | 484.01 | 287.25 | 6.44 |
183-
| 完全展开 | 477.23 | 291.77 | 6.53 |
175+
| 优化手段 | 运行时间(us) | 带宽(GB/s) | 加速比 |
176+
| ------------------ | ------------ | ---------- | ------ |
177+
| Baseline | 3118.4 | 42.503 | ~ |
178+
| 交错寻址 | 1904.4 | 73.522 | 1.64 |
179+
| 解决 bank conflict | 1475.2 | 97.536 | 2.29 |
180+
| 去除 idle 线程 | 758.38 | 189.78 | 4.11 |
181+
| 展开最后一个 Warp | 484.01 | 287.25 | 6.44 |
182+
| 完全展开 | 477.23 | 291.77 | 6.53 |
184183

185184

186185
## 3. 总结

0 commit comments

Comments
 (0)