cuda編程（一）基礎(chǔ)-創(chuàng)新互聯(lián)

• 基于c/c++的編程方法
• 支持異構(gòu)編程的擴(kuò)展方法
• 簡單明了的apis，能夠輕松的管理存儲系統(tǒng)
  cuda支持的編程語言：c/c++/python/fortran/java…

1、CUDA并行計(jì)算基礎(chǔ)

• 異構(gòu)計(jì)算
• CUDA 安裝
• CUDA 程序的編寫
• CUDA 程序編譯
• 利用NVProf查看程序執(zhí)行情況

gpu不是單獨(dú)的在計(jì)算機(jī)中完成任務(wù)，而是通過協(xié)助cpu和整個(gè)系統(tǒng)完成計(jì)算機(jī)任務(wù)，把一部分代碼和更多的計(jì)算任務(wù)放到gpu上處理，邏輯控制、變量處理以及數(shù)據(jù)預(yù)處理等等放在cpu上處理。

創(chuàng)新互聯(lián)公司是一家集網(wǎng)站建設(shè),環(huán)翠企業(yè)網(wǎng)站建設(shè),環(huán)翠品牌網(wǎng)站建設(shè),網(wǎng)站定制,環(huán)翠網(wǎng)站建設(shè)報(bào)價(jià),網(wǎng)絡(luò)營銷,網(wǎng)絡(luò)優(yōu)化,環(huán)翠網(wǎng)站推廣為一體的創(chuàng)新建站企業(yè)，幫助傳統(tǒng)企業(yè)提升企業(yè)形象加強(qiáng)企業(yè)競爭力。可充分滿足這一群體相比中小企業(yè)更為豐富、高端、多元的互聯(lián)網(wǎng)需求。同時(shí)我們時(shí)刻保持專業(yè)、時(shí)尚、前沿，時(shí)刻以成就客戶成長自我，堅(jiān)持不斷學(xué)習(xí)、思考、沉淀、凈化自己，讓我們?yōu)楦嗟钠髽I(yè)打造出實(shí)用型網(wǎng)站。

host 指的是cpu和內(nèi)存
device 指的是gpu和顯存
nvidia-smi 查看當(dāng)前gpu的運(yùn)行狀態(tài)

2、第一個(gè)cuda程序

系統(tǒng)中安裝了cuda但是執(zhí)行nvcc找不到命令。
添加環(huán)境變量。
vim ~/.bashrc
加入環(huán)境變量

export PATH="/usr/local/cuda-10.2/bin:$PATH"
export LD_LIBRARY_PATH="/usr/local/cuda-10.2/lib64:$LD_LIBRARY_PATH"

source ~/.bashrc

再次執(zhí)行nvcc -V結(jié)果如下。

2.1 helloword

nvcc也可以支持純c的代碼，所以先寫一個(gè)helloword的代碼進(jìn)行，使用nvcc進(jìn)行編譯！
cuda程序的編譯器驅(qū)動nvcc支持編譯純粹的c++代碼，一個(gè)標(biāo)準(zhǔn)的CUDA程序中既有C++代碼也有不屬于C++的cuda代碼。cuda程序的編譯器驅(qū)動nvcc在編譯一個(gè)cuda程序時(shí)，會將純粹的c++代碼交給c++的編譯器，他自己負(fù)責(zé)編譯剩下的部分（cuda）代碼。
創(chuàng)建hell.cu文件，cuda的代碼需要以cu為后綴結(jié)尾。

#includeint main()
{printf("helloword\n");
        return 0;
}
~      

nvcc hell.cu
./a.out
運(yùn)行結(jié)果如下。

2.2 核函數(shù)

cuda 中的核函數(shù)與c++中的函數(shù)是類似的，cuda的核函數(shù)必須被限定詞__global__修飾,核函數(shù)的返回類型必須是空類型，即void.

#include__global__ void hello_from_gpu()
{printf("hello word from the gpu!\n");
}

int main()
{hello_from_gpu<<<1,1>>>();
   cudaDeviceSynchronize();
   printf("helloword\n");
   return 0;
}
~     

運(yùn)行結(jié)果如下。

在核函數(shù)的調(diào)用格式上與普通C++的調(diào)用不同，調(diào)用核函數(shù)的函數(shù)名和（）之間有一對三括號，里面有逗號隔開的兩個(gè)數(shù)字。因?yàn)橐粋€(gè)GPU中有很多計(jì)算核心，可以支持很多個(gè)線程。主機(jī)在調(diào)用一個(gè)核函數(shù)時(shí)，必須指明需要在設(shè)備中指派多少個(gè)線程，否則設(shè)備不知道怎么工作。三括號里面的數(shù)就是用來指明核函數(shù)中的線程數(shù)以及排列情況的。核函數(shù)中的線程常組織為若干線程塊（thread block）。
三括號中的第一個(gè)數(shù)時(shí)線程塊的個(gè)數(shù)，第二個(gè)數(shù)可以看作每個(gè)線程中的線程數(shù)。一個(gè)核函數(shù)的全部線程塊構(gòu)成一個(gè)網(wǎng)格，而線程塊的個(gè)數(shù)記為網(wǎng)格大小，每個(gè)線程塊中含有同樣數(shù)目的線程，該數(shù)目稱為線程塊大小。所以核函數(shù)中的總的線程就等與網(wǎng)格大小乘以線程塊大小，即<<<網(wǎng)格大小，線程塊大小 >>>
核函數(shù)中的printf函數(shù)的使用方法和C++庫中的printf函數(shù)的使用方法基本上是一樣的，而在核函數(shù)中使用printf函數(shù)時(shí)也需要包含頭文件,核函數(shù)中不支持C++的iostream。
cudaDeviceSynchronize();這條語句調(diào)用了CUDA運(yùn)行時(shí)的API函數(shù)，去掉這個(gè)函數(shù)就打印不出字符了。因?yàn)閏uda調(diào)用輸出函數(shù)時(shí)，輸出流是先放在緩存區(qū)的，而這個(gè)緩存區(qū)不會核會自動刷新，只有程序遇到某種同步操作時(shí)緩存區(qū)才會刷新。這個(gè)函數(shù)的作用就是同步主機(jī)與設(shè)備，所以能夠促進(jìn)緩存區(qū)刷新。

3、cuda中的線程組織 3.1 使用多個(gè)線程的核函數(shù)

核函數(shù)中允許指派很多線程，一個(gè)GPU往往有幾千個(gè)計(jì)算核心，而總的線程數(shù)必須至少等與計(jì)算核心數(shù)時(shí)才有可能充分利用GPU的全部計(jì)算資源。實(shí)際上，總的線程數(shù)大于計(jì)算核心數(shù)時(shí)才能更充分地利用GPU中的計(jì)算資源，因?yàn)檫@會讓計(jì)算和內(nèi)存訪問之間及不同的計(jì)算之間合理地重疊，從而減小計(jì)算核心空閑的時(shí)間。
使用網(wǎng)格數(shù)為2，線程塊大小為4的計(jì)算核心，所以總的線程數(shù)就是2x4=8，所以核函數(shù)的調(diào)用將指派8個(gè)線程完成。
核函數(shù)中的代碼的執(zhí)行方式是“單指令-多線程”，即每一個(gè)線程都執(zhí)行同一指令的內(nèi)容。

#include__global__ void hello_from_gpu()
{printf("hello word from the gpu!\n");
}

int main()
{hello_from_gpu<<<2,4>>>();
   cudaDeviceSynchronize();
   printf("helloword\n");
   return 0;
}

運(yùn)行結(jié)果如下。

3.2 線程索引的使用

一個(gè)核函數(shù)可以指派多個(gè)線程，而這些線程的組織結(jié)構(gòu)是由執(zhí)行配置（<<<網(wǎng)格大小，線程塊大小 >>>）來決定的，這是的網(wǎng)格大小和線程塊大小一般來說是一個(gè)結(jié)構(gòu)體類型的變量，也可以是一個(gè)普通的整形變量。

一個(gè)核函數(shù)允許指派的線程數(shù)是巨大的，能夠滿足幾乎所有應(yīng)用程序的要求。但是一個(gè)核函數(shù)中雖然可以指派如此巨大數(shù)目的線程數(shù)，但在執(zhí)行時(shí)能夠同時(shí)活躍（不活躍的線程處于等待狀態(tài)）的線程數(shù)是由硬件（主要是CUDA核心數(shù)）和軟件（核函數(shù)的函數(shù)體）決定的。
每個(gè)線程在核函數(shù)中都有一個(gè)唯一的身份標(biāo)識。由于我們在三括號中使用了兩個(gè)參數(shù)制定了線程的數(shù)目，所以線程的身份可以由兩個(gè)參數(shù)確定。在程序內(nèi)部，程序是知道執(zhí)行配置參數(shù)grid_size和block_size的值的，這兩個(gè)值分別保存在內(nèi)建變量（built-in vari-
able）中。
gridDim.x ：該變量的數(shù)值等與執(zhí)行配置中變量grid_size的數(shù)值。
blockDim.x: 該變量的數(shù)值等與執(zhí)行配置中變量block_size的數(shù)值。
在核函數(shù)中預(yù)定義了如下標(biāo)識線程的內(nèi)建變量：
blockIdx.x :該變量指定一個(gè)線程在一個(gè)網(wǎng)格中的線程塊指標(biāo)。其取值范圍是從0到gridDim.x-1
threadIdx.x：該變量指定一個(gè)線程在一個(gè)線程塊中的線程指標(biāo)，其取值范圍是從0到blockDim.x-1
代碼如下。

#include__global__ void hello_from_gpu()
{const int bid = blockIdx.x;
   const int tid = threadIdx.x;
   printf("hello word from block %d and thread %d\n",bid,tid);
}
int main()
{hello_from_gpu<<<2,4>>>();
   cudaDeviceSynchronize(); 
   printf("helloword\n");
   return 0;
}

有時(shí)候線程塊的順序會發(fā)生改變，有時(shí)候是第1個(gè)先執(zhí)行有時(shí)候是第0個(gè)先執(zhí)行，這說明了cuda程序執(zhí)行時(shí)每個(gè)線程塊的計(jì)算都是相互獨(dú)立的，不管完成計(jì)算的次序如何，每個(gè)線程塊中間的每個(gè)線程都進(jìn)行一次計(jì)算。

3.3 cuda多維網(wǎng)格

上述四個(gè)內(nèi)建變量都使用了C++中的結(jié)構(gòu)體或者類的成員變量的語法，其中blockIdx和threadIdx是類型為uint3的變量，該類型是一個(gè)結(jié)構(gòu)體，具有x,y,z三個(gè)成員變量。所以blockIdx只是三個(gè)成員中的一個(gè)，threadIdx也有xyz三個(gè)成員變量。結(jié)構(gòu)體uint3在頭文件vector_types.h中定義有。
同樣的gridDim和blockDim是dim3類型的變量。也有xyz三個(gè)成員變量。
在前面三括號內(nèi)的網(wǎng)格大小和線程塊大小都是通過一維表示，可以通過dim3定義多維網(wǎng)格和線程塊，通過C++的構(gòu)造函數(shù)的方法實(shí)現(xiàn)。di3 grid_size(Gx.Gy,Gz);
如果第三個(gè)維度是1，可以省去不寫。
多維的網(wǎng)格和線程塊本質(zhì)上還是一維的，就像多維數(shù)組本質(zhì)上也是一維數(shù)組一樣。一個(gè)多維線程指標(biāo)threadIdx.x、threadIdx.y、threadIdx.z對應(yīng)的一維指標(biāo)為。

int tid = threadIdx.z * blockDim.x * blockDim.y +threadIdx.y * blockDim.x + threadIdx.x;
也就是說，x維度是最內(nèi)層的變化最快的，而z維度是最外層的變化最滿的。

代碼如下。

#include__global__ void hello_from_gpu()
{const int bid = blockIdx.x;
   const int tid = threadIdx.x;
   const int yid = threadIdx.y;
   printf("hello word from block %d and thread (%d,%d)\n",bid,tid,yid);
}
int main()
{const dim3 block_size(2,4);
   hello_from_gpu<<<1,block_size>>>();
   cudaDeviceSynchronize();
   printf("helloword\n");
   return 0;
}

因?yàn)榫€程塊的大小是2*4，所以在核函數(shù)中，blockDim.x的值為2，blokcDim.y值是4，threadIdx.x的取值是0到1，threadIdx.y的取值是0到3。
從結(jié)果可以看到。x維度是變量最快的是最內(nèi)層的，是因?yàn)榻Y(jié)果中，前兩行的x層發(fā)生了0-1的轉(zhuǎn)變，但是y層依然表示0。

3.3 網(wǎng)格與線程塊大小的限制

cuda中對能夠定義的網(wǎng)格大小和線程塊大小做了限制，一個(gè)線程塊最多只能有1024個(gè)線程。

4 cuda中的頭文件

cuda有自己的頭文件，但是在使用nvcc編譯器驅(qū)動.cu文件時(shí)，將自動包含必要的cuda頭文件，如個(gè)和。因?yàn)?cuda.h>包含了,所以在使用nvcc編譯的cuda程序也不需要包含stdlib頭文件。

5 nvcc編譯cuda程序

cuda的編譯器驅(qū)動nvcc會先將全部源代碼分離為主機(jī)代碼和設(shè)備代碼。設(shè)備代碼完全支持C++語法，但設(shè)備代碼只部分地支持C++。nvcc先將設(shè)備代碼編譯為PTX偽匯編代碼，再將PTX代碼編譯為二進(jìn)制的文件cubin目標(biāo)代碼。在將源代碼編譯為PTX代碼時(shí)，需要用選項(xiàng)-arch=compute_XY指定一個(gè)虛擬架構(gòu)的計(jì)算能力，用以確定代碼中共能夠使用的cuda功能。在將PTX代碼編譯為cubin代碼時(shí)，需要用選項(xiàng)-code=sm_ZW指定一個(gè)真實(shí)架構(gòu)的計(jì)算能力，用以確定可執(zhí)行文件能夠使用的GPU。真實(shí)架構(gòu)的計(jì)算能力必須等與或者大于虛擬架構(gòu)的計(jì)算能力。

你是否還在尋找穩(wěn)定的海外服務(wù)器提供商？創(chuàng)新互聯(lián)www.cdcxhl.cn海外機(jī)房具備T級流量清洗系統(tǒng)配攻擊溯源，準(zhǔn)確流量調(diào)度確保服務(wù)器高可用性，企業(yè)級服務(wù)器適合批量采購，新人活動首月15元起，快前往官網(wǎng)查看詳情吧

分享文章：cuda編程（一）基礎(chǔ)-創(chuàng)新互聯(lián)
鏈接URL：http://muchs.cn/article38/dcpopp.html

成都網(wǎng)站建設(shè)公司_創(chuàng)新互聯(lián)，為您提供微信公眾號、軟件開發(fā)、網(wǎng)站內(nèi)鏈、企業(yè)網(wǎng)站制作、搜索引擎優(yōu)化、電子商務(wù)

廣告

聲明：本網(wǎng)站發(fā)布的內(nèi)容（圖片、視頻和文字）以用戶投稿、用戶轉(zhuǎn)載內(nèi)容為主，如果涉及侵權(quán)請盡快告知，我們將會在第一時(shí)間刪除。文章觀點(diǎn)不代表本網(wǎng)站立場，如需處理請聯(lián)系客服。電話：028-86922220；郵箱：631063699@qq.com。內(nèi)容未經(jīng)允許不得轉(zhuǎn)載，或轉(zhuǎn)載時(shí)需注明來源： 創(chuàng)新互聯(lián)