本筆記主要是閱讀：譚升的博客的 GPU編程（CUDA）

1_0 并行計(jì)算與計(jì)算機(jī)架構(gòu)

【CUDA 基礎(chǔ)】1.0 并行計(jì)算與計(jì)算機(jī)架構(gòu)

并行計(jì)算其實(shí)設(shè)計(jì)到兩個(gè)不同的技術(shù)領(lǐng)域：

• 計(jì)算機(jī)架構(gòu)（硬件）：生產(chǎn)工具
• 并行程序設(shè)計(jì)（軟件）：用工具產(chǎn)生各種不同應(yīng)用

1.1 并行性

寫(xiě)并行程序主要是分解任務(wù)，一般把一個(gè)程序看成是指令和數(shù)據(jù)的組合，當(dāng)然并行也可以分為這兩種：

• 指令并行
• 數(shù)據(jù)并行

我們的任務(wù)更加關(guān)注數(shù)據(jù)并行。

任務(wù)并行多出現(xiàn)在各種管理系統(tǒng)，比如我們天天用的支付系統(tǒng)，基本上每時(shí)每刻都有很多人在同時(shí)使用，這時(shí)候就需要后臺(tái)的處理能夠并行執(zhí)行這些請(qǐng)求，不然全國(guó)人民排隊(duì)，那就比春運(yùn)還熱鬧了。

1.1.1 把數(shù)據(jù)依據(jù)線程進(jìn)行劃分

數(shù)據(jù)并行程序設(shè)計(jì)，第一步就是把數(shù)據(jù)依據(jù)線程進(jìn)行劃分：

塊劃分，把一整塊數(shù)據(jù)切成小塊，每個(gè)小塊隨機(jī)的劃分給一個(gè)線程，每個(gè)塊的執(zhí)行順序隨機(jī)（關(guān)于線程的概念可以去看《深入理解計(jì)算機(jī)系統(tǒng)》）

thread12345

block	1 2 3	4 5 6	7 8 9	10 11 12	13 14 15

周期劃分，線程按照順序處理相鄰的數(shù)據(jù)塊，每個(gè)線程處理多個(gè)數(shù)據(jù)塊，比如我們有五個(gè)線程，線程1執(zhí)行塊1，線程2執(zhí)行塊2……線程5執(zhí)行塊5，線程1執(zhí)行塊6

thread123451234512345

block	1	2	3	4	5	6	7	8	9	10	11	12	13	14	15

1.2 計(jì)算機(jī)架構(gòu)

1.2.1 Flynn’s Taxonomy

劃分不同計(jì)算機(jī)結(jié)構(gòu)的方法有很多，廣泛使用的一種被稱(chēng)為佛林分類(lèi)法Flynn’s Taxonomy，他根據(jù)指令和數(shù)據(jù)進(jìn)入CPU的方式分類(lèi)，分為以下四類(lèi)：

分別以數(shù)據(jù)和指令進(jìn)行分析：

• 單指令單數(shù)據(jù)SISD（傳統(tǒng)串行計(jì)算機(jī)，386）
• 單指令多數(shù)據(jù)SIMD（并行架構(gòu)，比如向量機(jī)，所有核心指令唯一，但是數(shù)據(jù)不同，現(xiàn)在CPU基本都有這類(lèi)的向量指令）
• 多指令單數(shù)據(jù)MISD（少見(jiàn)，多個(gè)指令圍毆一個(gè)數(shù)據(jù)）
• 多指令多數(shù)據(jù)MIMD（并行架構(gòu)，多核心，多指令，異步處理多個(gè)數(shù)據(jù)流，從而實(shí)現(xiàn)空間上的并行，MIMD多數(shù)情況下包含SIMD，就是MIMD有很多計(jì)算核，計(jì)算核支持SIMD）

為了提高并行的計(jì)算能力，我們要從架構(gòu)上實(shí)現(xiàn)下面這些性能提升：

• 降低延遲：是指操作從開(kāi)始到結(jié)束所需要的時(shí)間，一般用微秒計(jì)算，延遲越低越好。

• 提高帶寬：帶寬是單位時(shí)間內(nèi)處理的數(shù)據(jù)量，一般用MB/s或者GB/s表示。

• 提高吞吐量：是單位時(shí)間內(nèi)成功處理的運(yùn)算數(shù)量，一般用gflops來(lái)表示（十億次浮點(diǎn)計(jì)算）。

  吞吐量和延遲有一定關(guān)系，都是反應(yīng)計(jì)算速度的：

  • 一個(gè)是時(shí)間除以運(yùn)算次數(shù)，得到的是單位次數(shù)用的時(shí)間–延遲；
  • 一個(gè)是運(yùn)算次數(shù)除以時(shí)間，得到的是單位時(shí)間執(zhí)行次數(shù)–吞吐量。

1.2.2 根據(jù)內(nèi)存劃分

計(jì)算機(jī)架構(gòu)也可以根據(jù)內(nèi)存進(jìn)行劃分：

分布式內(nèi)存的多節(jié)點(diǎn)系統(tǒng)：更大，通常叫做集群，就是一個(gè)機(jī)房好多機(jī)箱，每個(gè)機(jī)箱都有內(nèi)存處理器電源等一些列硬件，通過(guò)網(wǎng)絡(luò)互動(dòng)，這樣組成的就是分布式。

共享內(nèi)存的多處理器系統(tǒng)：單個(gè)主板有多個(gè)處理器，他們共享相同的主板上的內(nèi)存，內(nèi)存尋址空間相同，通過(guò)PCIe和內(nèi)存互動(dòng)。

注意：多個(gè)處理器可以分多片處理器，和單片多核（眾核many-core），也就是有些主板上掛了好多片處理器，也有的是一個(gè)主板上就一個(gè)處理器，但是這個(gè)處理器里面有幾百個(gè)核。

最后：

• CPU適合執(zhí)行復(fù)雜的邏輯，比如多分支，其核心比較重（復(fù)雜）
• GPU適合執(zhí)行簡(jiǎn)單的邏輯，大量的數(shù)據(jù)計(jì)算，其吞吐量更高，但是核心比較輕（結(jié)構(gòu)簡(jiǎn)單）

1_1 異構(gòu)計(jì)算與CUDA

【CUDA 基礎(chǔ)】1.1 異構(gòu)計(jì)算與CUDA

Abstract: 介紹異構(gòu)計(jì)算和CUDA概述，完成GPU輸出Hello world！

Keywords: 異構(gòu)計(jì)算，CUDA

---

1.1 異構(gòu)計(jì)算

同構(gòu)計(jì)算：是使用相同類(lèi)型指令集和體系架構(gòu)的計(jì)算單元組成系統(tǒng)的計(jì)算方式。

異構(gòu)計(jì)算：主要是指使用不同類(lèi)型指令集和體系架構(gòu)的計(jì)算單元組成系統(tǒng)的計(jì)算方式，常見(jiàn)的計(jì)算單元類(lèi)別包括CPU、GPU、DSP、ASIC、FPGA等。

異構(gòu)計(jì)算用簡(jiǎn)單的公式可以表示為“CPU+XXX”。

由于術(shù)業(yè)有專(zhuān)攻，CPU、GPU、DSP、ASIC、FPGA各有所長(zhǎng)，在一些場(chǎng)景下，引入特定計(jì)算單元，讓計(jì)算系統(tǒng)變成混合結(jié)構(gòu)，就能讓CPU、GPU、DSP、FPGA執(zhí)行自己最擅長(zhǎng)的任務(wù)。

GPU原本是不可編程的，或者說(shuō)不對(duì)用戶開(kāi)放的，然后就有hacker開(kāi)始想辦法給GPU編程，來(lái)幫助他們完成規(guī)模較大的運(yùn)算，于是他們研究著色語(yǔ)言或者圖形處理原語(yǔ)來(lái)和GPU對(duì)話。后來(lái)黃老板發(fā)現(xiàn)了這個(gè)是個(gè)新的功能啊，然后就讓人開(kāi)發(fā)了一套平臺(tái)，CUDA，然后深度學(xué)習(xí)火了，順帶著，CUDA也火到爆炸。

1.2 異構(gòu)架構(gòu)

上面這張圖能大致反應(yīng)CPU和GPU的架構(gòu)不同。

• 左圖：一個(gè)四核CPU一般有四個(gè)ALU，ALU是完成邏輯計(jì)算的核心，也是我們平時(shí)說(shuō)四核八核的核，控制單元，緩存也在片上，DRAM是內(nèi)存，一般不在片上，CPU通過(guò)總線訪問(wèn)內(nèi)存。
• 右圖：GPU，綠色小方塊是ALU，我們注意紅色框內(nèi)的部分SM，這一組ALU公用一個(gè)Control單元和Cache，這個(gè)部分相當(dāng)于一個(gè)完整的多核CPU，但是不同的是ALU多了，control部分變小，可見(jiàn)計(jì)算能力提升了，控制能力減弱了，所以對(duì)于控制（邏輯）復(fù)雜的程序，一個(gè)GPU的SM是沒(méi)辦法和CPU比較的，但是對(duì)于邏輯簡(jiǎn)單，數(shù)據(jù)量大的任務(wù)，GPU更高效，并且，注意，一個(gè)GPU有好多個(gè)SM，而且越來(lái)越多。

一個(gè)異構(gòu)應(yīng)用包含兩種以上架構(gòu)，所以代碼也包括不止一部分：

• 主機(jī)代碼：在主機(jī)端運(yùn)行，被編譯成主機(jī)架構(gòu)的機(jī)器碼；
  • 主機(jī)端代碼主要是控制設(shè)備，完成數(shù)據(jù)傳輸?shù)瓤刂祁?lèi)工作，
• 設(shè)備代碼：設(shè)備端的在設(shè)備上執(zhí)行，被編譯成設(shè)備架構(gòu)的機(jī)器碼；
  • 設(shè)備端主要的任務(wù)就是計(jì)算。

所以主機(jī)端的機(jī)器碼和設(shè)備端的機(jī)器碼是隔離的，自己執(zhí)行自己的，沒(méi)辦法交換執(zhí)行。

1.3 【函數(shù)】cudaDeviceReset()

/* *hello_world.cu */ #include<stdio.h> __global__ void hello_world(void) {printf("GPU: Hello world!\n"); }int main(int argc,char **argv) {printf("CPU: Hello world!\n");hello_world<<<1,10>>>();cudaDeviceReset();//if no this line ,it can not output hello world from gpureturn 0; }

如果沒(méi)有cudaDeviceReset()，則不能正常的運(yùn)行。

答：這句話包含了隱式同步，GPU和CPU執(zhí)行程序是異步的，核函數(shù)調(diào)用后成立刻會(huì)到主機(jī)線程繼續(xù)，而不管GPU端核函數(shù)是否執(zhí)行完畢，所以上面的程序就是GPU剛開(kāi)始執(zhí)行，CPU已經(jīng)退出程序了，所以我們要等GPU執(zhí)行完了，再退出主機(jī)線程。

1.4 CUDA程序一般步驟

一般CUDA程序分成下面這些步驟：

分配host內(nèi)存，并進(jìn)行數(shù)據(jù)初始化;

分配device內(nèi)存，并從host將數(shù)據(jù)拷貝到device上;

調(diào)用CUDA的核函數(shù)在device上完成指定的運(yùn)算;

將device上的運(yùn)算結(jié)果拷貝到host上(性能)

釋放device和host上分配的內(nèi)存。

2_0 CUDA編程模型概述（一）

【CUDA 基礎(chǔ)】2.0 CUDA編程模型概述（一）

Abstract: 本文介紹CUDA編程模型的簡(jiǎn)要結(jié)構(gòu)，包括寫(xiě)一個(gè)簡(jiǎn)單的可執(zhí)行的CUDA程序，一個(gè)正確的CUDA核函數(shù)，以及相應(yīng)的調(diào)整設(shè)置內(nèi)存，線程來(lái)正確的運(yùn)行程序。

Keywords: CUDA編程模型，CUDA編程結(jié)構(gòu)，內(nèi)存管理，線程管理，CUDA核函數(shù)，CUDA錯(cuò)誤處理

1.1 CUDA編程模型概述

CUDA編程模型為應(yīng)用和硬件設(shè)備之間的橋梁，所以CUDA C是編譯型語(yǔ)言，不是解釋型語(yǔ)言。

GPU中大致可以分為：

• 核函數(shù)
• 內(nèi)存管理
• 線程管理
• 流

從宏觀上我們可以從以下幾個(gè)環(huán)節(jié)完成CUDA應(yīng)用開(kāi)發(fā)：

領(lǐng)域?qū)?#xff1a;在領(lǐng)域?qū)?#xff08;也就是你所要解決問(wèn)題的條件）分析數(shù)據(jù)和函數(shù)，以便在并行運(yùn)行環(huán)境中能正確，高效地解決問(wèn)題。

邏輯層：分析設(shè)計(jì)完程序就進(jìn)入了編程階段，我們關(guān)注點(diǎn)應(yīng)轉(zhuǎn)向如何組織并發(fā)進(jìn)程，這個(gè)階段要從邏輯層面思考。

硬件層：通過(guò)理解線程如何映射到機(jī)器上，能充分幫助我們提高性能。

1.2 CUDA編程結(jié)構(gòu)

一個(gè)異構(gòu)環(huán)境，通常有多個(gè)CPU多個(gè)GPU，他們都通過(guò)PCIe總線相互通信，也是通過(guò)PCIe總線分隔開(kāi)的。所以我們要區(qū)分一下兩種設(shè)備的內(nèi)存：

• 主機(jī)：CPU及其內(nèi)存
• 設(shè)備：GPU及其內(nèi)存

注意這兩個(gè)內(nèi)存從硬件到軟件都是隔離的（CUDA6.0 以后支持統(tǒng)一尋址），我們目前先不研究統(tǒng)一尋址.

一個(gè)完整的CUDA應(yīng)用可能的執(zhí)行順序如下圖：

從host的串行到調(diào)用核函數(shù)（核函數(shù)被調(diào)用后控制馬上歸還主機(jī)線程，也就是在第一個(gè)并行代碼執(zhí)行時(shí)，很有可能第二段host代碼已經(jīng)開(kāi)始同步執(zhí)行了）。

1.3 內(nèi)存管理

內(nèi)存管理在傳統(tǒng)串行程序是非常常見(jiàn)的：

• 棧空間內(nèi)的內(nèi)存由機(jī)器自己管理；
• 堆空間由用戶控制分配和釋放；

CUDA程序同樣，只是CUDA提供的API可以分配管理設(shè)備上的內(nèi)存，當(dāng)然也可以用CDUA管理主機(jī)上的內(nèi)存，主機(jī)上的傳統(tǒng)標(biāo)準(zhǔn)庫(kù)也能完成主機(jī)內(nèi)存管理。

下面表格有一些主機(jī)API和CUDA C的API的對(duì)比：

標(biāo)準(zhǔn)C函數(shù)CUDA C 函數(shù)說(shuō)明

malloc	cudaMalloc	內(nèi)存分配
memcpy	cudaMemcpy	內(nèi)存復(fù)制
memset	cudaMemset	內(nèi)存設(shè)置
free	cudaFree	釋放內(nèi)存

2_1 CUDA編程模型概述（二）

【CUDA 基礎(chǔ)】2.1 CUDA編程模型概述（二）

當(dāng)主機(jī)啟動(dòng)了核函數(shù)，控制權(quán)馬上回到主機(jī)，而不是主機(jī)等待設(shè)備完成核函數(shù)的運(yùn)行。即：所有CUDA核函數(shù)的啟動(dòng)都是異步的，這點(diǎn)與C語(yǔ)言是完全不同的。

這一點(diǎn)我們上一篇文章也有提到過(guò)（就是等待hello world輸出的那段代碼后面要加一句）

1.1 限定符：__global__/__device__/__host__

---

限定符執(zhí)行調(diào)用備注

global	設(shè)備端執(zhí)行	可以從主機(jī)調(diào)用也可以從計(jì)算能力3以上的設(shè)備調(diào)用	必須有一個(gè)void的返回類(lèi)型
device	設(shè)備端執(zhí)行	設(shè)備端調(diào)用
host	主機(jī)端執(zhí)行	主機(jī)調(diào)用	可以省略

1.2 核函數(shù)編寫(xiě)時(shí)的限制

Kernel核函數(shù)編寫(xiě)有以下限制

只能訪問(wèn)設(shè)備內(nèi)存

必須有void返回類(lèi)型

不支持可變數(shù)量的參數(shù)

不支持靜態(tài)變量

顯示異步行為

1.3 錯(cuò)誤處理

所有編程都需要對(duì)錯(cuò)誤進(jìn)行處理，早期的編碼錯(cuò)誤，編譯器會(huì)幫我們搞定，內(nèi)存錯(cuò)誤也能觀察出來(lái)，但是有些邏輯錯(cuò)誤很難發(fā)現(xiàn)，甚至到了上線運(yùn)行時(shí)才會(huì)被發(fā)現(xiàn)。

例如我們代碼庫(kù)頭文件里面的這個(gè)宏：

#define CHECK(call)\ {\const cudaError_t error=call;\if(error!=cudaSuccess)\{\printf("ERROR: %s:%d,",__FILE__,__LINE__);\printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\exit(1);\}\ }

就是獲得每個(gè)函數(shù)執(zhí)行后的返回結(jié)果，然后對(duì)不成功的信息加以處理，CUDA C 的API每個(gè)調(diào)用都會(huì)返回一個(gè)錯(cuò)誤代碼，這個(gè)代碼我們就可以好好利用了，當(dāng)然在release版本中可以去除這部分，但是開(kāi)發(fā)的時(shí)候一定要有的。

1.4 nvcc編譯執(zhí)行

編譯指令：nvcc xxxx.cu -o xxxx

2_2 給核函數(shù)計(jì)時(shí)

【CUDA 基礎(chǔ)】2.2 給核函數(shù)計(jì)時(shí)

我們可以大概分析下核函數(shù)啟動(dòng)到結(jié)束的過(guò)程：

主機(jī)線程啟動(dòng)核函數(shù)

核函數(shù)啟動(dòng)成功

控制返回主機(jī)線程

核函數(shù)執(zhí)行完成

主機(jī)同步函數(shù)偵測(cè)到核函數(shù)執(zhí)行完

如下圖所示：

我們要測(cè)試的是2~4的時(shí)間。

1.1 用CPU計(jì)時(shí)

主要講了三種：

• clock()：對(duì)于核函數(shù)計(jì)時(shí)不準(zhǔn)確；
• gettimeofday：對(duì)于核函數(shù)計(jì)時(shí)較準(zhǔn)確；
• nvprof 工具：該工具不僅給出了kernel執(zhí)行的時(shí)間，還有其他cuda函數(shù)的執(zhí)行時(shí)間。三者之中，最準(zhǔn)確。

1.1.1 使用clock()

使用cpu計(jì)時(shí)的方法是測(cè)試時(shí)間的一個(gè)常用辦法，我們?cè)趯?xiě)C程序的時(shí)候最多使用的計(jì)時(shí)方法是：

clock_t start, finish; start = clock(); // 要測(cè)試的部分 finish = clock(); duration = (double)(finish - start) / CLOCKS_PER_SEC; // CLOCKS_PER_SEC這個(gè)宏，就是每秒中多少clocks，在不同的系統(tǒng)中值可能不同。

不準(zhǔn)確原因：

• clock()是個(gè)關(guān)鍵的函數(shù)，clock函數(shù)測(cè)出來(lái)的時(shí)間為進(jìn)程運(yùn)行時(shí)間，單位為滴答數(shù)(ticks)。必須注意的是，并行程序這種計(jì)時(shí)方式有嚴(yán)重問(wèn)題！如果想知道具體原因，可以查詢clock的源代碼（c語(yǔ)言標(biāo)準(zhǔn)函數(shù)）

1.1.2 使用gettimeofday()

具體使用方法：https://linuxhint.com/gettimeofday_c_language/

#include <sys/time.h> double cpuSecond() // 例子 {struct timeval tp;gettimeofday(&tp,NULL);return((double)tp.tv_sec+(double)tp.tv_usec*1e-6); }

gettimeofday()是linux下的一個(gè)庫(kù)函數(shù)，創(chuàng)建一個(gè)cpu計(jì)時(shí)器，從1970年1月1日0點(diǎn)以來(lái)到現(xiàn)在的秒數(shù)，需要頭文件sys/time.h.

不準(zhǔn)確原因：使用gettimeofday()函數(shù)只能測(cè)試1~5的時(shí)間，所以測(cè)試得到的時(shí)間偏長(zhǎng)。

1.2 用nvprof計(jì)時(shí)

CUDA 5.0后有一個(gè)工具叫做nvprof的命令行分析工具，(后面還要介紹一個(gè)圖形化的工具)。

nvprof的用法如下：

nvprof [nvprof_args] <application>[application_args]

2_3 組織并行線程

【CUDA 基礎(chǔ)】2.3 組織并行線程

1.1 線程ID計(jì)算

首先，需要知道block、grid三維布局中，x/y/z的先后順序：

1.2 網(wǎng)格大小如何確定？

當(dāng)我們確定好 線程塊大小為a時(shí)，我們應(yīng)該如何確定網(wǎng)格的尺寸呢？

答：

• 在一維中，假設(shè)我們需要啟動(dòng)共N個(gè)線程，線程塊尺寸為a。最容易想到的是，網(wǎng)格尺寸為：N/a。但是我們需要注意，這可能除不盡（如N=10,a=3），此時(shí)算出來(lái)的是，網(wǎng)格尺寸為3。如果按照結(jié)果來(lái)，我們實(shí)際上只啟動(dòng)了3*3=9個(gè)線程。所以正確的做法是：

  • $\frac{(N+(a?1))}{(a)}$，公式可以理解為：計(jì)算 ≥N 的a的最小倍數(shù)。

• 在二維中，假設(shè)啟動(dòng)共(Nx * Ny)個(gè)線程，線程塊尺寸為(x, y)。則其網(wǎng)格尺寸：

  • $(\frac{(N_x+(x?1))}{(x)},\frac{(N_y+(y?1))}{(y)})$ // 即，其每一維的計(jì)算方式都是一樣的。

• 三維中，以此類(lèi)推。

2.4 GPU設(shè)備信息

一個(gè)命令nvidia-smi，也可以查詢一些信息。

nvidia-smi是nvidia驅(qū)動(dòng)程序內(nèi)帶的一個(gè)工具，可以返回當(dāng)前環(huán)境的設(shè)備信息。具體參數(shù)使用使用查文檔。

一些函數(shù)：

• cudaGetDeviceCount()
• cudaGetDeviceProperties()
• 例子中，沒(méi)用到的函數(shù)：
  • cudaSetDevice()：設(shè)置GPU執(zhí)行所使用的設(shè)備。
  • cudaDriverGetVersion()：返回驅(qū)動(dòng)程序支持的CUDA的最新版本。
  • cudaRuntimeGetVersion()：返回CUDA運(yùn)行時(shí)版本。

__host__ __device__ cudaError_t cudaGetDeviceCount ( int* count )

返回可計(jì)算設(shè)備的數(shù)量。

Parameters

• count：返回計(jì)算能力大于或等于2.0的設(shè)備數(shù)量。

Returns

• cudaSuccess

__host__ cudaError_t cudaSetDevice ( int device )

設(shè)置GPU執(zhí)行所使用的設(shè)備。

Parameters

• device：活動(dòng)主機(jī)線程應(yīng)該在其上執(zhí)行設(shè)備代碼的設(shè)備。

returns

• cudaSuccess
• cudaErrorInvalidDevice
• cudaErrorDeviceAlreadyInUse

查詢?cè)O(shè)備步驟：

使用cudaGetDeviceCount()來(lái)獲得系統(tǒng)中支持CUDA架構(gòu)的設(shè)備數(shù)目；

調(diào)用第一步的函數(shù)后，我們可以對(duì)每個(gè)設(shè)備進(jìn)行迭代。CUDA運(yùn)行時(shí)將返回一個(gè)cudaDeviceProp類(lèi)型的結(jié)構(gòu)（具體結(jié)構(gòu)，詳見(jiàn)），其中包含了設(shè)備的相關(guān)屬性。

所以，第二步應(yīng)該：

• 創(chuàng)建一個(gè)cudaDeviceProp類(lèi)型的變量。如：cudaDeviceProp prop;
• 然后使用函數(shù)cudaGetDeviceProperties ( cudaDeviceProp* prop, int device )來(lái)獲得設(shè)備信息。

在執(zhí)行完第2步后，該設(shè)備的信息以存入變量 prop中。所以，我們可以通過(guò)該變量來(lái)獲取設(shè)備具體信息。

// 詢問(wèn)設(shè)備信息的主要代碼為： int deviceCount; cudaGetDeviceCount(&deviceCount);cudaDeviceProp prop;for(int i =0; i<count; i++) {cudaGetDeviceProperties(&prop, i); // 將第i個(gè)設(shè)備的信息，存入變量prop中。// 接下來(lái)便可以通過(guò)變量訪問(wèn)第i個(gè)設(shè)備的信息了，如：printf("設(shè)備名字：\n",prop.name);.......更多屬性，可查閱cudaDeviceProp類(lèi)型的結(jié)構(gòu) }

3_1 CUDA執(zhí)行模型概述

https://face2ai.com/CUDA-F-3-1-CUDA%E6%89%A7%E8%A1%8C%E6%A8%A1%E5%9E%8B%E6%A6%82%E8%BF%B0/

Abstract: 本文介紹CUDA執(zhí)行模型，只比硬件高一層的抽象

Keywords: CUDA SM，SIMT，SIMD，Fermi，Kepler

GPU架構(gòu)是圍繞一個(gè)流式多處理器（SM）的擴(kuò)展陣列搭建的。通過(guò)復(fù)制這種結(jié)構(gòu)來(lái)實(shí)現(xiàn)GPU的硬件并行。

上圖包括關(guān)鍵組件：

• CUDA核心
• 共享內(nèi)存/一級(jí)緩存
• 寄存器文件
• 加載/存儲(chǔ)單元
• 特殊功能單元
• 線程束調(diào)度器

1.1 SM

一個(gè)GPU含有多個(gè)SM；

一個(gè)SM支持?jǐn)?shù)百個(gè)線程并發(fā)執(zhí)行；

當(dāng)一個(gè)核函數(shù)的網(wǎng)格被啟動(dòng)的時(shí)候，多個(gè)block會(huì)被同時(shí)分配給可用的SM上執(zhí)行。

當(dāng)一個(gè)blcok被分配給一個(gè)SM后，他就只能在這個(gè)SM上執(zhí)行了，不可能重新分配到其他SM上了;

多個(gè)線程塊可以被分配到同一個(gè)SM上。

SM上同一個(gè)塊內(nèi)的多個(gè)線程進(jìn)行線程級(jí)別并行；同一線程內(nèi)，指令利用指令級(jí)并行將單個(gè)線程處理成流水線。

1.2 warp(線程束)

CUDA 采用單指令多線程SIMT架構(gòu)管理執(zhí)行線程；

不同設(shè)備有不同的線程束大小，但是到目前為止基本所有設(shè)備都是維持在32；

• 每個(gè)SM上有多個(gè)block；一個(gè)block有多個(gè)線程（數(shù)量不能超過(guò)最大值）；
• 從機(jī)器的角度，在某時(shí)刻T，SM上只執(zhí)行一個(gè)線程束，即這個(gè)線程束中的線程同時(shí)同步的執(zhí)行。
  • 即：warp是基本的執(zhí)行單元。
• warp中的每一個(gè)線程執(zhí)行同一條指令，包括有分支的部分（只是有的線程可能不會(huì)執(zhí)行分支部分）。

1.3 SIMD vs SIMT

雖然兩者都是將相同指令廣播給多個(gè)執(zhí)行單元，同一時(shí)刻所有線程被分配給相同的指令，但是：

• SIMD單指令多數(shù)據(jù)：SIMD規(guī)定所有人必須執(zhí)行，分支必須同時(shí)執(zhí)行相同的指令，必須執(zhí)行沒(méi)有例外。SIMD更像是指令級(jí)別的并行。
• SIMT單指令多線程：規(guī)定有些人可以根據(jù)需要不執(zhí)行，這樣SIMT就保證了線程級(jí)別的并行。（如分支部分有些線程就不執(zhí)行分支部分的操作）

SIMT相比于SIMD，其有以下特征：

• 每個(gè)線程都有自己的指令地址計(jì)數(shù)器；
• 每個(gè)線程都有自己的寄存器狀態(tài)；
• 每個(gè)線程可以有一個(gè)獨(dú)立的執(zhí)行路徑；

而上面這三個(gè)特性在編程模型可用的方式就是給每個(gè)線程一個(gè)唯一的標(biāo)號(hào)（blckIdx,threadIdx），并且這三個(gè)特性保證了各線程之間的獨(dú)立。

1.4 只有warp內(nèi)的線程有相同的進(jìn)度

因?yàn)镾M有限，雖然我們的編程模型層面看所有線程都是并行執(zhí)行的。

但是在微觀上看，所有線程塊也是分批次的在物理層面的機(jī)器上執(zhí)行，線程塊里不同的線程可能進(jìn)度都不一樣，但是同一個(gè)線程束內(nèi)的線程擁有相同的進(jìn)度。

并行就會(huì)引起競(jìng)爭(zhēng)，多線程以未定義的順序訪問(wèn)同一個(gè)數(shù)據(jù)，就導(dǎo)致了不可預(yù)測(cè)的行為，CUDA只提供了一種塊內(nèi)同步的方式，塊之間沒(méi)辦法同步！

同一個(gè)SM上可以有不止一個(gè)常駐的線程束，有些在執(zhí)行，有些在等待，他們之間狀態(tài)的轉(zhuǎn)換是不需要開(kāi)銷(xiāo)的。即：每個(gè)線程束在同一時(shí)間執(zhí)行同一指令，同一個(gè)塊內(nèi)的線程束互相切換是沒(méi)有時(shí)間消耗的。

1.5 Fermi架構(gòu)、Kepler架構(gòu)

這一小節(jié)請(qǐng)閱讀：https://face2ai.com/CUDA-F-3-1-CUDA%E6%89%A7%E8%A1%8C%E6%A8%A1%E5%9E%8B%E6%A6%82%E8%BF%B0/#Fermi-架構(gòu)

1.6 CUDA執(zhí)行的順序

這一節(jié)以Fermi架構(gòu)作為硬件基礎(chǔ)。

每個(gè)SM有兩個(gè)線程束調(diào)度器，兩個(gè)指令調(diào)度單元。

當(dāng)一個(gè)線程塊被指定給一個(gè)SM時(shí)，線程塊內(nèi)的所有線程被分成warp，兩個(gè)線程束調(diào)度器會(huì)選擇其中的兩個(gè)線程束(warp)，然后用指令調(diào)度器存儲(chǔ)兩個(gè)線程束要執(zhí)行的指令。這些線程束在SM上交替執(zhí)行。

總結(jié)

以上是生活随笔為你收集整理的CUDA编程：笔记1的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問(wèn)題。

如果覺(jué)得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。