?CUDA編程技術(shù)匯總

本文主要羅列了CUDA編程的基本概念、基本原理、基礎(chǔ)框架等相關(guān)內(nèi)容，同時結(jié)合筆者工作實踐，總結(jié)了若干CUDA程序的優(yōu)化技巧。比較適合CUDA初學(xué)者。對于想要系統(tǒng)地、深入地研究CUDA的朋友，建議參閱<CUDA C++ Programming Guide>、<CUDA C++ Best Practices Guide>。

限于筆者認知水平與研究深度，難免有不當指出，歡迎批評指正！

注1：文章內(nèi)容會不定期更新。

注2：文中引用了<CUDA C++ Programming Guide>的圖片，但未逐一標注其出處。

一、高性能計算

高性能計算(High performance computing， 縮寫HPC) 通常指使用很多處理器或者某一集群中組織的幾臺計算機（作為單個計 算資源操作）的計算系統(tǒng)和環(huán)境來完成計算任務(wù)。

目前，各大超算中心以及國內(nèi)那幾個大廠所搭建的HPC集群基本上都是基于Beowulf架構(gòu)及其派生架構(gòu)。雖然國內(nèi)也有不少公司靠此“發(fā)家致富”，但是具有軟硬件核心技術(shù)的公司其實比較少。造成這種現(xiàn)象的一個原因，就是很多公司實施“暴力式基礎(chǔ)建設(shè)”，但在關(guān)鍵的、核心的軟硬件研發(fā)投入、產(chǎn)業(yè)化推廣等方面，還存在很多問題。

二、軟件層面

2.1 CUDA

CUDA是一個并行計算開發(fā)環(huán)境，主要針對基于NVIDIA?GPU硬件。

除了CUDA，還有OpenCL、OpenACC、DirectCompute等類似軟件開發(fā)平臺。

2.2 編程模型

將系統(tǒng)分成主機端與設(shè)備端；設(shè)備端劃分為粗細兩層的并行結(jié)構(gòu)；kernel函數(shù)由主機端調(diào)用，但在設(shè)備端運行。

核函數(shù)：由__global__修飾;返回值類型為void；主機端調(diào)用需要"<<<gridDim,blockDim>>>"指定執(zhí)行配置。

2.3 線程模型

核函數(shù)的全部線程塊構(gòu)成一個網(wǎng)格(Grid)，線程網(wǎng)格由多個線程塊（Block）組成；每個線程塊包含多個線程。在軟件層面，kernel以線程塊為單位并發(fā)執(zhí)行。

CUDA 兩級線程模型主要是為了使同一CUDA程序可以在不同代次的GPU上自動伸縮運行。

同時，這種線程模型也提供了粗細兩種粒度的并行模式（至少軟件層面上，可以這樣理解），支持線程塊內(nèi)線程同步與數(shù)據(jù)通信。

kernel函數(shù)中內(nèi)建變量：a) uint3類型 blockIdx、threadIdx; b）dim3類型gridDim、blockDim

線程配置限制因素：a) gridDim.x、gridDim.y、gridDim.z大小限制; b）blockDim.x、blockDim.y、blockDim.z大小限制; c) blockDim.x*blockDim.y*blockDim.z乘積大小限制；d)線程塊內(nèi)線程總數(shù)通常要設(shè)置為32的倍數(shù)，可取64、128、256; e) 線程網(wǎng)格的維度通常由問題規(guī)模決定。

2.4 內(nèi)存模型

CUDA為了降低訪存延遲，進一步提高計算效率，采用了一種由多種類型內(nèi)存構(gòu)成的分級內(nèi)存模型。這些內(nèi)存主要包括寄存器、共享內(nèi)存、L1緩存、L2緩存、全局內(nèi)存、常量內(nèi)存、紋理內(nèi)存、表面內(nèi)存等。

每一個線程擁有自己的私有存儲器寄存器和局部存儲器；每一個線程塊擁有一塊共享存儲器（shared_memory）；最后，grid中所有的線程都可以訪問同一塊全局存儲器（global memory）。

除此之外，還有兩種可以被所有線程訪問的只讀存儲器：常量存儲器（constant memory）和紋理存儲器（Texture memory），他們分別為不同的應(yīng)用進行了優(yōu)化。

全局存儲器、常數(shù)存儲器和紋理存儲器中的值在一個內(nèi)核函數(shù)執(zhí)行完成后將被繼續(xù)保持，可以被同一程序中的其他內(nèi)核函數(shù)調(diào)用。

存儲器	位置	擁有緩存	訪問權(quán)限	變量生命周期	帶寬	延遲 （cycles）
register	GPU片內(nèi)	N/A	device可讀/寫	與thread相同	~8TB/s	1
local_memory	板載顯存	無	device可讀/寫	與thread相同
shared_memory	GPU片內(nèi)	N/A	device可讀/寫	與block相同	~1.5TB/s	1 - 32
constant_memory	板載顯存	有	device可讀,host可讀/寫	可在程序中保持	~200GB/s	~400-600
texture_memory	板載顯存	有	device可讀,host可讀/寫	可在程序中保持	~200GB/s	~400-600
global_memory	板載顯存	無	device可讀/寫,host可讀/寫	可在程序中保持	~200GB/s	~400-600
host_memory	host內(nèi)存	無	host可讀/寫	可在程序中保持
pinned_memory	host內(nèi)存	無	host可讀/寫	可在程序中保持

在計算能力2.x之前的設(shè)備，全局內(nèi)存的訪問會在L1\L2 cache上緩存；在計算能力3.x以上的設(shè)備，全局內(nèi)存的訪問只在L2 cache上緩存。對于L1 cache，每次按照128字節(jié)進行緩存；對于L2 cache，每次按照32字節(jié)進行緩存。

3.1 GPU硬件架構(gòu)
三、硬件層面

GPU是顯卡上的核心組件，顯卡一般插在主板PCI-E插槽上，經(jīng)北橋與CPU相連。

GPU可以簡單看成由多個流多處理器(Streaming Multiprocessor, SM)構(gòu)成的陣列，每個SM具有多個流處理器(Streaming Processor,SP)。

3.2 Warp

在硬件層面，線程塊內(nèi)的線程按照線程號劃分成不同的warp，kernel以warp為單位執(zhí)行。線程束(Warp)是同一線程塊中相鄰的warpSize個線程。對于目前所有的GPU架構(gòu)warpSize是32。

四、優(yōu)化策略

4.1 線程塊的數(shù)量應(yīng)該是SM數(shù)量的整數(shù)倍，總線程數(shù)要大于GPU計算核心數(shù)，這樣可以有效地重疊計算與訪存，使GPU 計算核心保持忙碌，從而最大程度的發(fā)揮GPU的計算性能。

4.2 減少CPU<->GPU的數(shù)據(jù)傳輸；可以的話，考慮合并kernel函數(shù)。

4.3 訪問一次全局內(nèi)存，將耗費400~600個cycle，成本是非常高的，所以需要考慮提升全局內(nèi)存訪問效率的方法。

4.3.1 當硬件檢測到同一個warp中的線程訪問全局存儲器中連續(xù)的存儲單元時，便會將這些訪存合并成一個訪存。合并訪問可以提高DRAM的帶寬利用率，使DRAM在傳輸數(shù)據(jù)時的速度接近全局存儲器的峰值。

參考書籍

樊哲勇. CUDA編程：基礎(chǔ)與實踐. 清華大學(xué)出版社, 2020.

張舒. GPU高性能計算之CUDA. 中國水利水電出版社, 2009.

蘇統(tǒng)華. CUDA并行程序設(shè)計：GPU編程指南. 機械出版社, 2014.

網(wǎng)絡(luò)資源

Linux高性能計算集群 -- Beowulf集群http://www.jointforce.com.cn/page/hardware_linux.html

CUDA C++ Programming Guidehttps://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html

周斌：NVIDIA CUDA初級教程視頻https://www.bilibili.com/video/BV1kx411m7Fk?from=search&seid=11175574756932230442&spm_id_from=333.337.0.0

蘇統(tǒng)華：CUDA并行程序設(shè)計https://www.bilibili.com/video/BV15E411x7yT?from=search&seid=11175574756932230442&spm_id_from=333.337.0.0

總結(jié)

以上是生活随笔為你收集整理的CUDA编程技术汇总的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。