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庫！