首先要搞清楚的就是執行緒網格(grid),執行緒塊(block)和執行緒(thread)之間的關係.
在前面的文章裡面就已經看到了核函式kernel<<>>但是並不知道這個核函式啟動的背後的一些原理,接下來就結合這幅圖來說一說基礎的東西.
圖中有乙個執行緒網格,網格中有2*3=6個執行緒塊,而每乙個執行緒塊裡面又有3*4=12個執行緒.這時候注意他們的索引.很容易算出這裡一共有6*12=72個執行緒.這裡說這個沒有什麼其他的意思,僅僅是想讓你形象的體會一下….
cuda內部定義了一些量來儲存乙個模型的各種索引:
griddim.x:執行緒網路x維度上線程塊的數量
griddim.y-執行緒網路y維度上線程塊的數量
blockdim.x-乙個執行緒塊x維度上的執行緒數量
blockdim.y-乙個執行緒塊y維度上的執行緒數量
blockidx.x-執行緒網路x維度上的執行緒塊索引
blockidx.y-執行緒網路y維度上的執行緒塊索引
threadidx.x-執行緒塊x維度上的執行緒索引
threadidx.y-執行緒塊y維度上的執行緒索引
griddim.x=3
griddim.y=2
blockdim.x=4
blockdim.y=3
有了上面的一些基本的概念(不需要完全掌握),接下來就正式開始講執行緒模型.
說到這裡,我們回到核函式kernel<<<.....>>>,我們可以把核函式寫為更加通用的方式,kernel<<>>:
dg：定義整個grid的維度，型別dim3，但是實際上目前顯示卡支援兩個維度，所以，dim3<1>>>第z維度預設只能為1，而從最開始的圖中可以看出grid的維度主要是決定block的數量的.而且,有最大數量限制.後面回來具體分析數量上面的限制.通俗一點,這個引數就是放的執行緒塊的形狀
db：定義了每個block的維度，block的維度是決定thread數量和分布的.型別dim3，比如512*512*64，這個可以定義3維尺寸，但是，這個地方是有講究了，三個維度的積是有上限的，後面會細講.通俗一點,這個引數就是放的執行緒的形狀
ns：是可選引數，設定最多能動態分配的共享記憶體大小，比如16k，單不需要是，這個值可以省略或寫0。
s：也是可選引數，表示流號，預設為0。流這個概念以後講。
我們還是以上面那幅圖為例子,
首先定義乙個blocks的dim3:dim3 blocks(2,3).發現什麼規律沒有?沒有發現的話,繼續,定義乙個threads的dim3:dim3 threads(3,4).定義好這兩個之後,直接丟到kernel中啟動kernel<<>>.然後裝置就會幫你啟動像圖裡面那樣型別的乙個執行緒模型.
說到這裡,你肯定對於基本的執行緒模型和核函式的意義有了一定的了解.也許了解不能夠算是很深入也沒有關係.這裡本來就是乙個很繞的地方.
然後接下來就是最後乙個,是想告訴你還有乙個東西叫做各種索引和他們的應用.但是要是這裡直接列出來的話,基本上是不知道意義在**,所以我準備在後面根據具體的例子來描述各種索引的使用.然後最後再總結.
首先要宣告,之前我們講了很多很」理論」的東西,所以後面的內容就是通過非常經典的例子來講怎麼在之戰中使用那些思想.
背景是是這樣的.假設有兩個長度相同的陣列a和b,把兩個陣列的元素分別相加.放到陣列res中去.
假如你不多多執行緒,就用基本的c/c++,你的第一想法是不是用for迴圈?哈哈哈,但是for迴圈你想都想得到一次只能夠操作乙個元素.所以,向量求和的問題是乙個非常簡單非常經典的並行問題.從這個問題說起,能夠引出很多很多基本的概念.所以理解這個經典的框架是很有意義的.
在這裡,就簡單的把兩個陣列設為
a= b=
我們應該得到的答案應該是
res=
**:
#include
using
namespace
std;
const
int max=10;
void add(int *a,int *b,int *res)
}int main()
解釋:
**很簡單,就是設了三個陣列,其中乙個是用來接收答案的,還有兩個是待相加的陣列.然後就分別各自位置上面的相加.add()函式寫成那個樣子也是模擬多cpu的情況,為了對於接下來的gpu的並行有乙個更好的了解….
這裡和之前一樣，先直接給出**，然後在分析體會其中的概念。
**:
#include
#include
using
namespace
std;
const
int max = 10;
//核函式建立
__global__ void add(int *a, int *b, int *res)
int main()
分析：
其實這裡關於前面開闢記憶體的一些前面已經講到過了,等一下再分析.最重要的也是變化最大的首先講.那就是核函式的意義和呼叫核函式時候的樣式.
首先看核函式的定義方式.
//核函式建立
__global__ void add(int
*a, int
*b, int
*res)
__global__什麼的這些就不講了,這些都是表明這是在裝置上面執行的**.函式的意義也是很明顯,就是傳遞3個位址進去(位址肯定是位於gpu上面),這些位址是我們之前在gpu上面開闢的記憶體的位址.而且有兩塊位址我們已經賦值了.這些都很簡單.
最不同的就是res[blockidx.x] = a[blockidx.x] + b[blockidx.x];這句話.初學者會覺得很疑惑.但是你能夠看到blockidx.x這個東東是我們前面提到過的.表示乙個執行緒塊的索引.目前為止,你知道這些就夠了.
接下來結合呼叫的語句:add << > > (dev_a, dev_b, dev_res);這裡其實之前也講到過了第乙個引數是啟動的執行緒塊的數量,第二個引數是每個執行緒塊中線程的數量.也就是說,這裡我們一次性」建立」了10(max)個執行緒塊.我們這時候可以認為,我們建立了10*1個核函式的副本.不負責任的說,你可以認為有10個一模一樣的函式在同時跑,之前乙個函式要執行10次的操作,現在10個函式一下就跑完了.
在第一部分我們就知道blockidx.x是起到了區分不同的執行緒塊的作用,因為我們這裡只用到了一維,所以索引只需要blockidx.x就行了.是不是突然覺得很簡單?而且對於每乙個執行緒塊(副本),執行緒塊的blockidx.x是系統弄好了的.你不用瞎操心.知道每個都有乙個索引就行.比如」建立」4個執行緒塊.那麼索引就是從0到3..建立了2個執行緒塊,索引就是從0到2.這個太簡單了,看第一部分複習一下.
最終,把不同的執行緒塊的索引作為乙個陣列的索引.自然就使得所有的10個加法一次性完成了.這就是他的本質思想了.
再回過頭去看一下,是不是有體會?
那既然可以建立10個包含乙個執行緒的執行緒塊,那能不能建立乙個執行緒塊,其中包含10個執行緒?答案是可以的!
直接給**:
#include
#include
using
namespace
std;
const
int max = 10;
//核函式建立
__global__ void add(int *a, int *b, int *res)
int main()
res[threadidx.x] = a[threadidx.x] + b[threadidx.x];把執行緒塊的索引改為了執行緒的索引.思想其實是一樣的.
另外乙個就是呼叫核函式的時候add << <1, max >> > (dev_a, dev_b, dev_res);建立乙個執行緒塊,其中有10個執行緒.得到的結果也是一樣的.就不是解釋更多了,自己去體會.
講到這裡,這個例子就講完了.你可以當做是在一維的情況下面對於索引方式的乙個實戰.可以啟發以後碰到一維的並行怎麼去設計.
未完待續…..
Windwos核心程式設計 （三）執行緒
執行緒有兩個組成部分 1 執行緒的核心物件，作業系統用它來管理執行緒 2 執行緒棧，用於維護執行緒執行時所需的所有函式引數和區域性變數。執行緒只有乙個核心物件和乙個棧，執行緒描述了程序內部的一條執行線路，每次初始化執行緒時，系統都會建立乙個主線程。系統從程序的位址空間中分配記憶體給執行緒棧使用，新執...
Windows系統程式設計（三） 執行緒
csdn部落格日期 2012年11月5日 handle createthread lpsecurity attributes lpthreadattributes,size t dwstacksize,lpthread start routine lpstartaddress,lpvoid lppa...
C 多執行緒程式設計 三 執行緒間通訊
七 執行緒間通訊 一般而言,應用程式中的乙個次要執行緒總是為主執行緒執行特定的任務,這樣,主線程和次要執行緒間必定有乙個資訊傳遞的渠道,也就是主線程和次要執行緒間要進行通訊。這種執行緒間的通訊不但是難以避免的，而且在多執行緒程式設計中也是複雜和頻繁的，下面將進行說明。使用全域性變數進行通訊 由於屬於...