File tree Expand file tree Collapse file tree 1 file changed +7
-7
lines changed Expand file tree Collapse file tree 1 file changed +7
-7
lines changed Original file line number Diff line number Diff line change 6
6
7
7
8
8
9
- ## 2. 使用 Gpu 实现 Reduce
9
+ ## 2. 使用 GPU 实现 Reduce
10
10
11
- ### 2.1 Cpu 实现
11
+ ### 2.1 CPU 实现
12
12
13
13
14
14
``` cpp
@@ -22,11 +22,11 @@ int reduce(int *arr, int len) {
22
22
}
23
23
```
24
24
25
- 可以看到,cpu 实现非常简单,一层循环就可以解决问题。但是,这个实现并不是纯函数式的,因为它有副作用,即修改了 `sum` 的值。我们可以用 Gpu 来实现一个纯函数式的 `Reduce`。
25
+ 可以看到,cpu 实现非常简单,一层循环就可以解决问题。但是,这个实现并不是纯函数式的,因为它有副作用,即修改了 `sum` 的值。我们可以用 GPU 来实现一个纯函数式的 `Reduce`。
26
26
27
- ### 2.2 Gpu 层次结构
27
+ ### 2.2 GPU 层次结构
28
28
29
- 首先我们先来回归一下 Gpu 的层次结构。也就是代码中的 `block` 和 `grid`。`block` 是一个线程块,`grid` 是一个线程网格。`block` 中的线程可以通过 `threadIdx.x` 来获取自己的线程 id,`grid` 中的线程可以通过 `blockIdx.x` 来获取自己的线程 id。`block` 中的线程可以通过 `blockDim.x` 来获取 `block` 的大小,`grid` 中的线程可以通过 `gridDim.x` 来获取 `grid` 的大小。
29
+ 首先我们先来回归一下 GPU 的层次结构。也就是代码中的 `block` 和 `grid`。`block` 是一个线程块,`grid` 是一个线程网格。`block` 中的线程可以通过 `threadIdx.x` 来获取自己的线程 id,`grid` 中的线程可以通过 `blockIdx.x` 来获取自己的线程 id。`block` 中的线程可以通过 `blockDim.x` 来获取 `block` 的大小,`grid` 中的线程可以通过 `gridDim.x` 来获取 `grid` 的大小。
30
30
31
31
对于 `Reduce` 来说,我们可以按照下面这个图设计:
32
32
@@ -40,7 +40,7 @@ GPU 的计算过程如下图所示:
40
40
41
41
以上图为例,我们来看一下 `Reduce` 的计算过程。首先,我们把数组分成了 `3` 个 `block`,每个 `block` 中有 `8` 个线程。在第一轮计算中,奇数线程会把自己的值累加到偶数线程中。在第二轮计算中,`block` 中的第 `0` 个线程会把 `4` 号线程的值累加到自己的值中。 每个 `block` 的值都计算完之后还需要对 `block` 的值进行累加,下面我们来看一下代码要如何实现。
42
42
43
- ### 2.3 Gpu 实现
43
+ ### 2.3 GPU 实现
44
44
45
45
我们首先看 `Kernel` 的实现,代码如下:
46
46
@@ -160,7 +160,7 @@ nvcc -o reduce_naive reduce_naive.cu
160
160
161
161
## 3. 总结
162
162
163
- 本文主要讨论了如何用 Gpu 来实现 ` Reduce ` 。我们首先回顾了 Gpu 的层次结构,然后我们用 Gpu 来实现了一个纯函数式的 ` Reduce ` 。最后,我们对比了 cpu 和 gpu 的实现,发现 gpu 的实现更加简洁,而且可以充分利用 Gpu 的并行计算能力。下一篇文章我们将讨论如何优化 ` Reduce ` 的实现。
163
+ 本文主要讨论了如何用 GPU 来实现 ` Reduce ` 。我们首先回顾了 GPU 的层次结构,然后我们用 GPU 来实现了一个纯函数式的 ` Reduce ` 。最后,我们对比了 cpu 和 GPU 的实现,发现 GPU 的实现更加简洁,而且可以充分利用 GPU 的并行计算能力。下一篇文章我们将讨论如何优化 ` Reduce ` 的实现。
164
164
165
165
## Reference
166
166
You can’t perform that action at this time.
0 commit comments