劉朝華 李小花 張紅強(qiáng)
摘 要:本文將基于CUDA平臺(tái)的GPU并行程序設(shè)計(jì)模式引入自動(dòng)控制系統(tǒng)并行編程實(shí)踐教學(xué)中。通過介紹CUDA體系結(jié)構(gòu)與編程模式,實(shí)例教學(xué)使自動(dòng)控制專業(yè)學(xué)生對(duì)CUDA并行程序設(shè)計(jì)有著深刻的認(rèn)識(shí),將有助于同學(xué)們掌握CUDA并行計(jì)算平臺(tái)在工業(yè)控制程序開發(fā)中的應(yīng)用,增強(qiáng)同學(xué)們對(duì)自動(dòng)控制專業(yè)學(xué)習(xí)的興趣和實(shí)踐動(dòng)手能力。
關(guān)鍵詞:CUDA平臺(tái);自動(dòng)控制;并行編程;實(shí)踐教學(xué)
一、引言
隨著工業(yè)控制系統(tǒng)規(guī)模日益增大,控制算法日益復(fù)雜,工業(yè)數(shù)據(jù)日益激增,工業(yè)現(xiàn)場(chǎng)對(duì)系統(tǒng)控制與優(yōu)化調(diào)度的實(shí)時(shí)性要求日益提高,給現(xiàn)代控制系統(tǒng)程序設(shè)計(jì)提出挑戰(zhàn)。工業(yè)控制系統(tǒng)中存在大量的重復(fù)性計(jì)算,如快速傅里葉變換(FFT)、濾波、矩陣與數(shù)值計(jì)算等,給傳統(tǒng)控制系統(tǒng)帶來巨大時(shí)間開銷。以多核計(jì)算機(jī)為代表的并行計(jì)算技術(shù)為工業(yè)控制系統(tǒng)設(shè)計(jì)與開發(fā)提供了技術(shù)支持。特別是GPU(Graphic Processing Unit)通用并行計(jì)算技術(shù)的推廣[1],由于其高存儲(chǔ)器帶寬、低功耗等諸多優(yōu)點(diǎn),已被廣泛應(yīng)用于科學(xué)與工業(yè)等領(lǐng)域中,在提升實(shí)際問題的求解效率方面發(fā)揮了作用。CUDA(Compute Unified Device Architecture,計(jì)算統(tǒng)一設(shè)備架構(gòu))由英偉達(dá)公司于2007年發(fā)布的一種將GPU作為并行計(jì)算的軟硬件體系架構(gòu),可用于解決工程科學(xué)中的復(fù)雜計(jì)算,開啟了GPU通用并行計(jì)算時(shí)代[2]。研究表明在浮點(diǎn)運(yùn)算、數(shù)值計(jì)算等密集型計(jì)算方面,基于CUDA平臺(tái)的GPU計(jì)算系統(tǒng)可獲得數(shù)十倍于傳統(tǒng)CPU的加速性能[3]。因此,相比于傳統(tǒng)微機(jī)控制方式,基于CUDA的GPU并行計(jì)算適用于數(shù)據(jù)量大,實(shí)時(shí)性要求高的工業(yè)控制領(lǐng)域。
現(xiàn)有高校自控專業(yè)程序設(shè)計(jì)課程主要是基于串行編程方式的教學(xué),難以適應(yīng)工業(yè)信息化與知識(shí)自動(dòng)化對(duì)控制系統(tǒng)實(shí)時(shí)性能與數(shù)據(jù)存儲(chǔ)的需求。本文將基于CUDA平臺(tái)的GPU并行程序設(shè)計(jì)模式引入自動(dòng)控制系統(tǒng)并行程序設(shè)計(jì)實(shí)踐教學(xué)中。有利于同學(xué)掌握先進(jìn)計(jì)算技術(shù)在自動(dòng)控制系統(tǒng)中的開發(fā)與應(yīng)用,拓寬其專業(yè)視野。最后對(duì)自動(dòng)控制程序設(shè)計(jì)與CUDA并行編程實(shí)踐教學(xué)進(jìn)行了思考與總結(jié)。
二、CUDA體系架構(gòu)和編程模式
(一)CUDA體系結(jié)構(gòu)。CUDA框架包括硬件和軟件兩部分,拓寬了GPU程序開發(fā)與運(yùn)行的可操作性。GPU的硬件包括運(yùn)算核心和存儲(chǔ)器,其中運(yùn)算核心流處理器構(gòu)成流多處理器(Streaming Multi-Processor,SM),每個(gè)SM都帶共享存儲(chǔ)器,所有SM共享GPU的全局、常量與紋理存儲(chǔ)器。CUDA的軟件包括CPU代碼和GPU代碼,其中CPU負(fù)責(zé)處理邏輯性的串行事務(wù)控制,GPU負(fù)責(zé)處理大量重復(fù)性計(jì)算任務(wù)。在GPU上執(zhí)行的函數(shù)稱為核函數(shù),當(dāng)核函數(shù)被CPU代碼激活時(shí),GPU中在邏輯上的兩層線程組(頂層為線程網(wǎng)格(Grid),下層為線程塊(Block,包含若干個(gè)線程(Thread))并行執(zhí)行。一個(gè)內(nèi)核函數(shù)網(wǎng)格中的線程塊并行和線程塊中的線程兩個(gè)層次的并行方式來提高數(shù)據(jù)吞吐量和執(zhí)行效率。各個(gè)線程塊并行執(zhí)行,線程塊間無法通信,也沒有執(zhí)行順序,但同一個(gè)線程塊中的線程可以通過共享存儲(chǔ)器方式交換數(shù)據(jù)。
(二)CUDA編程模式。在CUDA架構(gòu)中,CPU與GPU是一種異構(gòu)協(xié)同并行計(jì)算模式。其中CPU為主機(jī),GPU為設(shè)備。CPU與GPU各自擁有獨(dú)立的存儲(chǔ)器地址空間:顯存和內(nèi)存。程序員可以應(yīng)用高級(jí)語言(C/C++/fortan)對(duì)GPU內(nèi)部計(jì)算資源進(jìn)行訪問。CUDA編程基本步驟為:步驟1:加載C(或C++)和CUDA頭文件,啟動(dòng)cudaSetDevice()配置GPU設(shè)備;初始化CPU和GPU數(shù)據(jù)空間;步驟2:調(diào)用設(shè)備端(GPU)的內(nèi)核函數(shù)計(jì)算;①從顯存讀取數(shù)據(jù)到共享存儲(chǔ)器或寄存器內(nèi);②對(duì)數(shù)據(jù)進(jìn)行并行計(jì)算和處理;③將處理后的數(shù)據(jù)寫回顯存。
步驟3:將顯存中的結(jié)果回讀到內(nèi)存;步驟4:使用CPU對(duì)數(shù)據(jù)進(jìn)行處理并輸出結(jié)果;步驟5:釋放內(nèi)存和顯存空間并退出CUDA。
三、實(shí)例講解
以一個(gè)簡(jiǎn)單的CUDA程序設(shè)計(jì)實(shí)例來講解,加深自動(dòng)控制專業(yè)同學(xué)對(duì)CUDA平臺(tái)GPU并行程序設(shè)計(jì)過程的理解??紤]給一個(gè)向量V{1:N}上每一位矢量的加上常數(shù)。其中CUDA程序C代碼如下:
#define N 100;
_global_void Vadd(int*V,int m)
{int index=blockIdx.x;if(indexint main(void){ int host_V[N]; int *device_V; cudaMalloc((void**)&device_V, L * sizeof(int)); for (int k=0; k