基于GPU的并行Turbo乘積碼譯碼器*

李榮春，周 鑫，潘衡岳，牛 新，高 蕾，竇 勇

(國(guó)防科技大學(xué)計(jì)算機(jī)學(xué)院，湖南 長(zhǎng)沙 410073)

1 引言

Turbo乘積碼TPC(Turbo Product Code)是一類高性能前向糾錯(cuò)碼FEC(Forward Error Correction)。1994年，Pyndiah等[1]提出了一種基于軟譯碼和軟判決輸出的迭代譯碼算法，該算法在誤比特率BER(Bit Error Rate)方面接近最優(yōu)譯碼性能。文獻(xiàn)[2]中介紹的Turbo譯碼算法被稱為Chase-Pyndiah算法,它顯著降低了乘積碼的譯碼復(fù)雜性。因?yàn)檐涊斎胲涊敵鯯ISO(Soft-In Soft-Out)譯碼方案與卷積Turbo碼CTC(Convolution Turbo Code)的譯碼方法非常相似，所以TPC也稱為塊Turbo碼BTC(Block Turbo Code)。之后，TPC因其在高碼率下具有良好的BER性能和合理的譯碼復(fù)雜度而得到廣泛應(yīng)用。

TPC在現(xiàn)代通信網(wǎng)絡(luò)中發(fā)揮著重要作用，出現(xiàn)在幾個(gè)通信標(biāo)準(zhǔn)[3]中，例如IEEE 802.16[4]和IEEE 802.20[5]。此外，TPC還廣泛應(yīng)用于光通信系統(tǒng)[6]、衛(wèi)星通信應(yīng)用和數(shù)據(jù)存儲(chǔ)系統(tǒng)[7]。

許多研究都聚焦于提高譯碼性能和降低TPC譯碼器的譯碼復(fù)雜度[8 - 10]。但是，大多數(shù)工作都基于硬件平臺(tái)，例如專用集成電路(ASIC)[11]。基于圖形處理單元(GPU)的TPC譯碼器的實(shí)現(xiàn)很少。可以看出，一組二維乘積碼通常以二維矩陣形式出現(xiàn)，并且在譯碼過程中涉及許多矩陣的加法和乘法運(yùn)算。由于GPU具有強(qiáng)大的并行處理能力和浮點(diǎn)運(yùn)算能力，因此通常用于加速大規(guī)模矩陣運(yùn)算。因此，基于GPU譯碼是提高TPC譯碼效率的可行方案。與專用硬件譯碼器相比，在GPU上譯碼具有更好的靈活性和通用性，還可以節(jié)省開發(fā)時(shí)間。

軟件定義無線電SDR(Soft Defined Radio)的提出可以通過軟件配置支持各種通信標(biāo)準(zhǔn)，而無需改變硬件平臺(tái)[12]的理念。基于GPU的譯碼也是SDR的解決方案，在一些實(shí)例中，GPU譯碼器的吞吐率甚至優(yōu)于ASIC[13]。已有研究提出了許多高效的基于GPU的譯碼器，例如Viterbi譯碼器[14]、Turbo 譯碼器[15]和LDPC 譯碼器[16]。因此，基于GPU的TPC譯碼器也能夠滿足SDR的要求。

Chase-Pyndiah算法是一種迭代算法，每輪迭代之間存在數(shù)據(jù)依賴。盡管文獻(xiàn)[10,17]提出了幾種并行方法，但很少有并行實(shí)現(xiàn)基于諸如GPU之類的軟件平臺(tái)。文獻(xiàn)[18]提出了由BCH(Bose-Chaudhuri-Hocquenghem)碼組成的TPC的高吞吐率譯碼，實(shí)現(xiàn)了TPC的并行譯碼,且引入了幾種有效的存儲(chǔ)器控制方法來提高譯碼效率。

在本文中，TPC由漢明碼構(gòu)成，本文的目標(biāo)是簡(jiǎn)化譯碼過程并有效利用GPU上的資源。本文解決方案中考慮了細(xì)粒度和粗粒度并行性，不僅能夠同時(shí)譯碼TPC矩陣的每一行(或列)，而且并行計(jì)算擾動(dòng)序列和有效碼字。實(shí)驗(yàn)結(jié)果表明，與基于CPU的TPC譯碼器相比，本文方法譯碼延遲減少，而吞吐率顯著提高。

2 乘積碼

乘積碼是由Elias于1954年推出的。如文獻(xiàn)[19]中所述，可以使用2個(gè)短分組碼高效地構(gòu)建乘積碼。假設(shè)C1(n1,k1,d1)和C2(n2,k2,d2)為2個(gè)線性分組碼,其中ni,ki和di(i=1,2) 代表碼長(zhǎng)度、信息位數(shù)和最小漢明距離。C1和C2可用于構(gòu)建乘積碼P。使P=C1?C2，P的參數(shù)為(n1×n2,k1×k2,d1×d2)。乘積碼P的構(gòu)建過程如下所示:

(1) 將k1×k2信息位放在碼字矩陣的前k1行和k2列中;

(2) 使用碼C1編碼k1行;

(3) 使用碼C2編碼k2列;

(4) 計(jì)算并將擴(kuò)展位附加到相應(yīng)的行和列。

如上所述可以發(fā)現(xiàn)，通過使用擴(kuò)展的線性分組碼，僅通過非常小的碼率損失，可以大大增加TPC的最小漢明距離。圖1所示是二維TPC碼的結(jié)構(gòu)。

Figure 1 Structure of 2D Turbo product code圖1 二維TPC碼結(jié)構(gòu)

3 Chase-Pyndiah譯碼算法

文獻(xiàn)[2]中描述的Turbo譯碼算法適用于基于線性分組碼的任何乘積碼。因此，我們用擴(kuò)展?jié)h明碼構(gòu)造乘積碼，每行和列的譯碼方法對(duì)應(yīng)漢明碼譯碼方法。本文提出的基于GPU并行譯碼算法就是基于這種方法譯碼的。

3.1 漢明碼譯碼

假設(shè)C是二進(jìn)制漢明碼，其中C=(c1,c2,…,cn)。設(shè)H為C的校驗(yàn)矩陣,H中的所有列都不相同；y是經(jīng)過硬判決的二進(jìn)制向量。C的syndrome由式(1)得到：

s=yH

(1)

如果C中第i位出錯(cuò),那么s等于hi,其中hi是H的第i列。漢明碼的該特性可用于定位接收數(shù)組y中的錯(cuò)誤位置。

由于H在譯碼之前就已經(jīng)確定，可以構(gòu)造糾錯(cuò)表以定位錯(cuò)誤位置。hi是一個(gè)二進(jìn)制序列，可以轉(zhuǎn)換為十進(jìn)制值。設(shè)這個(gè)十進(jìn)制值是糾錯(cuò)表的索引，相應(yīng)的列索引是表的值。然后，一旦計(jì)算出syndrome，就可以找到y(tǒng)的錯(cuò)誤位置。

3.2 Chase-Pyndiah算法

通常，碼字通過二進(jìn)制相移鍵控BPSK(Binary Phase Shift Keyind modulation)調(diào)制，并且由{0,1}序列映射到{-1,+1}序列。然后它們通過加性高斯白噪聲AWGN(Additive White Gallssian Noise)信道傳輸。假設(shè)收到的序列是r=(r1,r2,…,rn)。以乘積碼行的譯碼為例，以下步驟描述了主要的譯碼過程：

步驟1可靠性數(shù)組|K|={|r1|,|r2|,…,|rn|},其中|ri|為接收值ri的絕對(duì)值，每行的輸入序列y={y1,y2,…,yn},其中：

(2)

其中,i=1,…,n。

步驟2生成2p個(gè)測(cè)試樣例。通過獲取|K|中p個(gè)最小值的索引，找到p個(gè)最不可靠的位置。對(duì)于每個(gè)測(cè)試樣例ti,i=1,2,…,2p,ti中的p個(gè)位置設(shè)置為2p個(gè)由0和1組成的可能組合，ti中的其他位置設(shè)置為0。

步驟3確定2p個(gè)擾動(dòng)序列:

zi=ti⊕y,i=1,2,…,2p

(3)

然后對(duì)2p個(gè)擾動(dòng)序列進(jìn)行譯碼:

(1) 首先，通過式(4)計(jì)算擾動(dòng)序列的syndrome:

si=ziHT,i=1,2,…,2p

(4)

其中,H是漢明碼的奇偶校驗(yàn)矩陣。

(2) 矩陣H在編碼某個(gè)漢明碼時(shí)可事先確定，因此可以確定由syndrome索引的查找表(LUT),有效地找到錯(cuò)誤位置。因此，可以通過式(5)糾正錯(cuò)誤位:

(5)

LUT(si)是zi的錯(cuò)誤位置，由si索引。

(3)計(jì)算奇偶校驗(yàn)位zi:

(6)

前3個(gè)步驟描述了Chase算法的搜索過程[20]，旨在獲得以y為中心,(di-1)，i=1,2 為半徑的范圍內(nèi)最可能的碼字。

步驟4計(jì)算已譯碼字的模擬權(quán)重為：

wi=- 〈zi,r〉

(7)

d=zm

(8)

(9)

步驟6更新外部信息E:

(10)

其中，β是可靠性因子。

步驟7更新列的軟輸入:

(11)

其中，α是縮放因子。

(12)

注意到，在列計(jì)算結(jié)束后，譯碼器的軟輸出會(huì)按式(13)更新:

(13)

完整的譯碼過程如圖2所示。

Figure 2 Process of Chase-Pyndiah algorithm圖2 Chase-Pyndiah算法計(jì)算過程

4 TPC譯碼器GPU實(shí)現(xiàn)

GPU上的流處理器(SM)提供強(qiáng)大的計(jì)算能力，成千上萬的線程同時(shí)在這些SM上工作。通常，通過盡可能多地并行完成工作，可以實(shí)現(xiàn)更低的延遲和更高的吞吐率。雖然GPU具有強(qiáng)大的計(jì)算能力，但GPU上的資源有限。例如，存儲(chǔ)器和寄存器是GPU上的有限資源。因此，在優(yōu)化算法時(shí)，需要了解每個(gè)線程的工作負(fù)載和進(jìn)行合理的存儲(chǔ)器分配。

本節(jié)研究并行TPC譯碼算法，以高效利用GPU上的資源。為了簡(jiǎn)化譯碼過程，對(duì)算法的某些部分進(jìn)行了修改。經(jīng)過優(yōu)化串行計(jì)算的部分，減少了譯碼延遲。此外，通過利用二維線程塊和多通道譯碼方法實(shí)現(xiàn)了更高的并行度。

4.1 GPU和CUDA的架構(gòu)

SM是GPU上的基本單元，啟動(dòng)核函數(shù)后，會(huì)將線程塊分配給某個(gè)SM。每個(gè)SM都有寄存器文件、緩存、CUDA核心和處理器所需的其他必要單元。可以一次將多個(gè)線程塊分配給同一SM,根據(jù)SM空閑資源情況來安排線程塊[21]。計(jì)算統(tǒng)一設(shè)備架構(gòu)(CUDA)為用戶提供了一種在GPU上利用計(jì)算資源的簡(jiǎn)便方法。CUDA將GPU上的硬件計(jì)算單元映射到邏輯線程網(wǎng)格和塊，在CUDA編程時(shí)，所有工作都由線程組織。

4.2 多線程計(jì)算和內(nèi)存層次結(jié)構(gòu)

Chase-Pyndiah算法是一種迭代譯碼算法，本文聚焦于優(yōu)化一輪迭代中的譯碼過程。通常，組成乘積碼的2組漢明碼C1和C2是同源碼字，這意味著n1=n2,k1=k2,d1=d2。乘積碼碼字P被組織為二維矩陣，矩陣的每行或每列是漢明碼字，并由基本譯碼器譯碼。在并行譯碼器中，n行或n列被同時(shí)譯碼。

令每行碼字或列碼字的參數(shù)為(n,k,d)。在并行譯碼器中，分配用于譯碼一行或一列的線程數(shù)始終等于碼字長(zhǎng)度，即n，所以有n×n個(gè)線程用于并行譯碼。因?yàn)槊總€(gè)塊的最大線程數(shù)是1 024(CUDA計(jì)算能力為7.5)，所以通過將每個(gè)行或列的譯碼計(jì)算分配給不同的線程塊以避免過載。因此，分配n個(gè)線程塊以同時(shí)譯碼n行或列。注意到，對(duì)一列的譯碼需要來自所有行的譯碼結(jié)果，因此在2次半迭代之間需要同步。該并行方法的優(yōu)點(diǎn)是在不改變譯碼程序的情況下支持任意大小的碼字。

GPU上的內(nèi)存層次結(jié)構(gòu)包括全局內(nèi)存、共享內(nèi)存和常量?jī)?nèi)存。全局內(nèi)存通常用于GPU和CPU之間的通信，因此輸入矩陣和估計(jì)結(jié)果存儲(chǔ)在全局存儲(chǔ)器中。GPU上的共享內(nèi)存速度更快，但容量更小，中間結(jié)果存儲(chǔ)在共享存儲(chǔ)器中，以減少訪存的開銷。在本文的TPC譯碼器中，輸入矩陣R存儲(chǔ)在全局存儲(chǔ)器中,在譯碼過程開始時(shí)它被加載到共享存儲(chǔ)器，以便進(jìn)行譯碼計(jì)算，并且在譯碼結(jié)束時(shí)將結(jié)果更新為R并存入全局存儲(chǔ)。

4.3 并行基本譯碼器

算法1是并行基本譯碼器算法。基本譯碼器是用于譯碼字輸入矩陣的一行或一列的基本譯碼單元。但是，如第3節(jié)介紹，此譯碼過程需要分階段執(zhí)行7個(gè)步驟。本節(jié)介紹了一種并行基本譯碼器，以簡(jiǎn)化該譯碼過程。

算法1并行基本譯碼器

輸入：p,z,r,lrp。

輸出：dist。

1.tx←threadIdx.x;

2.ty←threadIdx.y;

3.iftx

4.patternNum←(ty>>(p-tx-1));

5.zty[lrp[tx]]←patternNum⊕hdata[lrp[tx]];

6.end if

7.iftx

8.syndromeBinty←ztyHmod 2;

9.end if

10.iftx==0 then

11. forj=0→d-1 do

12.stv+=syndromeBinty<<(V-j-1);

13. end for

14. ifsty≠0 then

15.errPosty←LUT[sty];

16.zty[errPosty]←1-zty[errPostv];

17. end if

18.end if

盡管可以通過n個(gè)線程塊同時(shí)譯碼TPC，但在這種分配方案下，一輪迭代過程中的譯碼過程仍未被優(yōu)化。在計(jì)算得到最不可靠的位置之后，在這些位置處更新硬輸入數(shù)組，以構(gòu)建2p個(gè)測(cè)試樣例。這些測(cè)試樣例通過漢明譯碼方法譯碼，然后計(jì)算這些已譯碼字與軟輸入陣列之間的相關(guān)性。根據(jù)相關(guān)性，從2p個(gè)碼字中選擇最大似然碼字。在原始算法中，上述過程是串行的，這意味要串行處理測(cè)試樣例的構(gòu)造、漢明譯碼和相關(guān)的計(jì)算。但是，不同測(cè)試樣例之間不存在依賴關(guān)系，完全可以并行處理此過程，因此，本節(jié)研究并行基本譯碼器，以簡(jiǎn)化該譯碼過程。在優(yōu)化并行基本譯碼器時(shí)，我們修改了原始算法以提高譯碼效率，并實(shí)現(xiàn)了幾種優(yōu)化方法。

在并行基本譯碼器中，分配2p×n大小的二維線程塊用于譯碼。線程的垂直方向用于計(jì)算2p個(gè)測(cè)試樣例并將它們并行譯碼。用ty表示塊的列索引，用tx表示塊的行索引。在算法1中,z表示擾動(dòng)序列矩陣，它是大小為2p×n的short int矩陣；hdata是硬判決輸入{0,1}序列；lrp是已經(jīng)在先前步驟中計(jì)算過的最不可靠位置的數(shù)組。

首先，使用lrp來生成測(cè)試序列t。擾動(dòng)序列z是t與硬判決數(shù)組hdataxor運(yùn)算的結(jié)果。在我們的并行算法中，可以通過樣例數(shù)patternNum高效地計(jì)算擾動(dòng)序列。如圖3所示，patternNum的大小為0～2p-1。它的二進(jìn)制形式用于生成2p個(gè)不同的序列，這些序列由2p個(gè)0和1組合構(gòu)成。該方法用int數(shù)替換n-length數(shù)組，節(jié)省了2p×n×4字節(jié)共享內(nèi)存。

Figure 3 Constructing test patterns圖3 測(cè)試樣例的構(gòu)建

其次， 使用H和patternNum計(jì)算2p擾動(dòng)序列的syndrome，其中H是(n-1)×d奇偶校驗(yàn)矩陣。syndromeBin是2p×d矩陣， 是由z中每一行的syndrome值的二進(jìn)制表示組成。V表示校驗(yàn)位的數(shù)量。一旦計(jì)算出syndrome，就根據(jù)由syndrome索引的查找表校正錯(cuò)誤位。

最后，z中每一行的最后一位，即前n-1位的奇偶校驗(yàn)位，可通過并行歸約來進(jìn)行更新。偽代碼的最后一行計(jì)算zty和r之間2p的相關(guān)性。

本文所提出的并行基本譯碼器利用二維線程塊來譯碼每個(gè)行或列，這種并行方法在譯碼過程中減少了循環(huán)次數(shù)。例如，令p=4，然后在計(jì)算測(cè)試樣例和有效碼字時(shí)迭代次數(shù)為16。循環(huán)帶來的開銷使得這一過程在GPU上計(jì)算效率極低，因此可以利用并行基本譯碼器來提高譯碼效率。

4.4 多通道譯碼

在優(yōu)化TPC的譯碼算法時(shí)，不僅要考慮細(xì)粒度并行性，還要在GPU上實(shí)現(xiàn)粗粒度并行。

本文設(shè)計(jì)了多通道譯碼方法，以便有效地利用GPU上的資源并提高譯碼吞吐率。本文的并行基本譯碼器的塊大小為2p×n，n個(gè)線程塊被分配用于譯碼TPC的n行或列。設(shè)n為一個(gè)譯碼組，分配N個(gè)線程塊對(duì)來自N通道的碼字同時(shí)譯碼。

圖4描述了多通道譯碼器的主要思想。其中，N通道的輸入矩陣由N組塊譯碼。L代表最不可靠位，LUT是事先構(gòu)造好的查找表。在行譯碼器中，每個(gè)塊譯碼輸入矩陣的一行。對(duì)于列譯碼器，每列輸入矩陣被并行譯碼。H表示奇偶校驗(yàn)矩陣。s是測(cè)試樣例和H的乘積。ErrPos是通過搜索由syndromes索引的查找表獲得的錯(cuò)誤位置,其中FEC對(duì)碼字進(jìn)行糾錯(cuò)。

在圖4所示的的多通道TPC譯碼器中，通道數(shù)量可根據(jù)GPU的計(jì)算能力進(jìn)行配置和自適應(yīng)。

Figure 4 Multi-channel parallel TPC decoder圖4 多通道并行TPC譯碼器

4.5 并行TPC譯碼器

并行TPC譯碼器由2部分組成：一部分是行譯碼核函數(shù)，它在前半個(gè)迭代時(shí)啟動(dòng)；另一部分是列譯碼核函數(shù)，一旦行譯碼核函數(shù)完成就會(huì)啟動(dòng)。

算法2是整個(gè)迭代譯碼過程。輸入為TPC矩陣R、最不可靠位數(shù)p、迭代次數(shù)maxIter和syndromeNum。因?yàn)榱凶g碼器按列譯碼并更新R，所以它必須等到行譯碼器更新完R的所有行才能開始執(zhí)行。碼字估計(jì)值由最后一步硬判決結(jié)果決定。

算法2迭代譯碼算法

輸入：R,p,sydromeNum,maxIter。

輸出：EstimatedCodes。

1.foriter=0 →maxIterdo

4.end for

算法3并行行譯碼算法

輸入：R,p,syndromeNum。

1.tx←threadIdx.x;

2.ty←threadIdx.y;

3.r←Rblockidx.x;//R的第blockIdx.x行

4.E←|r|;

5.ifr[tx]≥0 then

6.hdata[tx]←1

7.else

8.hdata[ty]←0

9.end if

10.lrp←BitonicSort(E);

11.zty←hdata;//z的第ty行

12.dist←ElementaryDecoder(p,z,r,lrp)

13.iftx==0 &ty==0 then

14.tmp←MAXVAL;

15. fori=0→2p-1 do

16. iftmp

17.min←i

18. end if

19. end for

20.end if

21.wc←MAXVAL;

22.fors=0→2p-1 do

23. ifzty==(1-zty) &dist[s]

24.wc←dist[s]

25. end if

26.end for

27.ifwc

28.E←(2zmin-1)(wc-dist[min])-r

29.else

30.E←(2zmin-1)β

31.end if

5 實(shí)驗(yàn)結(jié)果

本節(jié)測(cè)試所提出的并行TPC譯碼器在GPU上的延遲和吞吐率，為了驗(yàn)證優(yōu)化方法的效果，測(cè)試了不同版本的程序，此外，還測(cè)試了基于不同GPU的譯碼吞吐率。測(cè)試結(jié)果表明，本文的并行譯碼器比串行譯碼器具有更低的延遲和更高的吞吐率。實(shí)驗(yàn)用的CPU是Intel i7-8700K，3.7 GHz，GPU是NVIDIA RTX 2080 Ti。CUDA版本是10.1。

5.1 延遲和吞吐率

測(cè)試樣例中的TPC由(64,57)的擴(kuò)展?jié)h明碼構(gòu)成，最后一位是奇偶校驗(yàn)位，用于檢查碼字是否有效。參數(shù)p為4，迭代次數(shù)為6，縮放因子α=1和β=0.5，信噪比(SNR)為3 dB。實(shí)驗(yàn)驗(yàn)證了本文GPU基譯碼器的正確性，結(jié)果表明BER性能接近Matlab中的仿真結(jié)果。

當(dāng)使用CPU譯碼時(shí)， 延遲為5 709 μs， 吞吐率為0.7 Mbps。加速比的計(jì)算方法如式(14)所示：

(14)

其中,tCPU是CPU的譯碼延遲，tGPU是GPU的平均譯碼延遲。

測(cè)試結(jié)果如表1所示。表1中的縮寫解釋如下：SCSE代表單通道和串行基本譯碼器，SCPE代表單通道和并行基本譯碼器，MCSE代表多通道和串行基本譯碼器，MCPE代表單通道和并行基本譯碼器。由表1可知，通過GPU譯碼比使用未經(jīng)優(yōu)化的SCSE譯碼器的CPU譯碼快近2倍。優(yōu)化并行方法的加速比大于12。多通道方法減少了平均延遲并進(jìn)一步提高了吞吐率。最高吞吐率超過30 Mbps，是CPU吞吐率的44倍。測(cè)試樣例和有效碼字的計(jì)算在SCSE中是串行的，但在SCPE中是并行的。因此，SCPE的譯碼吞吐率比SCSE高6倍。這種改進(jìn)主要?dú)w功于基本譯碼器的優(yōu)化。并行基本譯碼器的譯碼塊是二維的，這個(gè)二維線程塊包含2p×n個(gè)線程，在(64,57)的情況下為1 024。此外，RTX 2080 Ti的計(jì)算能力為7.5，每塊的最大線程數(shù)為1 024，因此并行基本譯碼器在GPU上實(shí)現(xiàn)了較高的占用率。

Table 1 Performance comparison of various methods on RTX 2080 Ti

5.2 通道數(shù)量

從表1中可以發(fā)現(xiàn)，在多信道的情況下，吞吐率的提高不如單信道的情況那樣顯著。MCPE的吞吐率僅比MCSE高2倍。圖5顯示了吞吐率隨通道數(shù)量的增長(zhǎng)情況，并非所有1～2 048通道數(shù)的結(jié)果都在直方圖中給出，本文選擇了8個(gè)典型數(shù)字來表示其增長(zhǎng)趨勢(shì)。用N表示通道數(shù)，結(jié)果表明，當(dāng)N>256時(shí)，多通道譯碼的吞吐率增加速度變緩。這種變緩的一個(gè)原因是數(shù)據(jù)傳輸?shù)拈_銷。由于輸入數(shù)據(jù)和譯碼輸出的大小隨N增加，CPU和GPU之間的通信延遲變得更高。因此，譯碼吞吐率受到該延遲的影響。另一個(gè)原因是GPU上資源的限制。RTX 2080 Ti具有68個(gè)SM和4 352個(gè)CUDA核心,每個(gè)SM的最大駐留塊數(shù)是32,因此理論上GPU上的最大駐留塊總數(shù)是2 176。然而，256通道TPC譯碼器需要16 384個(gè)塊用于并行譯碼,太大的通道數(shù)可能導(dǎo)致大多數(shù)塊必須等到SM空閑時(shí)才能計(jì)算。此外，每個(gè)SM上能使用的寄存器和存儲(chǔ)器也是有限的。例如，如果每個(gè)SM上的64K 32位寄存器用完，大多數(shù)中間數(shù)據(jù)將被存儲(chǔ)在本地存儲(chǔ)器中，訪問本地存儲(chǔ)器的速度比訪問寄存器的速度慢得多，因此寄存器耗盡也可能導(dǎo)致嚴(yán)重的訪問開銷。

Figure 5 Throughput with various channel numbers圖5 多通道下變化的吞吐率

由于上述原因可知，更多的通道數(shù)不一定帶來更高的吞吐率，選取合適的通道數(shù)取決于GPU的計(jì)算能力。

5.3 CUDA流

多個(gè)核函數(shù)可以在GPU上同時(shí)執(zhí)行。CUDA應(yīng)用程序通過CUDA流管理核函數(shù)啟動(dòng)和內(nèi)存?zhèn)鬏數(shù)膱?zhí)行順序。如果同時(shí)啟動(dòng)多個(gè)核函數(shù)則會(huì)為每個(gè)核函數(shù)分配流ID，具有相同流ID的核函數(shù)按順序啟動(dòng)，具有不同流ID的核函數(shù)同時(shí)執(zhí)行。存儲(chǔ)器傳輸和核函數(shù)執(zhí)行的重疊可以提高譯碼吞吐率。

本節(jié)測(cè)試了MCSE和MCPE的吞吐率，結(jié)果如表2所示，并發(fā)執(zhí)行和內(nèi)存?zhèn)鬏數(shù)耐掏侣试黾恿舜蠹s6%。帶來的加速效果不夠明顯，這是因?yàn)榈g碼過程引入了同步開銷。

Table 2 Performance using CUDA stream on RTX 2080 Ti

5.4 不同GPU上的譯碼

本文并行譯碼器的性能是在不同的GPU上測(cè)量的，用于測(cè)試的CPU平臺(tái)有NVIDIA GeForce GTX Titan V、RTX 2080 Ti和GTX 1080 Ti。表3給出了這些GPU平臺(tái)的SM數(shù)量、CUDA核心數(shù)和內(nèi)存大小。這些GPU具有不同的計(jì)算能力，可用于衡量本文并行譯碼器在不同規(guī)模的計(jì)算資源下的執(zhí)行的性能。結(jié)果如表4所示。

基于GTX 1080 Ti的SCSE吞吐率小于基于CPU的吞吐率。此結(jié)果表明SCSE無法有效利用GPU上的計(jì)算資源。

Table 3 Performance comparison on various GPU platforms

與RTX 2080 Ti相比，GTX Titan V上的CUDA核心數(shù)量和SM數(shù)量增加了大約20%以上。因此，GTX Titan V上的MCSE和MCPE在表4中具有更高的吞吐率和加速比。此外，SCPE的吞吐率是SCSE吞吐率的16倍。結(jié)果表明，GTX Titan V上更多的SM和CUDA核數(shù)有助于提高譯碼吞吐率。

Table 4 Throughput of different methods on different GPU platforms

6 結(jié)束語

本文提出了一種基于GPU的并行TPC譯碼器，可同時(shí)譯碼矩陣的所有行或列。研究了并行基本譯碼器，以實(shí)現(xiàn)更高的吞吐率和更低的譯碼延遲。所有測(cè)試樣例和有效碼字都是并行生成的。此外，本文還提出了一種多通道TPC譯碼器，以進(jìn)一步提高譯碼吞吐率。

本文測(cè)試了4種不同譯碼器SCSE、SCPE、MCSE和MCPE的延遲和吞吐率，還分析了多通道譯碼器的局限性。最后，我們測(cè)試了本文譯碼器在不同GPU上的性能，結(jié)果表明，該方法具有良好的可擴(kuò)展性。