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

1、CUDA并行計算基礎

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

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

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

host 指的是cpu和內存
device 指的是gpu和顯存
nvidia-smi 查看當前gpu的運行狀態(tài)

2、第一個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結果如下。

2.1 helloword

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

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

nvcc hell.cu
./a.out
運行結果如下。

2.2 核函數

cuda 中的核函數與c++中的函數是類似的，cuda的核函數必須被限定詞__global__修飾,核函數的返回類型必須是空類型，即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;
}
~     

運行結果如下。

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

3、cuda中的線程組織 3.1 使用多個線程的核函數

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

#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;
}

運行結果如下。

3.2 線程索引的使用

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

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

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

3.3 cuda多維網格

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

int tid = threadIdx.z * blockDim.x * blockDim.y +threadIdx.y * blockDim.x + threadIdx.x;
也就是說，x維度是最內層的變化最快的，而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;
}

因為線程塊的大小是2*4，所以在核函數中，blockDim.x的值為2，blokcDim.y值是4，threadIdx.x的取值是0到1，threadIdx.y的取值是0到3。
從結果可以看到。x維度是變量最快的是最內層的，是因為結果中，前兩行的x層發(fā)生了0-1的轉變，但是y層依然表示0。

3.3 網格與線程塊大小的限制

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

4 cuda中的頭文件

cuda有自己的頭文件，但是在使用nvcc編譯器驅動.cu文件時，將自動包含必要的cuda頭文件，如個和。因為包含了,所以在使用nvcc編譯的cuda程序也不需要包含stdlib頭文件。

5 nvcc編譯cuda程序

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

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

本文題目：cuda編程（一）基礎-創(chuàng)新互聯
鏈接URL：http://weahome.cn/article/dcpopp.html