CUDA 編程模型

概述

一、 主機與設備

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

二、運行在 GPU 上的 CUDA 并行計算函數稱為 kernel( 內核函數 ) 。一個 kernel 函數并不是一個完整的程序，而是整個 CUDA 程序中的一個可以被并行執行的步驟。

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

Kernel 函數的定義與調用

一、 運行在 GPU 上的程序稱為 kernel( 內核函數 ) 。

內核函數必須通過 _global_ 函數類型限定符定義，并且只能在主機端代碼中調用。在調用時，必須聲明內核函數的執行參數。

注意： __global__ 下劃線在Visual studio中的寫法。

使用一種新 <<<…>>> 執行配置語法指定執行某一指定內核調用的線程數。必須先為 Kernel 中用到的數組或變量分配好足夠的空間，再調用 kernel 函數

二、 在設備端運行的線程之間是并行執行的，每個線程有自己的 blockID 和 threadID 用于與其他線程相區分。 BlockID 和 threadID 只能在 kernel 中通過內建變量訪問。

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

線程結構

Kernel 是以 block 為單位執行的， CUDA 引入 grid 來表示一系列可以被并行執行的 block 的集合。各 block 是并行執行的， block 之間無法通信，也沒有執行順序。

?

block 內通信原理

在同一個 block 中的線程通過共享存儲器 (shared memory) 交換數據，并通過柵欄同步保證線程間能夠正確地共享數據。具體來說，可以在 kernel 函數中需要同步的位置調用 _syncthreads() 函數。

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

計算單元

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

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

目前，一個內核函數只有一個 grid ，但在支持 DirectX 11 的硬件中，這一限制將會解除。

真正的執行單元

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

執行模型

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