并行課件備注

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

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

這些占用的資源不包括使用malloc分配的內(nèi)存和ipc資源

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

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

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