CUDA 編程模型

概述

一、 主機(jī)與設(shè)備

CUDA 編程模型將 CPU 作為主機(jī) (Host) ， GPU 作為協(xié)處理器 (co-processor) 或者設(shè)備 (Device). 在一個(gè)系統(tǒng)中可以存在一個(gè)主機(jī)和多個(gè)設(shè)備。 CPU 主要負(fù)責(zé)進(jìn)行邏輯性強(qiáng)的事物處理和串行計(jì)算， GPU 則專注于執(zhí)行高度線程化的并行處理任務(wù)。 CPU 、 GPU 各自擁有相互獨(dú)立的存儲(chǔ)器地址空間：主機(jī)端的內(nèi)存和設(shè)備端的顯存。 CUDA 對(duì)內(nèi)存的操作與一般的 C 程序基本相同，但增加了一種新的 pinned memory ；

二、運(yùn)行在 GPU 上的 CUDA 并行計(jì)算函數(shù)稱為 kernel( 內(nèi)核函數(shù) ) 。一個(gè) kernel 函數(shù)并不是一個(gè)完整的程序，而是整個(gè) CUDA 程序中的一個(gè)可以被并行執(zhí)行的步驟。

三、 CPU 串行代碼完成的工作包括在 kernel 啟動(dòng)前進(jìn)行數(shù)據(jù)準(zhǔn)備和設(shè)備初始化的工作，以及在 kernel 之間進(jìn)行一些串行計(jì)算。理想情況是， CPU 串行代碼的作用應(yīng)該只是清理上一個(gè)內(nèi)核函數(shù)，并啟動(dòng)下一個(gè)內(nèi)核函數(shù)。這樣，就可以在設(shè)備上完成盡可能多的工作，減少主機(jī)與設(shè)備之間的數(shù)據(jù)傳輸。

Kernel 函數(shù)的定義與調(diào)用

一、 運(yùn)行在 GPU 上的程序稱為 kernel( 內(nèi)核函數(shù) ) 。

內(nèi)核函數(shù)必須通過(guò) _global_ 函數(shù)類型限定符定義，并且只能在主機(jī)端代碼中調(diào)用。在調(diào)用時(shí)，必須聲明內(nèi)核函數(shù)的執(zhí)行參數(shù)。

注意： __global__ 下劃線在Visual studio中的寫(xiě)法。

使用一種新 <<<…>>> 執(zhí)行配置語(yǔ)法指定執(zhí)行某一指定內(nèi)核調(diào)用的線程數(shù)。必須先為 Kernel 中用到的數(shù)組或變量分配好足夠的空間，再調(diào)用 kernel 函數(shù)

二、 在設(shè)備端運(yùn)行的線程之間是并行執(zhí)行的，每個(gè)線程有自己的 blockID 和 threadID 用于與其他線程相區(qū)分。 BlockID 和 threadID 只能在 kernel 中通過(guò)內(nèi)建變量訪問(wèn)。

三、 內(nèi)建變量不需由程序員定義，是由設(shè)備中的專用寄存器提供的。所以，內(nèi)建變量是只讀的，并且只能在 GPU 端得 kernel 函數(shù)中使用。

線程結(jié)構(gòu)

Kernel 是以 block 為單位執(zhí)行的， CUDA 引入 grid 來(lái)表示一系列可以被并行執(zhí)行的 block 的集合。各 block 是并行執(zhí)行的， block 之間無(wú)法通信，也沒(méi)有執(zhí)行順序。

?

block 內(nèi)通信原理

在同一個(gè) block 中的線程通過(guò)共享存儲(chǔ)器 (shared memory) 交換數(shù)據(jù)，并通過(guò)柵欄同步保證線程間能夠正確地共享數(shù)據(jù)。具體來(lái)說(shuō)，可以在 kernel 函數(shù)中需要同步的位置調(diào)用 _syncthreads() 函數(shù)。

為了保證線程塊中的各個(gè)線程能夠有效協(xié)作，訪問(wèn)共享存儲(chǔ)器的延遲必須很小。所以在 GPU 中，共享存儲(chǔ)器與執(zhí)行單元的物理距離必須很小，處于同一個(gè)處理核心中。而為了在硬件上用很小的代價(jià)就能實(shí)現(xiàn) _syncthreads() 函數(shù)，一個(gè) block 中所有線程的數(shù)據(jù)都必須交由同一處理核心進(jìn)行處理。所以，這導(dǎo)致每個(gè)線程塊中的線程數(shù)量、共享存儲(chǔ)器大小、寄存器數(shù)量都要受到處理核心硬件資源的限制。目前，每個(gè) block 里最多只能有 512 個(gè)線程。

計(jì)算單元

在 GPU 內(nèi)部， SM 代表流多處理器，即計(jì)算核心。每個(gè) SM 中又包含 8 個(gè)標(biāo)量流處理器 SP 以及少量的其他計(jì)算單元。實(shí)際上， SP 只是執(zhí)行單元，并不是完整的處理核心。處理核心必須包含取指、解碼、分發(fā)邏輯和執(zhí)行單元。隸屬同一 SM 的 8 個(gè) SP 共用同一套取指和發(fā)射單元，也共用一塊共享存儲(chǔ)器。

一個(gè) block 必須被分配到一個(gè) SM 中，但是一個(gè) SM 中同一時(shí)刻可以有多個(gè)活動(dòng)線程塊等待執(zhí)行。這可以更好地利用執(zhí)行單元的資源，當(dāng)一個(gè) block 進(jìn)行同步或者訪問(wèn)顯存等高延遲操作時(shí)，另一個(gè) block 就可以占用 GPU 執(zhí)行資源。

目前，一個(gè)內(nèi)核函數(shù)只有一個(gè) grid ，但在支持 DirectX 11 的硬件中，這一限制將會(huì)解除。

真正的執(zhí)行單元

在實(shí)際運(yùn)行中， block 會(huì)被分割為更小的線程束 (warp) ， warp 的大小由硬件的計(jì)算能力版本決定。在采用 Tesla 架構(gòu)的 GPU 中，一個(gè)線程束由連續(xù)的 32 個(gè)線程組成。 Warp 中的線程只與線程 thread? ID 有關(guān)。在每發(fā)射一條 warp 指令， SM 中的 8 個(gè) SP 將執(zhí)行這條指令 4 遍。

執(zhí)行模型

CUDA 采用了 SIMT( 單指令多線程 ) 執(zhí)行模型。在 SIMT 模型中，如果需要控制單個(gè)線程的行為，這會(huì)大大降低效率。