日本搞逼视频_黄色一级片免费在线观看_色99久久_性明星video另类hd_欧美77_综合在线视频

國內最全IT社區平臺 聯系我們 | 收藏本站
阿里云優惠2
您當前位置:首頁 > php框架 > 框架設計 > CUDA系列學習(五)GPU基礎算法: Reduce, Scan, Histogram

CUDA系列學習(五)GPU基礎算法: Reduce, Scan, Histogram

來源:程序員人生   發布時間:2015-07-03 08:53:05 閱讀次數:6864次
喵~不知不覺到了CUDA系列學習第5講,前幾講中我們主要介紹了基礎GPU中的軟硬件結構,內存管理,task類型等;這1講中我們將介紹3個基礎的GPU算法:reduce,scan,histogram,它們在并行算法中非常經常使用,我們在本文中分別就其功能用途,串行與并行實現進行論述。

1. Task complexity

task complexity包括step complexity(可以并行成幾個操作) & work complexity(總共有多少個工作要做)。
e.g. 下面的tree-structure圖中每一個節點表示1個操作數,每條邊表示1個操作,同層edge表示相同操作,問該圖表示的task的step complexity & work complexity分別是多少。

tree operation

Ans:
step complexity: 3;
work complexity: 6。
下面會有更具體的例子。




## 2. Reduce 引入:我們斟酌1個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: 兩個輸入1個輸出 2. 操作滿足結合律: (a@b)@c = a@(b@c), 其中@表示operator e.g +, 按位與 都符合;a^b(expotentiation)和減法都不是 ![2. add_tree.png](http://img.blog.csdn.net/20150213145544805)


### 2.1.1 Serial implementation of Reduce: reduce的每步操作都依賴于其前1個操作的結果。比如對前面那個例子,n個數相加,work complexity 和 step complexity都是O(n)(緣由不言自明吧~)我們的目標就是并行化操作,降下來step complexity. e.g add serial reduce -> parallel reduce。

### 2.1.2 Parallel implementation of Reduce: ![3. parallel_add.png](http://img.blog.csdn.net/20150213145641025) 也就是說,我們把step complexity降到了log2n 那末如果對210個數做parallel reduce add,其step complexity就是10. 那末在這個parallel reduce的第1步,我們需要做512個加法,這對modern gpu不是啥大問題,但是如果我們要對220個數做加法呢?就需要斟酌到gpu數量了,如果說gpu最多能并行做512個操作,我們就應將220個數分成1024*1024(共1024組),每次做210個數的加法。這類斟酌task范圍和gpu數量關系的做法有個理論叫Brent’s Theory. 下面我們具體來看: ![4. brent’s theory.png](http://img.blog.csdn.net/20150213145804704) 也就是進行兩步操作,第1步分成1024個block,每一個block做加法;第2步將這1024個結果再用1個1024個thread的block進行求和。kernel code:
__global__ void parallel_reduce_kernel(float *d_out, float* d_in){ int myID = threadIdx.x + blockIdx.x * blockDim.x; int tid = threadIdx.x; //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 for(unsigned int s = blockDim.x / 2 ; s>0; s>>1){ if(tid<s){ d_in[myID] += d_in[myID + s]; } __syncthreads(); //ensure all adds at one iteration are done } if (tid == 0){ d_out[blockIdx.x] = d_in[myId]; } }


Quiz: 看1下上面的code可以從哪里進行優化?
Ans:我們在上1講中提到了global,shared & local memory的速度,那末這里對global memory的操作可以更改成shared memory,從而進行提速:
__global__ void parallel_shared_reduce_kernel(float *d_out, float* d_in){ int myID = threadIdx.x + blockIdx.x * blockDim.x; int tid = threadIdx.x; extern __shared__ float sdata[]; sdata[tid] = d_in[myID]; __syncthreads(); //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 for(unsigned int s = blockDim.x / 2 ; s>0; s>>1){ if(tid<s){ sdata[myID] += sdata[myID + s]; } __syncthreads(); //ensure all adds at one iteration are done } if (tid == 0){ d_out[blockIdx.x] = sdata[myId]; } }

優化的代碼中還有1點要注意,就是聲明的時候記得我們第3講中說過的kernel通用表示情勢:
kernel<<<grid of blocks, block of threads, shmem>>>
最后1項要在call kernel的時候聲明好,即:
parallel_reduce_kernel<<<blocks, threads, threads*sizeof(float)>>>(data_out, data_in);


好,那末問題來了,對這兩個版本(parallel_reduce_kernel 和 parallel_shared_reduce_kernel), parallel_reduce_kernel比parallel_shared_reduce_kernel多用了幾倍的global memory帶寬? Ans: 分別斟酌兩個版本的讀寫操作:
parallel_reduce_kernel
Times Read Ops Write Ops
1 1024 512
2 512 256
3 256 128
n 1 1
parallel_shared_reduce_kernel
Times Read Ops Write Ops
1 1024 1

所以,parallel_reduce_kernel所需的帶寬是parallel_shared_reduce_kernel的3倍。


3. Scan

3.1 what is scan?

  • Example:

    • input: 1,2,3,4
    • operation: Add
    • ouput: 1,3,6,10(out[i]=sum(in[0:i]))
  • 目的:解決難以并行的問題

拍拍腦袋想一想上面這個問題O(n)的1個解法是out[i] = out[i⑴] + in[i].下面我們來引入scan。

Inputs to scan:

  1. input array
  2. 操作:binary & 滿足結合律(和reduce1樣)
  3. identity element [I op a = a], 其中I 是identity element
    quiz: what is the identity for 加法,乘法,邏輯與,邏輯或?
    Ans:
op Identity
加法 0
乘法 1
邏輯或|| False
邏輯與&& True



3.2 what scan does?

I/O content
input [a0 a1 a2 an]
output [I a0 a0?a1 a0?a1??an]

其中?是scan operator,I 是?的identity element




3.2.1 Serial implementation of Scan

很簡單:

int acc = identity; for(i=0;i<elements.length();i++){ acc = acc op elements[i]; out[i] = acc; }

work complexity: O(n)
step complexity: O(n)

那末,對scan問題,我們怎樣對其進行并行化呢?



3.2.1 Parallel implementation of Scan

斟酌scan的并行化,可以并行計算n個output,每一個output元素i相當于a0?a1??ai,是1個reduce operation。

Q: 那末問題的work complexity和step complexity分別變成多少了呢?
Ans:

  • step complexity:
    取決于n個reduction中耗時最長的,即O(log2n)
  • work complexity:
    對每一個output元素進行計算,總計算量為0+1+2+…+(n⑴),所以復雜度為O(n2).

可見,step complexity降下來了,惋惜work complexity上去了,那末怎樣解決呢?這里有兩種Scan算法:

more step efficiency more work efficiency
hillis + steele (1986)
blelloch (1990)




  1. Hillis + Steele

    對Scan加法問題,hillis+steele算法的解決方案以下:

hillis + steele

即streaming’s
step 0: out[i] = in[i] + in[i⑴];
step 1: out[i] = in[i] + in[i⑵];
step 2: out[i] = in[i] + in[i⑷];
如果元素不存在(向下越界)就記為0;可見step 2的output就是scan 加法的結果(想一想為何,我們1會再分析)。

那末問題來了。。。
Q: hillis + steele算法的work complexity 和 step complexity分別為多少?

Hillis + steele Algorithm complexity
log(n) O(n) O(n) O(nlogn) O(n^2)
work complexity
step complexity

解釋:

為了無妨礙大家思路,我在表格中將答案設為了白色,選中表格可見答案。

  1. step complexity:
    由于第i個step的結果為上1步輸出作為in, out[idx] = in[idx] + in[idx - 2^i], 所以step complexity = O(log(n))
  2. work complexity:
    workload = (n?1)+(n?2)+(n?4)+... ,共有log(n)項元素相加,所以可以近似看作1個矩陣,對應上圖,長log(n), 寬n,所以復雜度為 nlog( 生活不易,碼農辛苦
    如果您覺得本網站對您的學習有所幫助,可以手機掃描二維碼進行捐贈
    程序員人生
------分隔線----------------------------
分享到:
------分隔線----------------------------
關閉
程序員人生
主站蜘蛛池模板: 亚洲国产精品久久久久久 | 免费网站黄 | 国产精品区一区二区三 | 国产伊人精品 | 91网站免费看 | 精品国产欧美一区二区三区成人 | 亚州国产| 三级av毛片| 国产精品极品 | 午夜毛片 | 日韩欧美在线免费观看 | 亚洲 欧美 视频 | 综合久久一区二区 | 看黄在线观看 | 91国自产精品中文字幕亚洲 | 99在线观看 | 日韩城人免费 | 欧美 日韩 国产 在线 | 中文字幕在线观看一区二区三区 | 激情在线视频 | av在线一区二区 | 日韩一区二区在线视频 | 玖玖色资源 | 玖玖精品视频 | 欧美a性 | 色av影视| 美女黄18| 国产精品美女久久久久 | 国产精品国产a级 | a级毛片毛片免费很很综合 91久久 | 亚洲综合久久久 | 久久精品123 | 亚洲综合无码一区二区 | av一区二区三区在线播放 | 亚洲三级免费 | 国产免费小视频 | 国产黄色大片免费 | 国产免费久久 | av在线网站观看 | av片在线看免费高清网站 | 亚洲毛茸茸少妇高潮呻吟 |