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

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

(國防科技大學計算機學院，湖南 長沙 410073)

1 引言

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

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

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

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

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

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

2 乘積碼

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

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

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

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

(4) 計算并將擴展位附加到相應的行和列。

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

Figure 1 Structure of 2D Turbo product code圖1 二維TPC碼結構

3 Chase-Pyndiah譯碼算法

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

3.1 漢明碼譯碼

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

s=yH

(1)

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

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

3.2 Chase-Pyndiah算法

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

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

(2)

其中,i=1,…,n。

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

步驟3確定2p個擾動序列:

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

(3)

然后對2p個擾動序列進行譯碼:

(1) 首先，通過式(4)計算擾動序列的syndrome:

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

(4)

其中,H是漢明碼的奇偶校驗矩陣。

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

(5)

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

(3)計算奇偶校驗位zi:

(6)

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

步驟4計算已譯碼字的模擬權重為：

wi=- 〈zi,r〉

(7)

d=zm

(8)

(9)

步驟6更新外部信息E:

(10)

其中，β是可靠性因子。

步驟7更新列的軟輸入:

(11)

其中，α是縮放因子。

(12)

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

(13)

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

Figure 2 Process of Chase-Pyndiah algorithm圖2 Chase-Pyndiah算法計算過程

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

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

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

4.1 GPU和CUDA的架構

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

4.2 多線程計算和內(nèi)存層次結構

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

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

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

4.3 并行基本譯碼器

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

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

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

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

Figure 3 Constructing test patterns圖3 測試樣例的構建

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

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

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

4.4 多通道譯碼

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

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

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

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

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

4.5 并行TPC譯碼器

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

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

算法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 實驗結果

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

5.1 延遲和吞吐率

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

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

(14)

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

測試結果如表1所示。表1中的縮寫解釋如下：SCSE代表單通道和串行基本譯碼器，SCPE代表單通道和并行基本譯碼器，MCSE代表多通道和串行基本譯碼器，MCPE代表單通道和并行基本譯碼器。由表1可知，通過GPU譯碼比使用未經(jīng)優(yōu)化的SCSE譯碼器的CPU譯碼快近2倍。優(yōu)化并行方法的加速比大于12。多通道方法減少了平均延遲并進一步提高了吞吐率。最高吞吐率超過30 Mbps，是CPU吞吐率的44倍。測試樣例和有效碼字的計算在SCSE中是串行的，但在SCPE中是并行的。因此，SCPE的譯碼吞吐率比SCSE高6倍。這種改進主要歸功于基本譯碼器的優(yōu)化。并行基本譯碼器的譯碼塊是二維的，這個二維線程塊包含2p×n個線程，在(64,57)的情況下為1 024。此外，RTX 2080 Ti的計算能力為7.5，每塊的最大線程數(shù)為1 024，因此并行基本譯碼器在GPU上實現(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ù)量的增長情況，并非所有1～2 048通道數(shù)的結果都在直方圖中給出，本文選擇了8個典型數(shù)字來表示其增長趨勢。用N表示通道數(shù)，結果表明，當N>256時，多通道譯碼的吞吐率增加速度變緩。這種變緩的一個原因是數(shù)據(jù)傳輸?shù)拈_銷。由于輸入數(shù)據(jù)和譯碼輸出的大小隨N增加，CPU和GPU之間的通信延遲變得更高。因此，譯碼吞吐率受到該延遲的影響。另一個原因是GPU上資源的限制。RTX 2080 Ti具有68個SM和4 352個CUDA核心,每個SM的最大駐留塊數(shù)是32,因此理論上GPU上的最大駐留塊總數(shù)是2 176。然而，256通道TPC譯碼器需要16 384個塊用于并行譯碼,太大的通道數(shù)可能導致大多數(shù)塊必須等到SM空閑時才能計算。此外，每個SM上能使用的寄存器和存儲器也是有限的。例如，如果每個SM上的64K 32位寄存器用完，大多數(shù)中間數(shù)據(jù)將被存儲在本地存儲器中，訪問本地存儲器的速度比訪問寄存器的速度慢得多，因此寄存器耗盡也可能導致嚴重的訪問開銷。

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

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

5.3 CUDA流

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

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

Table 2 Performance using CUDA stream on RTX 2080 Ti

5.4 不同GPU上的譯碼

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

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

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倍。結果表明，GTX Titan V上更多的SM和CUDA核數(shù)有助于提高譯碼吞吐率。

Table 4 Throughput of different methods on different GPU platforms

6 結束語

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

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