?CUDA編程技術匯總

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

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

注1：文章內容會不定期更新。

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

一、高性能計算

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

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

二、軟件層面

2.1 CUDA

CUDA是一個并行計算開發環境，主要針對基于NVIDIA?GPU硬件。

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

2.2 編程模型

將系統分成主機端與設備端；設備端劃分為粗細兩層的并行結構；kernel函數由主機端調用，但在設備端運行。

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

2.3 線程模型

核函數的全部線程塊構成一個網格(Grid)，線程網格由多個線程塊（Block）組成；每個線程塊包含多個線程。在軟件層面，kernel以線程塊為單位并發執行。

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

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

kernel函數中內建變量：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)線程塊內線程總數通常要設置為32的倍數，可取64、128、256; e) 線程網格的維度通常由問題規模決定。

2.4 內存模型

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

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

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

全局存儲器、常數存儲器和紋理存儲器中的值在一個內核函數執行完成后將被繼續保持，可以被同一程序中的其他內核函數調用。

存儲器	位置	擁有緩存	訪問權限	變量生命周期	帶寬	延遲 （cycles）
register	GPU片內	N/A	device可讀/寫	與thread相同	~8TB/s	1
local_memory	板載顯存	無	device可讀/寫	與thread相同
shared_memory	GPU片內	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內存	無	host可讀/寫	可在程序中保持
pinned_memory	host內存	無	host可讀/寫	可在程序中保持

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

3.1 GPU硬件架構
三、硬件層面

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

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

3.2 Warp

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

四、優化策略

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

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

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

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

參考書籍

樊哲勇. CUDA編程：基礎與實踐. 清華大學出版社, 2020.

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

蘇統華. CUDA并行程序設計：GPU編程指南. 機械出版社, 2014.

網絡資源

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

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

總結

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

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。