《基於CUDA的並行程式設計》學習筆記 三 中

2021-08-11 05:46:02 字數 3527 閱讀 6326

cuda架構第一次引入了主機(host)端與裝置(device)端的概念。如下圖所示,乙個完整的cuda程式由主機**和裝置**兩部分組成。主機**在主機端cpu上序列執行,是普通的c**;裝置**部分在裝置端gpu上並行執行,稱為核心(kernel)。kernel函式不是乙個完整的程式,而是任務中能被分解為並行執行的步驟的集合。cpu執行的序列程式負責kernel啟動之前進行資料準備和裝置初始化的工作,以及在kernel之間進行一些序列計算。gpu執行的並行部分是在被稱為grid和block的兩個層次並行中完成的,即每個kernel函式存在兩個層次的並行:網路(grid)中的執行緒塊(block)間並行金額執行緒塊中的執行緒(thread)間並行。

3.3.1 執行結構

cuda程式在執行過程中,主機**與裝置**交替執行,程式從主機端的序列**開始執行,當執行至裝置**時,呼叫核心函式,並切換到裝置端,啟動多個gpu執行緒(thread)並行執行裝置**。每乙個執行緒塊(block)均包含多個執行緒,執行相同的指令,達到執行緒塊內的並行,同時在每個執行緒網格(grid)中,實現不同執行緒塊之間的並行。裝置完成計算後,返回主機執行緒,主機繼續執行序列操作。重複以上過程,直到程度執行完畢。

3.3.2 核心函式

先看乙個簡單的c檔案,實現長度均為n的陣列a和資料b的相加,結果儲存在陣列c中。

void vectoradd(float *a, float *b, float *c, int n)

}int main()

在主函式main中呼叫函式vectoradd,使用while迴圈進行a[x]與b[x]的相加並將結果存入c[x]中。

// 裝置端**

__global__ void vectoradd(float

*a, float

*b, float

*c)// 主機端**

int main()

在主函式main中呼叫核心函式vectoradd,使得每乙個執行緒平行計算每乙個a[x]與b[x]的加法並將結果存入c[x]中。

在cuda程式中,程式設計師可以自定義稱為核心的c語言函式,如同上面的vectoradd函式,用__global__宣告說明符來指定核心程式。和普通的c語言程式一樣,cuda程式的執行入口仍是main函式。乙個完整的cuda程式包括了在cpu端執行的序列**和在gpu端執行的並行**,其流程如下圖所示,主要包括以下幾個步驟。

3.3.3 執行緒層次

在核心函式中,通過執行緒的索引來訪問執行緒id。在一維陣列中,可以用一維執行緒塊中的threadidx直接指向相應id的執行緒,但是在二維資料、三維資料中卻不相同。對於大小為(dx,dy)的二維塊,索引為(x,y)的執行緒的id是(x+ydx);對於大小為(dx,dy,dz)的三維塊,索引為(x,y,z)的執行緒的id是(x+ydx+zdxdy)。

有了針對不同維度的陣列的索引方式,可以定義出二維、三維的執行緒塊,去應對不同情況的資料並行方式。

上乙個例子中,使用執行緒的並行操作實現了高效的陣列相加,現在可以以執行緒塊的並行方式,進一步分析矩陣加法的cuda程式設計。

__global__ void matrixadd(float a[n][n], float b[n][n], float c[n][n])

int main()

cuda程式使用類似kernelfunc<<>>(d_a, d_b, d_c);的語句來啟動kernel函式,其中,<<<>>>運算子中的n和m是主機端告訴裝置執行時如何啟動kernel函式的引數,n表示乙個grid中有多少個並行的block,m表示乙個block中多少個並行thread;(d_a, d_b, d_c)則為kernel函式的函式引數,和普通c函式一樣。在上面**中,執行核心的每個執行緒都會被分配乙個獨特的執行緒id,可通過threadidx變數在核心中訪問此id。

3.3.4 儲存器結構與執行緒對映機制

cuda程式執行過程中,如下圖所示,每乙個執行緒作為執行緒塊的一部分,都單獨擁有乙個私有的本地儲存器,同時也可訪問多個其他儲存器空間。每乙個執行緒塊作為網格的一部分,都有乙個共享儲存器,可供執行緒塊內的每乙個執行緒訪問,並且與塊具有相同的生命週期。所有執行緒都可訪問相同的全域性儲存器。另外,所有執行緒也可以訪問固定儲存器和紋理儲存器,這兩個儲存器均為唯讀。與本地儲存器和共享儲存器不同的是,全域性儲存器、固定儲存器和紋理儲存器有持久的生命週期。

3.3.5 通訊機制

目前,gpu與計算機晶元組的各個部件,如cpu、記憶體等的通訊介面目前主要是pci-e,它應日益膨脹的通訊資料量要求而產生,代替了傳統的agp介面。當gpu與晶元組通訊的資料超過了這個數,就會出現等待響應的情況,影響資料傳輸效率。因此,為了提高程式的執行效率,在cuda程式中應盡量少地進行主機端與裝置端之間的資料傳輸。

3.3.6 cuda的軟體體系

cuda的軟體體系由以下三部分構成:cuda庫函式(library)、cuda執行時api(runtime api)和cuda驅動api(drive api)構成。cuda軟體體系機構如下圖所示。

cuda應用程式是用cuda c語言編寫的,cuda提供了nvcc編譯器對其進行編譯。對cuda c語言進行編譯得到的只是gpu端的**。而gpu資源的管理、gpu上分配視訊記憶體並啟動kernel函式,則由cuda執行時候api或者cuda驅動api負責。這兩者不能混合使用,乙個程式中只能使用其中乙個。

3.4.1 nvcc編譯流程

kernel函式可用cuda c語言編寫,也可以用ptx編寫。ptx就是cuda指令集架構,效率高於像c語言一樣的高階語言。但無論使用ptx還是高階語言,kernel函式都必須通過nvcc編譯成可執行的二進位制檔案後才可以在裝置上執行。

nvcc是cuda程式編譯器驅動,簡化了c語言或ptx的編譯流程:它提供了簡單的命令列選項,呼叫一系列不同的編譯工具集來執行它們。nvcc可編譯同時包含主機**(在主機上執行的**)和裝置**(在裝置上執行的**)的原始檔。nvcc的基本流程如下圖所示,包括了以下4個步驟:

3.4.2 相容性分析

CUDA程式設計第一章 基於CUDA的異構平行計算

本部分主要講解並行的基礎知識,在前頭的mpi openmp並行程式設計基礎基本掌握過了,直接pass 這裡的記憶體頻寬應該指的是視訊記憶體頻寬 有兩個考量點 並行性資料規模 如果乙個問題有較小的資料規模 複雜的控制邏輯和 或很少的並行性,那麼最好選擇cpu處理該問題,因為它有處理複雜邏輯和指令級並行...

CUDA並行程式設計學習心得1

正在學習cuda中,寫一些自己學習中的心得,如有錯誤,歡迎指正 對於普通變數 如int型,float型等 不必在device端事先分配空間。在global函式宣告時,不要宣告為指標,直接宣告為int或float。global函式呼叫時也是。如果要在device端以指標形式使用這個變數,如下 floa...

CUDA學習筆記(2) 執行緒並行和塊並行

有些顯示卡支援cuda有些不支援,那麼如何確定主機的顯示卡裝置是否支援cuda呢。可以使用下面的函式獲取顯示卡的相關資訊。下面是關於這幾個函式的簡單使用 cudaerror t cudastatus int number 0 cudadeviceprop prop cudastatus cudage...