[架构设计] CUDA系列学习（五）GPU基础算法: Reduce, Scan, Histogram

喵~不知不觉到了CUDA系列学习第五讲，前几讲中我们主要介绍了基础GPU中的软硬件结构，内存管理，task类型等；这一讲中我们将介绍3个基础的GPU算法：reduce，scan，histogram，它们在并行算法中非常常用，我们在本文中分别就其功能用处，串行与并行实现进行阐述。 
———-

1. Task complexity

task complexity包括step complexity（可以并行成几个操作） & work complexity（总共有多少个工作要做）。 
e.g. 下面的tree-structure图中每个节点表示一个操作数，每条边表示一个操作，同层edge表示相同操作，问该图表示的task的step complexity & work complexity分别是多少。

Ans: 
step complexity: 3； 
work complexity: 6。 
下面会有更具体的例子。

---

2. Reduce

引入：我们考虑一个task：1+2+3+4+… 
1) 最简单的顺序执行顺序组织为((1+2)+3)+4… 
2) 由于operation之间没有依赖关系，我们可以用Reduce简化操作，它可以减少serial implementation的步数。 

2.1 what is reduce?

Reduce input:

1. set of elements
2. reduction operation 
   
   1. binary: 两个输入一个输出
   2. 操作满足结合律： (a@b)@c = a@(b@c), 其中@表示operator 
      e.g +, 按位与 都符合；a^b(expotentiation)和减法都不是

 

2.1.1 Serial implementation of Reduce:

reduce的每一步操作都依赖于其前一个操作的结果。比如对于前面那个例子，n个数相加，work complexity 和 step complexity都是O(n)（原因不言自明吧~）我们的目标就是并行化操作，降下来step complexity. e.g add serial reduce -> parallel reduce。 

2.1.2 Parallel implementation of Reduce:

也就是说，我们把step complexity降到了 log2n

举个栗子，如下图所示： 

那么如果对 210 个数做parallel reduce add，其step complexity就是10. 那么在这个parallel reduce的第一步，我们需要做512个加法，这对modern gpu不是啥大问题，但是如果我们要对 220 个数做加法呢？就需要考虑到gpu数量了，如果说gpu最多能并行做512个操作，我们就应将 220 个数分成1024*1024(共1024组)，每次做 210 个数的加法。这种考虑task规模和gpu数量关系的做法有个理论叫Brent’s Theory. 下面我们具体来看：

也就是进行两步操作，第一步分成1024个block，每个block做加法；第二步将这1024个结果再用1个1024个thread的block进行求和。kernel code：

<code class="hljs objectivec has-numbering" style="display: block; padding: 0px; color: inherit; box-sizing: border-box; font-family: 'Source Code Pro', monospace;font-size:undefined; white-space: pre; border-top-left-radius: 0px; border-top-right-radius: 0px; border-bottom-right-radius: 0px; border-bottom-left-radius: 0px; word-wrap: normal; background: transparent;">__global__ <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">void</span> parallel_reduce_kernel(<span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">float</span> *d_out, <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">float</span>* d_in){
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">int</span> myID = threadIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span> + blockIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span> * blockDim<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span>;
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">int</span> tid = threadIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span>;

    <span class="hljs-comment" style="color: rgb(136, 0, 0); box-sizing: border-box;">//divide threads into two parts according to threadID, and add the right part to the left one, lead to reducing half elements, called an iteration; iterate until left only one element</span>
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">for</span>(<span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">unsigned</span> <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">int</span> s = blockDim<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span> / <span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">2</span> ; s><span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">0</span>; s>>=<span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">1</span>){
        <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">if</span>(tid<s){
            d_in[myID] += d_in[myID + s];
        }
        __syncthreads(); <span class="hljs-comment" style="color: rgb(136, 0, 0); box-sizing: border-box;">//ensure all adds at one iteration are done</span>
    }
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">if</span> (tid == <span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">0</span>){
        d_out[blockIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span>] = d_in[myId];
    }
}</code><ul class="pre-numbering" style="box-sizing: border-box; position: absolute; width: 50px; top: 0px; left: 0px; margin: 0px; padding: 6px 0px 40px; border-right-width: 1px; border-right-style: solid; border-right-color: rgb(221, 221, 221); list-style: none; text-align: right; background-color: rgb(238, 238, 238);"><li style="box-sizing: border-box; padding: 0px 5px;">1</li><li style="box-sizing: border-box; padding: 0px 5px;">2</li><li style="box-sizing: border-box; padding: 0px 5px;">3</li><li style="box-sizing: border-box; padding: 0px 5px;">4</li><li style="box-sizing: border-box; padding: 0px 5px;">5</li><li style="box-sizing: border-box; padding: 0px 5px;">6</li><li style="box-sizing: border-box; padding: 0px 5px;">7</li><li style="box-sizing: border-box; padding: 0px 5px;">8</li><li style="box-sizing: border-box; padding: 0px 5px;">9</li><li style="box-sizing: border-box; padding: 0px 5px;">10</li><li style="box-sizing: border-box; padding: 0px 5px;">11</li><li style="box-sizing: border-box; padding: 0px 5px;">12</li><li style="box-sizing: border-box; padding: 0px 5px;">13</li><li style="box-sizing: border-box; padding: 0px 5px;">14</li><li style="box-sizing: border-box; padding: 0px 5px;">15</li></ul>

Quiz: 看一下上面的code可以从哪里进行优化？

Ans：我们在上一讲中提到了global，shared & local memory的速度，那么这里对于global memory的操作可以更改为shared memory，从而进行提速：

<code class="hljs objectivec has-numbering" style="display: block; padding: 0px; color: inherit; box-sizing: border-box; font-family: 'Source Code Pro', monospace;font-size:undefined; white-space: pre; border-top-left-radius: 0px; border-top-right-radius: 0px; border-bottom-right-radius: 0px; border-bottom-left-radius: 0px; word-wrap: normal; background: transparent;">__global__ <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">void</span> parallel_shared_reduce_kernel(<span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">float</span> *d_out, <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">float</span>* d_in){
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">int</span> myID = threadIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span> + blockIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span> * blockDim<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span>;
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">int</span> tid = threadIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span>;
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">extern</span> __shared__ <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">float</span> sdata[];
    sdata[tid] = d_in[myID];
    __syncthreads();

    <span class="hljs-comment" style="color: rgb(136, 0, 0); box-sizing: border-box;">//divide threads into two parts according to threadID, and add the right part to the left one, lead to reducing half elements, called an iteration; iterate until left only one element</span>
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">for</span>(<span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">unsigned</span> <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">int</span> s = blockDim<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span> / <span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">2</span> ; s><span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">0</span>; s>>=<span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">1</span>){
        <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">if</span>(tid<s){
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads(); <span class="hljs-comment" style="color: rgb(136, 0, 0); box-sizing: border-box;">//ensure all adds at one iteration are done</span>
    }
    <span class="hljs-keyword" style="color: rgb(0, 0, 136); box-sizing: border-box;">if</span> (tid == <span class="hljs-number" style="color: rgb(0, 102, 102); box-sizing: border-box;">0</span>){
        d_out[blockIdx<span class="hljs-variable" style="color: rgb(102, 0, 102); box-sizing: border-box;">.x</span>] = sdata[myId];
    }
}</code><ul class="pre-numbering" style="box-sizing: border-box; position: absolute; width: 50px; top: 0px; left: 0px; margin: 0px; padding: 6px 0px 40px; border-right-width: 1px; border-right-style: solid; border-right-color: rgb(221, 221, 221); list-style: none; text-align: right; background-color: rgb(238, 238, 238);"><li style="box-sizing: border-box; padding: 0px 5px;">1</li><li style="box-sizing: border-box; padding: 0px 5px;">2</li><li style="box-sizing: border-box; padding: 0px 5px;">3</li><li style="box-sizing: border-box; padding: 0px 5px;">4</li><li style="box-sizing: border-box; padding: 0px 5px;">5</li><li style="box-sizing: border-box; padding: 0px 5px;">6</li><li style="box-sizing: border-box; padding: 0px 5px;">7</li><li style="box-sizing: border-box; padding: 0px 5px;">8</li><li style="box-sizing: border-box; padding: 0px 5px;">9</li><li style="box-sizing: border-box; padding: 0px 5px;">10</li><li style="box-sizing: border-box; padding: 0px 5px;">11</li><li style="box-sizing: border-box; padding: 0px 5px;">12</li><li style="box-sizing: border-box; padding: 0px 5px;">13</li><li style="box-sizing: border-box; padding: 0px 5px;">14</li><li style="box-sizing: border-box; padding: 0px 5px;">15</li><li style="box-sizing: border-box; padding: 0px 5px;">16</li><li style="box-sizing: border-box; padding: 0px 5px;">17</li><li style="box-sizing: border-box; padding: 0px 5px;">18</li><li></li></ul>