從0開始學習《GPU高效能運算之CUDA》——3

6 規約思想和同步概念

擴大點說，並行計算是有一種基本思想的，這個演算法能解決很多很常規的問題，而且很實用，比如說累加和累積等——規約思想。對於基礎的、重要的，我想有必要系統的學習。

我覺得有必要重新複製下之前寫的這篇介紹：

http://www.cnblogs.com/viviman/archive/2012/11/21/2780286.html

並行程式的開發有其不同於單核程式的特殊性，演算法是重中之重。根據不同業務設計出不同的並行演算法，直接影響到程式的效率。因此，如何設計並行程式的演算法，似乎成為並編程的最大痛點。觀其演算法，包括cuda sdk的例子和網上的牛人，給出的一些例子，以矩陣和向量處理為主，深入點的包括fft和julia等數學公式，再進階一點的算是圖形處理方面的例子。學習這些演算法的思想，免不了有自己的一點點總結。之前學習過omp編程，結合現在的cuda，我覺得要理解並行編程，首先理解劃分和規約這兩個概念。也許你的演算法學的更加紮實。劃分是《演算法》裡面的一個重要思想，將一個大的問題或任務，分解成小問題小任務，各個擊破，最後歸併結果；規約是《cuda**》書上介紹的一個入門的重要思想，規約演算法(reduction)用來求連加、連乘、最值等，應用廣泛。每次迴圈參加運算的線程減少一半。不管演算法的思想如何花樣，萬變不離其中的一點--將一個大的任務分解成小的任務集合，分解原則是粒度合適盡量小、資料相關性盡量小。如此而已。因為，我們用GPU是為了加速，要加速必須提高執行任務的並行度！明白這個道理，那麼我們將絞盡腦汁地去想方設法分析自己手上的任務，分解、分解、分解！這裡拿規約來說事情，因為，規約這個東西，似乎可以拿來單做9*9乘法表來熟悉，熟悉了基礎的口訣，那麼99*99的難題也會迎刃而解。ex：向量加法，要實現N=64*256長度的向量的累加和。假設a+b計算一次耗時t。

cpu計算：顯然單核的話需要64*256*t。我們容忍不了。

gpu計算：最初的設想，我們如果有個gpu能同時跑N/2個線程，我們這N/2個線程同時跑，那麼不是只需要t時間就能將N個數相加編程N/2個數相加了嗎？對的。這一輪我們用了t時間；接著的想法，我們不斷的遞迴這個過程，能發現嗎？第二輪，我們用N/2/2個線程同時跑，剩下N/2/2個數相加，這一輪我們同樣用了t時間；一直這樣想下去吧，最後一輪，我們用了1個線程跑，剩下1個數啦，這就是我們的結果！每一輪時間都為t，那麼理想情況，我們用了多少輪這樣的計算呢？計算次數=log(N)=6*8=48，對的，只用了48輪，也就是說，我們花了48*t的時間！

規約就是這樣，很簡單，很好用，我們且不管程式後期的最佳化，單從這個演算法分析上來說，從時間複雜度N降到了logN，這在常規演算法上，要提高成這樣的效率，是不得了的，這是指數層級的效率提高！所以，你會發現，GPU有CPU無法取代的得天獨厚的優勢——處理單元真心多啊！

規約求和的核函數代碼如下：__global__ void RowSum(float* A, float* B){ int bid = blockIdx.x; int tid = threadIdx.x;

__shared__ s_data[128]; //read data to shared memory s_data[tid] = A[bid*128 + tid]; __synctheads(); //sync

for(int i=64; i>0; i/=2){ if(tid<i) s_data[tid] = s_data[tid] + s_data[tid+i] ; __synctheads(); } if(tid==0) B[bid] = s_data[0];}

 

這個例子還讓我學到另一個東西——同步！我先不說同步是什麼，你聽我說個故事：我們調遣了10個小組從南京去日本打仗，我們的約定是，10個組可以自己行動，所有組在第三天在上海機場會合，然後一起去日本。這件事情肯定是需要處理的，不能第1組到了上海就先去日本了，這些先到的組，唯一可以做的事情是——等待！這個先來後到的事情，需要統一管理的時候，必須同步一下，在上海這個地方，大家統一下步調，快的組等等慢的組，然後一起幹接下去的旅程。

是不是很好理解，這就是同步在生活中的例子，應該這樣說，電腦的所有機制和演算法很多都是源於生活！結合起來，理解起來會簡單一點。

在CUDA中，我們的同步機制用處大嗎？又是如何用的呢？我告訴你，一個正常規模的工程中，一般來說資料都會有先來後到的關係，這一個計算結果可能是提供給另一個線程用的，這種依賴關係存在，會造成同步的應用。

__synctheads()這句話的作用是，這個block中的所有線程都執行到此的時候，都聽下來，等所有都執行到這個地方的時候，再往下執行。

7 撬開編程的鎖

鎖是資料相關性裡面肯定要用到的東西，很好，生活中也一樣，沒鎖，家裡不安全；GPU中沒鎖，資料會被“盜”。

對於存在競爭的資料，CUDA提供了原子操作函數——ATOM操作。

先亮出使用的例子：

__global__ void kernelfun()

{

__shared__ int i=0;

atomicAdd(&i, 1);

}

如果沒有加互斥機制，則同一個half warp內的線程將對i的操作混淆林亂。

用原子操作函數，可以很簡單的編寫自己的鎖，SKD中有給出的鎖結構體如下：

#ifndef __LOCK_H__

#define __LOCK_H__

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include "atomic_functions.h"

struct Lock {

    int *mutex;

    Lock( void ) {

        HANDLE_ERROR( cudaMalloc( (void**)&mutex, sizeof(int) ) );

        HANDLE_ERROR( cudaMemset( mutex, 0, sizeof(int) ) );

    }

    ~Lock( void ) {

        cudaFree( mutex );

    }

    __device__ void lock( void ) {

        while( atomicCAS( mutex, 0, 1 ) != 0 );

    }

    __device__ void unlock( void ) {

        atomicExch( mutex, 0 );

    }

};

#endif

 

8 CUDA軟體體繫結構

 

9 利用好現有的資源

如果連開方運算都需要自己去編寫程式實現，那麼我相信程式員這個職業將會縮水，沒有人願意去幹這種活。我想，程式員需要學會“偷懶”，現有的資源必須學會高效率的使用。當c++出現了STL庫，c++程式員的開發效率可以說倍增，而且程式穩定性更高。

CUDA有提供給我們什麼了嗎？給了，其實給了很多。

先介紹幾個庫：CUFFT、CUBLAS、CUDPP。

這裡我先不詳細學習這些庫裡到底有哪些函數，但是，大方向是需要瞭解的，不然找都不知道去哪兒找。CUFFT是傅裡葉變換的庫，CUBLAS提供了基本的矩陣和向量運算，CUDPP提供了常用的並行排序、搜尋等。

CUDA4.0以上，提供了一個類似STL的模板庫，初步窺探，只是一個類似vector的模板類型。有map嗎？map其實是一個散列表，可以用hashtable去實現這項機制。

SDK裡面有很多例子，包括一些通用的基本操作，比如InitCUDA等，都可以固化成函數組件，供新程式的調用。

具體的一些可以固化的東西，我將在以後的學習中歸納總結，豐富自己的CUDA庫！