并行課件備注

5-10不需別人用pthread_join清理門戶，自己了斷了。由于線程也占用資源，如果你不設置成為detach狀態(tài)，那么當你的線程推出后，你必須執(zhí)行

pthread_join調用才能釋放這些被占用的資源，如果設置成detach狀態(tài)，線程再推出后將自動釋放自己占用的資源

這些占用的資源不包括使用malloc分配的內存和ipc資源

5-39其實函數(shù)的執(zhí)行過程非常簡單，在第一個線程執(zhí)行到pthread_cond_wait(&cond,&mut)時，此時如果X<=Y，則此函數(shù)就將mut互斥量解鎖，再將cond條件變量加鎖，此時第一個線程掛起（不占用任何CPU周期）。

而在第二個線程中，本來因為mut被第一個線程鎖住而阻塞，此時因為mut已經釋放，所以可以獲得鎖mut，并且進行修改X和Y的值，在修改之后，一個IF語句判定是不是X>Y，如果是，則此時pthread_cond_signal()函數(shù)會喚醒第一個線程，并在下一句中釋放互斥量mut。然后第一個線程開始從pthread_cond_wait()執(zhí)行，首先要再次鎖mut，如果鎖成功，再進行條件的判斷（至于為什么用WHILE，即在被喚醒之后還要再判斷，后面有原因分析），如果滿足條件，則被喚醒進行處理，最后釋放互斥量mut。

至于為什么在被喚醒之后還要再次進行條件判斷（即為什么要使用while循環(huán)來判斷條件），是因為可能有“驚群效應”。有人覺得此處既然是被喚醒的，肯定是滿足條件了，其實不然。如果是多個線程都在等待這個條件，而同時只能有一個線程進行處理，此時就必須要再次條件判斷，以使只有一個線程進入臨界區(qū)處理。6-20Sectionsaredistributedamongthethreadsintheparallelteam.Eachsectionisexecutedonlyonceandeachthreadmayexecutezeroormoresections.It’snotpossibletodeterminewhetherornotasectionwillbeexecutedbeforeanother.Therefore,theoutputofonesectionshouldnotserveastheinputtoanother.Instead,thesectionthatgeneratesoutputshouldbemovedbeforethesectionsconstruct.6-22DataScopeAttributesAlldataclausesapplytoparallelregionsandworksharingconstructsexcept“shared,”whichonlyappliestoparallelregions.6-23PrivateCauseFor-loopiterationvariableisPRIVATEbydefault.6-29AtomicConstructSinceindex[i]canbethesamefordifferentIvalues,theupdatetoxmustbeprotected.Useofacriticalsectionwouldserializeupdatestox.Atomicprotectsindividualelementsofxarray,sothatifmultiple,concurrentinstancesofindex[i]aredifferent,updatescanstillbedoneinparallel.7-4We’vealreadydefinedspeeduptobethesequentialexecutiontimedividedbytheparallelexecutiontime.Wejustpluginthenumeratorandthedenominatorfromthepreviousslide.Sincethevalueinthedenominatorisalowerbound,thequotient(i.e.,thespeedup)isanupperbound.7-2首先我們介紹一下為什么要使用gpu進行計算，gpu計算比傳統(tǒng)的cpu計算好在哪里。然后是gpu的基本架構，然后介紹一下有哪些常用的調用gpu的方法。后兩部分使我們的重點，首先介紹一下cuda的編程模型，這一章有4個小節(jié)，講到這的時候再細說。然后就是怎么用cuda編程，我們在這里介紹一些基本的編程方法，還有一些比較高級的用法，如果有時間的話就將，沒時間的話，大家看一下文檔。7-3三個整數(shù)數(shù)組，我們想要計算A[i]+B[i]然后將結果存入C[i]中，我們首先看一下傳統(tǒng)的編程方法7-4這是一個傳統(tǒng)的做法，需要迭代N次，時間復雜度O(N)。在這個循環(huán)中，我們發(fā)現(xiàn)，A[i]+B[i]和A[i-1]+B[i-1]是沒有關系的。這就意味著這些運算是可以并行處理的。這時候我們就可以用多線程。7-6一些大型的應用程序，例如計算化學，用來計算分子間是如何相互作用的，還有天氣和氣候的模擬程序，等等一些其它的大型程序，它們要處理的數(shù)據量都是百萬，上億級別的數(shù)據。需要數(shù)千個線程同時執(zhí)行，這已經遠遠超過了cpu的負載能力。而且操作系統(tǒng)也不支持。我們需要一個專門的硬件來處理這種大規(guī)模的數(shù)據。而gpu正是一個非常好的選擇。因為gpu本來是為游戲設計的，而游戲中最多的操作就是圖像的矩陣運算，而且是大規(guī)模的矩陣運算，和我們要處理的數(shù)據非常類似，因此gpu也就有了另外一個功能，用于高性能計算。英偉達也開發(fā)了cuda用于gpu編程。下面介紹一下gpu的基本架構。7-8一個cpu可以有2個，4個，8個，甚至更多個核，cpu可以做任何類型的計算，串行的，并行的，各種IO操作，能夠進行指令預取，指令流水，分支預測，亂序執(zhí)行等等一些其它的功能，在同樣的芯片面積上，gpu去掉了一些cpu中的功能，或者簡化了一些cpu中的功能，取而代之的是更多的計算核心，一個gpu一般都含有幾千個核，能夠同時運行10000個線程，因此，GPU是專門用來處理大規(guī)模的計算密集型程序。這些程序能夠高度的并行化。運行在gpu上會獲得非常高的加速比。7-9這個是英偉達特斯拉顯卡核心GP100的架構圖，這個核心由6個GPC組成，這些概念不重要，了解一下就行。每個GPC包含5個TPC，每個TPC包含2個SM單元，這里面最重要的就是SM概念。整個gp100核心包含了60個SM單元，這里展示的是一個完整的gp100核心，總共包含60個SM單元，不同的產品可能會有不同個數(shù)的SM單元。對于gp100核心的產品，最多是60個SM單元。下面詳細介紹一下SM單元。7-10每一個sm單元被分成了兩個處理器塊，每個處理器塊包含了32個單精度CUDA核（圖中綠色方塊）、一個指令緩存、一個warp調度（后面會介紹warp的概念）、兩個指令分發(fā)單元8個ld/st單元，用來計算訪存地址，8個sfu單元，specialfunctionunit，特殊函數(shù)單元，例如cos，sin，平方根等等。合起來一個SM單元里總共包含64個單精度浮點cuda核心，因此，也就有32個雙精度浮點cuda核心。整個gp100核心總過有3840個單精度和1920個雙精度cuda核心。一個gp100核心，它的雙精度浮點的峰值運算速度為5300GFLOPs。以上是gp100核心的一個簡單介紹，更詳細的架構信息可以再官方的白皮書中找到，在百度上搜pascal-architecture-whitepaper7-11下面我們介紹一下三種不同的調用gpu的方法7-12這是三種常用的調用gpu的方法，cudaoptimizedlibraries,openacc,programminglanguages，下面我們詳細的介紹一下7-13為了讓gpu能夠被廣泛的應用，英偉達用cuda重新編寫了一些常用的庫函數(shù)，例如BLAS(BasicLinerAlgebraSubprograms基本線性代數(shù)程序,FFTW(FastFourierTransformintheWest)快速傅里葉變換程序，等等一些其它的數(shù)學函數(shù)庫，雖然英偉達重寫了這些函數(shù)庫，但是他們沒有修改這些函數(shù)的調用接口，因此我們不需要要修改源程序，只需要在編譯程序時告訴編譯器使用cuda數(shù)學函數(shù)庫，而不是普通的數(shù)學函數(shù)庫。這樣我們編譯出來的程序就能調用gpu了。這種方法是最簡單的，但同時也是限制最多的，因為很多時候我們要做的計算不是一個數(shù)學函數(shù)就能解決的。所以這種方法非常有局限性，但是如果你的程序中需要用到數(shù)學函數(shù)，可以考慮使用cuda的數(shù)學函數(shù)庫。這個不是我們的重點。7-14openacc是一個基于預編譯指令的編程模型，我們需要在代碼中插入一些openacc的預編譯指令來指導編譯器進行并行化處理，因此，我們需要專門的編譯器，openacc來編譯程序。看一下左邊的例子，源程序就是一個簡單的for循環(huán)結構。我們如果想讓這個循環(huán)并行執(zhí)行呢，就需要插入圖中的那些預編譯指令，openacc會自動的對這些結構進行并行化處理。這種方法比較實用于那些已經完成的代碼，如果用cuda將這些代碼重寫一遍，會非常的耗時耗力，那么通過openacc這種自動化的并行處理，能夠省去很多的人力物力。但是這種方法只適用于一些簡單的結構，對于復雜的結構，例如dowhile循環(huán)，openacc就無能為力了，只能靠人來進行手工并行化了。這個也不是我們的重點。7-15cuda支持的編程語言很多，第一個是cuda對Fortran語言的擴展，雖然Fortran語言很古老，比你們熟悉的C語言還要早10幾年，但是在科學計算領域，fortran一直都被廣泛地應用著。第二個是對腳本語言python的擴展，在科學計算領域應用的比較少，第三個是opencl，opencl是為異構編程而設計的，什么是異構編程呢，假如我們有多個計算設備，例如cpu，英偉達的gpu，或者是amd的gpu，如何將這些設備統(tǒng)一起來呢，opencl便提供了一個標準，這個標準規(guī)定了軟硬件api的規(guī)范，但它不提供具體的實現(xiàn)，每個廠商根據這個規(guī)范來編寫具體的實現(xiàn)代碼。程序員只需要調用統(tǒng)一的api就行，不同的設備會調用不同的實現(xiàn)來完成。第四個是cuda對C和C++的擴展。我們后面只講如何利用c和c++進行cuda編程。其它的有興趣的可以自己去看。7-17在這一章，我們會介紹cudacc++的編程模型。首先我們介紹一下cuda編程時需要用到哪些軟硬件環(huán)境，然后在介紹cuda程序的執(zhí)行過程，cpu和gpu之間是如何協(xié)同工作的。然后我們詳細講一下cuda中線程的概念，以及和線程相關的一些其它重要的概念，這些都是和后面的編程緊密相關的。在這一章的最后，我們介紹一下cuda的內存模型，訪存的速度一直都是程序的瓶頸，要想寫出高效的程序，就一定要了解設備的存儲結構，充分的利用cache，提高程序的效率。7-18要想利用cuda進行編程，首先，你要有一塊英偉達的顯卡，AMD的不行，cuda不支持。這塊顯卡可以是專門的gpu計算顯卡，例如我們前面介紹的TeslaPascalGP100，不過這種專門計算的顯卡一般都非常貴，普通用戶是買不起的，普通用戶買來也沒什么用，這種gpu計算卡一般都是給高性能計算的集群使用的。但這并不意味著我們不能用cuda編程了，普通用戶可以利用英偉達的游戲顯卡來進行cuda編程，例如我們筆記本里面的GT系列的顯卡，還有更好的GTX系列的顯卡都可以用來進行gpu編程，但是一些太老的顯卡呢，就不行了。如果你想看一下你的顯卡是否支持cuda編程，可以上這個網站，然后點擊cuda-enabledgeforceproducts，如果在里面找到了你的顯卡型號，那么你的顯卡就可以用來cuda編程。7-19利用cuda編程，我們需要安裝cudatoolkit，這個軟件包會為我們安裝一些cuda程序編譯，運行會用到的庫，它還會更新顯卡驅動，以便能夠運行cuda程序。如果你的操作系統(tǒng)是windows的話，要想編譯cuda程序，你可能還要安裝visualstudio，vs是對cuda支持最好的ide。如果你不想安裝vs，你也可以直接從命令行調用cuda的編譯器nvcc。Nvcc.exe可能沒在你默認的系統(tǒng)路徑里面，你需要自己找到nvcc.exe的位置，一般在它的安裝目錄下都能找到。如果是在linux平臺的話，安裝完cudatoolkit之后，可以在終端里直接敲命令nvcc來編譯cuda程序。7-21一個cuda程序主要是由兩部分組成，串行部分和并行部分。串行部分在cpu上執(zhí)行，并行部分在gpu上執(zhí)行。大多數(shù)時候，我們稱cpu為host，gpu為device。串行部分主要進行邏輯控制，例如if語句，switch語句，還有輸入輸出，包括讀取數(shù)據，向屏幕顯示信息。并行部分最主要的工作就是是計算。例如矩陣相乘。Cuda程序在運行的過程中，普通指令就在cpu上執(zhí)行，當遇到cuda指令時，便將代碼和數(shù)據發(fā)送到gpu上，然后調用gpu進行進行計算。Gpu計算完成后，返回到cpu中，cpu繼續(xù)執(zhí)行。這個過程和普通的函數(shù)調用沒有什么區(qū)別，只不過被調用的函數(shù)是在gpu上執(zhí)行的。這是一個大致的cuda程序的執(zhí)行過程，下面我們看一下具體是如何工作的。7-22當cpu要調用gpu進行計算時，首先要將程序和數(shù)據拷貝到gpu的內存中。因為Gpu在進行計算時，需要從自己的內存中讀數(shù)據。因為gpu是通過pcie接口和cpu相連的，如果gpu直接從cpu的內存中讀數(shù)據，需要不斷的通過pcie來進行數(shù)據傳輸。這個過程消耗的時間要遠大于我們直接將數(shù)據拷貝到gpu內存的時間。而且也會造成gpu要等待數(shù)據這樣的時間浪費。因此，我們需要現(xiàn)將數(shù)據通過pcie拷貝到gpu的內存中，然后進行計算。7-23準備好數(shù)據之后，cpu發(fā)送指令，命令gpu開始進行計算，計算時，gpu就不需要和cpu進行通信了。直接在自己的內存中讀寫數(shù)據。7-24Gpu計算完之后，便將結果從自己的內存中拷貝到cpu內存中。具體的執(zhí)行過程會在后面體現(xiàn)出來。下面我們進入cuda編程中最重要的一部分，cudathreads。7-25與傳統(tǒng)的cpu線程相比，cuda線程更加的輕量化，能夠快速的創(chuàng)建幾千個線程，并且能夠快速地進行上下文切換，因此，盡量讓cuda線程做計算工作，讓cpu做邏輯工作。7-26在介紹cudathreads之前，先普及以下cudakernel地概念.程序中可以并行執(zhí)行地部分成為cudakernel，一般情況下，cudakernel都是一個函數(shù)，函數(shù)里面時可以并行執(zhí)行地代碼。Cpu調用這個函數(shù)，函數(shù)里面地cudaapi便在gpu上執(zhí)行。Gpu上所有地線程都執(zhí)行相同地代碼，但是可以選擇不同地路徑，比如遇到分支語句，可能偶數(shù)線程和奇數(shù)線程執(zhí)行地時不同地代碼。和cpu線程一樣，每個cuda線程都有一個唯一地標識。右圖是一個簡單的cudakernel例子。我們有四個線程，threadIdx.x表示每個線程地id，首先在輸入數(shù)組中取出數(shù)據，0號線程取得第0號數(shù)據，1號線程取得第1號數(shù)據，依次類推。每個線程調用func函數(shù)來處理數(shù)據，然后將結果存到輸出數(shù)組的相應位置。這是一個簡單的例子，我們只用到了四個進程，真正運行在gpu上的程序都會用到幾百個，上千個線程。為了高效地管理這些線程，Cuda采用了層次化的管理方式。7-27首先將線程分為warp，每32個線程分為一個warp，一個warp是cuda任務調度，程序執(zhí)行的最小單元。AwarpinCUDA,then,isagroupof32threads,whichistheminimumsizeofthedataprocessedinSIMDfashionbyaCUDAmultiprocessor.但是在編程時，我們不需要考慮warp。我們需要考慮的是blocks。一個或者多個warps構成了一個block7-28程序執(zhí)行時，至少使用一個warp，既32個線程。這是cuda里最小的調度和執(zhí)行單位。如果你的程序連32線程都不到，那么cpu完全能滿足你的需求，不需要將程序遷移到gpu上，這樣反而會使程序運行變慢。在一個warp里面的線程的線程號是連續(xù)遞增的。但是我們編程時并不會直接操作warp，我們操作的是block。7-29一個block內的線程可以通過共享內存來進行數(shù)據交換和同步操作，也可以通過調用同步api來進行同步。相比之下，通過共享內存來同步會更快一些。由于硬件資源的限制呢，一個線程塊最多可以包含1024個線程每個線程塊內的新城id都是從0開始的。并且在這個線程塊內是唯一的7-30一個或多個block構成了一個grid。具體包含多少個block，是在編程時指定的，每個block里包含多少個線程也是在編程時指定的，但必須是32的整數(shù)倍。因為block是由warp組成的。當我們提交一個kernel到gpu上運行時，需要指定一個grid里包含多少個block，一個block里包含多少個thread。然后整個kernel被當作一個grid載入到gpu上運行。好了，到現(xiàn)在我們已經了解了cudakernel，線程，線程塊，grid，這些概念了，那么接下來我們看一個更加具體的kernel的執(zhí)行過程。7-31當我們向gpu提交了一個kernel時，這個kernel被看做是一個grid，假設我們指定Grid里面有8個block，如果你的GPU含有兩個SM單元，向SM單元分配任務時，都是直接將一個block分配給一個SM單元。那么在這個例子里面，每個SM單元會被分配4個block，每個SM單元中的block順序執(zhí)行。而不同SM單元中的block并行執(zhí)行。Block0和block1就是并行執(zhí)行的。7-32如果我們有4個SM單元，每個SM單元執(zhí)行兩個block，每個SM單元內的block順序執(zhí)行。我們知道block是由warp組成的，那么這些warp是如何執(zhí)行的呢？7-33先回憶一下SM單元的結構。一個SM單元里包含了兩個warp調度器，每個warp調度器又有兩個指令分發(fā)單元，每個warp調度器還對應著32個單精度的浮點cuda核。當一個SM單元被分配了一個block執(zhí)行時，7-34當一個block在SM單元上運行時，它首先被分解成多個warp，然后warp調度器來決定哪個warp可以執(zhí)行。一個SM單元每次可以同時執(zhí)行兩個warp，當一個warp被調度運行時，warp中的所有線程都會執(zhí)行同一個指令，如果執(zhí)行到一條分支指令的話，分支指令的所有路徑都會被順序執(zhí)行，不符合當前路徑的線程會處于阻塞狀態(tài)。我們看一下下面的例子，7-36如果我們的kernel中有一段代碼，用來區(qū)別偶數(shù)進程和奇數(shù)進程。那么這段代碼會在block中差生兩個執(zhí)行路徑，一個用來執(zhí)行奇數(shù)進程，一個用來執(zhí)行偶數(shù)進程，如下圖所示7-37假設warp里面有8個線程，這8個線程執(zhí)行上面的代碼，從開始執(zhí)行到分支語句之前，這8個進程每次都是執(zhí)行相同的指令，但是當執(zhí)行到分支語句的時候，就出現(xiàn)問題了，這個if語句一共有兩個分支，它首先執(zhí)行第一個分支，既偶數(shù)進程分支，那么這8個進程中一共有4個進程符合次分支，因此，這4個進程首先執(zhí)行，而另外4個進程則等待。當這4個進程執(zhí)行結束后。再執(zhí)行第二個分支，既奇數(shù)進程分支，這時，另外4個處于等待的進程開始執(zhí)行，而首先執(zhí)行的那4個進程則進入等待狀態(tài)。當奇數(shù)進程執(zhí)行完后，這8個進程再同時執(zhí)行后面的指令。這中情況會造成50%的性能損失。那么如何避免呢。7-38這段代碼依然會在block產生兩個分支，但是與前一個不同的是，同一個warp內的線程，它們會選擇相同的路徑，這樣在一個warp內，就不產生branchdivergence了。當然了，實際情況可能不會這么簡單，我們可能還需要改變原有的數(shù)據結構和算法，這些都需要根據實際情況來定。但是我們遵循的一個原則是，讓warp內的線程做同樣的工作。以上便是cuda線程的所有內容了。下面我們介紹cuda的內存模型7-39與傳統(tǒng)的cpu線程相比，cuda線程更加的輕量化，能夠快速的創(chuàng)建幾千個線程，并且能夠快速地進行上下文切換，因此，盡量讓cuda線程做計算工作，讓cpu做邏輯工作。7-40GPU的全局內存是GPU的主要的存儲器，之所以是全局的，主要是因為GPU與CPU都可以對它進行寫操作。任何設備都可以通過PCI-E總線對其進行訪問。全局內存的功能類似于C語言程序中的堆。cudaMalloc()hastwoparameters:AddressofapointertotheallocatedobjectSizeofallocatedobjectintermsofbytesTheaddressofthepointervariableshouldbecastto(void**)becausethefunctionexpectsagenericpointer;thememoryallocationfunctionisagenericfunctionthatisnotrestrictedtoanyparticulartypeofobjectscudaFree()hasoneparameter:Pointertofreedobject7-41Oncethehostcodehasallocateddevicememoryforthedataobjects,itcanrequestthatdatabetransferredfromhosttodevice.ThisisaccomplishedbycallingoneoftheCUDAAPIfunctions,cudaMemory().PleasenotecudaMemcpycurrentlycannotbeusedtocopybetweendifferentGPU’sinmultiGPUsystems右圖是CPU和GPU之間傳輸關系圖，可以看出來，CPU和GPU之前傳輸速度相對很差。GPU和GPUmemory傳輸速度要快得多，所以，對于編程來說，要時刻考慮減少CPU和GPU之間的數(shù)據傳輸。7-42常量內存其實是全局內存的一種虛擬地址形式，并沒有特殊保留的常量內存塊。常量內存有幾個特性，第一個是高速緩存，第二個時只讀，第三個是它支持將單個值廣播到線程束中的每個線程。常量內存的大小比較小，一般被限制為64KB。常量內存的聲明方式有兩種，一種是在編譯時聲明，需要用到“__constant__”關鍵字；另一種是在運行時通過主機端定義為只讀內存，使用cudaMemcpyToSymbol函數(shù)。7-43textureMemory駐留在deviceMemory中，并且使用一個只讀cache（per-SM）。textureMemory實際上也是globalMemory在一塊，但是他有自己專有的只讀cache。這個cache在浮點運算很有用。textureMemory是針對2D或3D空間局部性的優(yōu)化策略，所以thread要獲取2D或3D數(shù)據就可以使用textureMemory來達到很高的性能.Globalmemory沒有Cache，訪問速度很慢，Sharedmemory訪問速度很快，但是容量很小，對于較大的數(shù)組，將其綁定至texturememory往往是個不錯的選擇。Texturememory可以cache，而且容量很大。7-44Sharedmemory可以用于block內線程之間的數(shù)據共享。Sharedmemory實際上是可受用戶控制的一級緩存。每個SM中的一級緩存與Sharedmemory共享一個64KB的內存段。其訪問速度僅次于registers，延遲較低。需要注意的是SM=Streamingmultiprocessor而不是SharedMemory。Sharedmemoryisanefficientmeansforthreadstocooperatebysharingtheirinputdataandintermediateresults.Canallocatesharedmemorystatically(sizeknownatcompiletime)ordynamically(sizenotknownuntilruntime)each“__’’consistsoftwo“_’’characters.Onecanalsoaddanoptional“__device__”infrontof“__shared__”inthedeclarationtoachievethesameeffect.7-45這里是計算強度的計算方法，所謂的計算強度，就是浮點計算次數(shù)/IO次數(shù)，就是平均每個數(shù)據所參與的浮點計算操作的次數(shù)。當計算強度>1時，說明每個數(shù)據參與超過一個浮點計算。這種情況下，就需要盡量使用sharedmemoryload，來減少訪存延遲。7-46這兩種內存屬于不能操作的內存，是由一套自動機制來達到很好的性能。7-47這里我們對CUDA內存模型進行總結。包括每一類內存的位置、訪問權限、變量的生存周期等。其中registers和localmemory是由編譯器控制和分配（Non-programmable），而shared，global，constant和texturememory可受程序員控制(Programmable)。Registers和local是On-chip內存，Loal,global,constant,和texturememory是設備內存。7-48下面我們開始將cuda編程的基本語法，其實和c語言的語法一樣，只是增加了一些cuda的api，7-49在這一章，我們只介紹一些基本的cuda編程的api，以及一些特殊內存的使用方法。7-50我們看第一個例子，數(shù)組求和。有兩個數(shù)組，a，b，將他們對應的元素相加，然后將結果存入c中。右圖所示的是一個傳統(tǒng)的單線程的編程方法。這是一個簡單串行的例子，可以看到程序中做主要工作的是add函數(shù)，并行化的工作也就集中在了add函數(shù)上。我們首先考慮一個雙核的cpu，如何讓兩個核同時工作呢。7-51其中一個方法是，一個核處理奇數(shù)索引的數(shù)據，另一個核處理偶數(shù)索引的數(shù)據。這兩個代碼只是兩個例子，實際寫多線程代碼時是不一樣。如果我們需要處理的數(shù)據量太大了，需要用到幾千個核，這時，我們需要將這個代碼改成cuda程序了。7-52我們首先看一下main函數(shù)要如何修改。我們先說一下大致過程，cudaapi的具體信息后面再介紹。首先是定義三個數(shù)組，然后時三個指針，每個指針都以dev開頭，很明顯這三個指針是要在gpu上用到的。不是一定要以dev開頭，只要符合c語言命名規(guī)則的都可以，以dev開頭是為了和cpu上的變量區(qū)分開。下一條語句，很明顯是cudaapi了，一般情況下cudaapi都會以cuda開頭。這個api是用來在gpu上開辟內存空間的。A,b,c這三個數(shù)組不僅在cpu上會用到，在gpu上也會用到，因此我們也需要在gpu上開辟空間，要注意的是于c語言的malloc不同，cuda的malloc函數(shù)需要多傳遞一個指針的指針最為參數(shù)，具體語言后面會說到。為b，c數(shù)組開辟空間的語句直接省略了。下一條語句是初始化，初始化只需要做一次就行，既可以在cpu上完成初始化，也可以在gpu上完成初始化。如果在gpu上初始化會更快一些，但是這不是我們關注的地方。所以我們在cpu上完成初始化。在cpu上完成初始化之后，我們需要將數(shù)據拷貝到gpu中，cudamemcpy就是用來再cpu和gpu之間傳遞數(shù)據的，最后一個參數(shù)是用來指定傳輸方向的，從cpu到gpu或者從gpu到cpu。下一條語句就是告訴gpu要再gpu上執(zhí)行哪個函數(shù)。我們這個例子中是add函數(shù)。我們前面講過一個kernel被當作一個grid來執(zhí)行?，F(xiàn)在呢，add函數(shù)就是一個kernel，尖括號里的內容就是grid的大小，第一個參數(shù)N指定了一個grid有多少個block，第二個參數(shù)1指定了每個block里有幾個線程。雖然我們指定了每個block里有一個線程，但是由于warp是執(zhí)行和調度的最小單元，因此，每個SM單元還是會開啟32個線程，只不過，32個線程只有一個做真正的任務，其余的一直處于等待狀態(tài)。這里我們指定每個block里有一個線程只是一個例子，你們寫真正的cuda代碼時千萬不要這樣寫，盡量讓每個block里的線程數(shù)能被32整除。在下一條語句是將結果從gpu拷貝到cpu中。最后一條語句是釋放gpu上的空間，b，c也要釋放，我這里注釋掉了。下面我們看一下這幾個cudaapi的具體使用方法。7-53C語言中的malloc函數(shù)會返回一個指針，但是因為cuda的api全都會返回一個錯誤碼，所以只能傳遞一個指針的指針，用來存儲開辟的內存空間的地址。如果只有__device__修飾符，表明了這個函數(shù)只能在gpu上調用，并且只能運行在gpu上，如果只有一個__host__修飾符，表明這個函數(shù)只能在cpu上調用和執(zhí)行。7-54在gpu上運行add函數(shù)時需要加幾個尖括號，這幾個尖括號被稱為executionconfiguration。Dg，db，都是3維數(shù)據結構，dg指明了grid的三個維度的大小，db指明了block三個維度的大小，Ns指明了每個block使用的共享內存是多少，這是一個可選參數(shù)，默認是0。S指明了與這個kernel相關的cudastream是什么，默認是0，cudastream是比較高級的內容，我們不會講到，感興趣的可以自己在網上查一下。我們重點看一下dim3這個類型，我們在調用add函數(shù)時，只傳遞了兩個整形變量，N和1.這時候，編譯器會自動將這兩個整型變量轉換成兩個dim3類型變量，分別代表grid和block的維度，grid的x維賦予N，其它兩個維度都是1，block的x維度賦予1，其它兩個維度都是1.如果我們不想讓其它兩個維度是一，可以用下面的代碼。7-55聲明兩個dim3的變量，然后對它們的各個維度分別賦值。7-56Cudamemcpy只有一個host修飾符，意味著這個cuda函數(shù)只能在cpu端調用。它不僅能在cpu和gpu之間傳遞數(shù)據，還能在cpu內存上進行數(shù)據復制，也能在gpu內存上進行數(shù)據復制。只要指定傳輸?shù)姆较蚣纯?。下面我們看一下add函數(shù)是如何定義的。7-57我們看到add函數(shù)定義的時候多了一個__global__修飾符，global修飾符表明了這個函數(shù)要在gpu上執(zhí)行，但是可以在cpu上調用。一般kernel函數(shù)都要加上__global__修飾符。使用global修飾符對函數(shù)也有一些要求。首先，global函數(shù)的返回值必須是void，其次，在調用global函數(shù)的地方必須指明它的executionconfiguration。最后，global函數(shù)的調用是異步的，意思是，一旦我們將這個函數(shù)交給gpu執(zhí)行后，調用語句馬上返回，不會等到gpu執(zhí)行完才返回。我們看一下add函數(shù)內部。7-58Add函數(shù)里用到了一個變量blockIdx.x，這是由cuda定義的全局變量，指明了當前線程所屬的block的索引，除了blockidx.x外，還有blockidx.y和blockidx.z，分別指明了x,y,z這三個維度的索引。除了blockidx還有threadidx，指明了當前thread在block內的索引。同樣也有三個，分別是threadidx.x,threadidx.y,threadidx.z。在我們這個例子里面呢，grid只有一維，其它兩個維度的大小都是一。而每個block里只有一個線程，所以blockidx.x的索引就相當于是線程的索引了。Add函數(shù)很簡單，對于小于N的線程做加法操作。為什么寫小于N呢，雖然我們可以肯定tid不會超過N，但是有可能由于我們的疏忽tid超過N了，這會在gpu上造成數(shù)組訪問越界。如果你確定不會超過N，也可以不加這個判斷。這一節(jié)到這里就介紹完了，下面我們講一下如何使用gpu的共享內存和同步操作。7-62下面我們詳細的講一下這個過程，假設我們由N個線程，兩個數(shù)組x，y的大小都是N*M，這兩個數(shù)組都是一維數(shù)組。這N個線程每次可以同時處理N組數(shù)據，我們一共有N*M組數(shù)據，總共需要M次迭代。第一次迭代，這N個線程分別處理的是x1*y1,x2*y2一直到xN*yN。第二次迭代就是從xN+1*yN+1開始，一直到x2N*y2N。每個進程迭代M次，迭代結束后，再將每個進程的累加和加起來，得到最后的結果。下面我們看一下具體的代碼怎么寫。7-63首先看一下main函數(shù)的結構，這是前半部分的代碼，一直到調用點乘函數(shù)。N代表了數(shù)據規(guī)模，blockspergrid表示grid里由多少個block，threadperblock表示每個block里由多少個thread。Main函數(shù)的結構呢和我們上一個數(shù)組相加的結構是一樣的，不同的是這個例子中的C數(shù)組的大小并不是N，而是blockspergrid，后面會講到為什么是這個，我們先看一下dot函數(shù)的結構7-64前面講過，cuda會為每個block來分配一塊共享內存，block里的所有線程共享這塊內存，因此，在dot函數(shù)的開頭，我們?yōu)閎lock里的每一個線程申請了一個共享內存，共享內存的總大小為threadsperblock。下一步是要計算線程的全局id，threadidx表示的是當前線程在block內的索引，我們取數(shù)據時需要用到當前thread在整個grid中的索引是多少，blockdim.x表示block的x維度的大小。整個計算過程就相當于將二維數(shù)組的坐標轉換成一維數(shù)組一樣，我們這里就不再說了。下一個是cacheindex，表示當前線程在共享內存中的位置，因為每個block有一塊共享內存，所以我們只需要threadidx來索引即可。下一步便是遍歷這個線程需要處理的數(shù)據，進行累加操作。以0號線程為例，第一次處理a[0],b[0]，第二次處理a[n],b[n],間隔是總的進程數(shù)，代碼里面總的進程數(shù)就是用blockdim.x*griddim.x來表示的，因為grid和block都是一維的。如果它們是二維或者三維的，還需要加上其它維度的大小。最后每個進程將累加和存入相應的共享內存中。下一步就是要對共享內存中的數(shù)據進行累加，計算出這個blcok的累加和是多少。因為只有當所有進程的數(shù)據都寫入共享內存之后，我們才能開始累加操作。所以在此處要有一個同步操作。只有當所有進程都到達這個同步點的時候，所有的進程再同時執(zhí)行下面的指令。執(zhí)行到這里的時候，是一個多線程的環(huán)境，我們可以充分的利用多線程來累加共享內存中的數(shù)據。下面我們看一下累加操作。7-65這是共享內存cache中的數(shù)據，假設cache一共有8個數(shù)據需要累加，也就是說我們由8個線程，每個線程對著一個數(shù)據，累加的過程就是前一半的線程將自己的數(shù)據和后一半的數(shù)據相加?，F(xiàn)在是8個數(shù)據，8個線程，第一次是迭代是前4個線程自己的4個數(shù)據和后4個數(shù)據相加，結果存入到前4個數(shù)據中。然后我們只需要累加4個數(shù)據了，不斷的重復次過程，直到結果都存在了第一個數(shù)據中。這里我們要注意需要同步一下。7-66Dot函數(shù)的最后一步操作，將cache[0]的數(shù)據存入數(shù)c中，7-67然后將數(shù)據傳輸?shù)絚pu端，在cpu上進行最后的累加操作。以上便是共享內存的使用。7-71Image中的每一個方格都是一個像素，我們假設每個像素都能發(fā)射出一個平行的光，當一束光遇到場景中的物體時，會產生一個或多個焦點，我們找到離image最遠的那個交點。7-72為了方便，我們場景中的物體全是球形物體7-73Hit函數(shù)用來計算這個表面與ox，oy這條射線有沒有相交，如果相交的話，交點的深度是多少，既交點到image的距離是多少。下面我們看一下沒有使用contantmemory的主程序7-74Cudaevent函數(shù)主要是用來計時的，cudaeventrecord函數(shù)里的0代表是是哪個cudastream，cudastream我們就不再講了。還有cpubitmap結構，這是一個像素矩陣，我們直到這是一個二維矩陣就夠了，這個不是cuda定義的。從代碼中可以看到，我們在cpu和gpu端分別定義了20個球體。For循環(huán)是對這20個球體進行初始化。7-75然后將初始化后的球體拷貝到gpu端。下面就是計算每個像素點與每個物體的交點了，kernel函數(shù)我們就不再展示了。我們看紅色框里的內容也是用來計時的，它和前面的計時的相對應，我們把這個用法記住就行了。這是普通的寫法，我們知道contantmemory是用來存儲常量的。那么程序里面哪些是不會改變的呢，球體這個數(shù)組，一旦我們初始化之后，球體數(shù)組就不再改變了。下面我們看一下怎么修改主程序，kernel函數(shù)不需要修改。7-76左邊的是修改后的代碼，右邊是原來的代碼，我們來看一下不同之處，我們在開頭多了一個constant的聲明，這表明這個球體數(shù)組存儲在了constantmemory上。并且constantmemory必須聲明維全局變量。Constantmemory的初始化還和以前一樣，當然了，為了更快的初始化，我們也可以在gpu上完成初始化。將constantmemory拷貝到gpu上時，要用特殊的拷貝函數(shù)，用法和傳統(tǒng)的cudamemcpy一樣，只不過cudamemcpytosymbol不需要指明拷貝的方向，因為只能從cpu端拷貝到gpu端。其它的沒有任何區(qū)別了，以上便是contantmemory的用法了。7-80我們用一個heatingmodel來說明如何利用texturememory7-81不同的顏色