基于GPU的高吞吐量QC-LDPC碼編碼器實(shí)現(xiàn)

段運(yùn)德，黎 勇

(1.重慶郵電大學(xué) 通信與信息工程學(xué)院，重慶 400065；2.重慶大學(xué) 計(jì)算機(jī)學(xué)院，重慶 400044)

0 引 言

1962年，R.Gallager[1]在他的博士論文中第一次提出了低密度奇偶校驗(yàn)(low-density parity-check, LDPC)碼。當(dāng)時(shí)由于計(jì)算機(jī)技術(shù)并不發(fā)達(dá)，LDPC碼譯碼復(fù)雜度高，導(dǎo)致其優(yōu)秀的性能未被發(fā)現(xiàn)。直到1996年D.J.C.Mackay等[2]重新研究LDPC碼，發(fā)現(xiàn)其性能可以接近香農(nóng)限，LDPC碼才重新開始受到重視，并迅速成為信道編碼研究中的重要方向。QC-LDPC碼是一類具有準(zhǔn)循環(huán)結(jié)構(gòu)的LDPC碼，因?yàn)槠溲h(huán)特性可以降低編譯碼復(fù)雜度，目前已被應(yīng)用在很多標(biāo)準(zhǔn)和系統(tǒng)中，比如IEEE 802.11n(WLAN)，IEEE 802.16e(WiMAX)[3]，量子秘鑰分發(fā)[4]，ATSC3.0[5]等。2017年，3GPP標(biāo)準(zhǔn)會(huì)議將QC-LDPC碼確定為5G通信系統(tǒng)的長碼標(biāo)準(zhǔn)[6]。

因?yàn)閷?shí)際系統(tǒng)往往選擇QC-LDPC碼，而不是其他LDPC碼，因此，提高QC-LDPC碼的編碼速率有重要的實(shí)際意義。為提高編碼效率，大家普遍選擇現(xiàn)場可編程邏輯門陣列(field programmable gate array, FPGA)、專有電路(application-specific integrated circuit, ASIC)等硬件解決方案[7]，硬件方案通常只針對(duì)特定碼型或碼率的LDPC碼有效，通用性差。相比而言，圖形處理單元(graphics processing unit, GPU)加速是一種更便宜、更靈活、更高效的軟件方案。本文通過把QC-LDPC碼的校驗(yàn)矩陣轉(zhuǎn)換為準(zhǔn)循環(huán)的生成矩陣，再基于編碼過程中的并行部分和GPU的內(nèi)存、線程結(jié)構(gòu)，給出一種吉比特速率的軟件編碼方案。該方案不僅通過單個(gè)碼字編碼并行化以及多個(gè)碼字同時(shí)編碼來提高編碼吞吐量，而且通過對(duì)片上常量內(nèi)存和共享內(nèi)存的合理使用進(jìn)一步加快了編碼速率。仿真結(jié)果表明，相較于已有的GPU編碼方案，本文的方案通用性更強(qiáng)，編碼吞吐量比文獻(xiàn)[8]高；相較于其他硬件快速編碼方案，本文的編碼器在通用性和編碼吞吐量上均明顯高于文獻(xiàn)[9]的FPGA編碼器和文獻(xiàn)[10]的CMOS編碼器。

1 通過校驗(yàn)矩陣計(jì)算QC-LDPC碼準(zhǔn)循環(huán)結(jié)構(gòu)的生成矩陣

QC-LDPC碼的校驗(yàn)矩陣由具有準(zhǔn)循環(huán)特性的子矩陣組成，表示為

(1)

(1)式中：c

rank(H)=rank(D)=r

(2)

設(shè)得到的D矩陣為

(3)

H中選擇q列后，剩下的t-q列構(gòu)成矩陣B，表示為

(4)

將H按列重排得到Hqc=[B|D]，Hqc同樣為該QC-LDPC碼的校驗(yàn)矩陣。設(shè)該碼的生成矩陣為Gqc，則有

(5)

(5)式中：Gqc為k×n階矩陣；n為該QC-LDPC碼碼長，n=tb；k為信息位長，k=tb-r；O是cb×(tb-r)階零矩陣。由(5)式和文獻(xiàn)[11]中的算法即可解出矩陣Gqc，且Gqc有如下結(jié)構(gòu)

(6)

(7)

Q=

(8)

(7)—(8)式中：I為b階單位陣；O為b階零矩陣；Gi,j為b階方陣，且具有準(zhǔn)循環(huán)結(jié)構(gòu)，即Gi,j中每行向量是其前一行向量向右循環(huán)移動(dòng)一位得到；Oi,j為全零矩陣；qrow代表有qrow行子矩陣。Oi,j，Qi,j列數(shù)都為b，在計(jì)算Gqc時(shí)可以得到Oi,j和Qi,j的行數(shù)。本文先將Qi,j的行數(shù)保存在數(shù)組qNum[qrow]中以備后面程序使用，表示為

qNum[qrow]={qNum0,qNum1,…,qNumqrow}

(9)

(9)式中，qNumi對(duì)應(yīng)Q的第i行子矩陣Qi,·的行數(shù)。這樣得到的Qi,j同樣具有準(zhǔn)循環(huán)結(jié)構(gòu)，即Qi,j中每行向量是其前一行向量向右循環(huán)移動(dòng)一位得到的。設(shè)gi,j為Gi,j的首行向量，qi,j為Qi,j的首行向量，則通過所有g(shù)i,j和qi,j的循環(huán)移位即可以得到矩陣G和Q，從而得到Gqc。這里，定義和gi,jqi,j為Gqc的生成向量。

通過上述計(jì)算得到的生成矩陣Gqc通常具有(6)式的結(jié)構(gòu)，即由子矩陣G和Q組成。但如果給定的QC-LDPC碼的校驗(yàn)矩陣H是非奇異矩陣，且q=c，此時(shí)通過上述計(jì)算得到的Gqc只由子矩陣G組成，即Gqc=[G]。

在求解Gqc的過程中，需要多次求GF(2)中矩陣的秩。為了加快求秩速度，本文采用了文獻(xiàn)[12]中的稀疏矩陣高斯消去算法求秩。

通過以上處理，從QC-LDPC碼的校驗(yàn)矩陣H得到了具有準(zhǔn)循環(huán)結(jié)構(gòu)的生成矩陣Gqc。根據(jù)其循環(huán)特性，即由首行向量gi,j和qi,j循環(huán)移位可得到Gqc。因此，在編碼程序中，僅需存儲(chǔ)生成向量gi,j和qi,j，節(jié)省了存儲(chǔ)空間。

2 QC-LDPC碼的GPU編碼

基于GPU的并行算法設(shè)計(jì)通常是一項(xiàng)具有挑戰(zhàn)性的工作。必須慎重考慮線程的數(shù)量、內(nèi)存的大小以及寄存器等資源的限制。同時(shí)為了獲得正確的結(jié)果，必須注意線程的同步。CUDA(compute unied device architecture)是NVIDIA的GPU提供并行計(jì)算的平臺(tái)和編程接口，本文將基于CUDA平臺(tái)編程仿真。在設(shè)計(jì)編碼時(shí)，本文將從以下一些方面進(jìn)行優(yōu)化。表1給出了下文敘述中用到的一些參數(shù)。

CUDA核函數(shù)初始化比CPU函數(shù)初始化需要更多時(shí)間，因此編碼時(shí)，核函數(shù)將同時(shí)輸入N=32×s(s為正整數(shù))個(gè)信息序列，然后在運(yùn)行一次核函數(shù)時(shí)同時(shí)對(duì)N個(gè)碼字進(jìn)行編碼，以加快編碼速度。CUDA執(zhí)行核函數(shù)的調(diào)度單位為線程束，每個(gè)線程束包含32個(gè)線程，因此，選擇N為32的倍數(shù)可以加快線程調(diào)度。

表1 文中重要參數(shù)

本文將定義結(jié)構(gòu)體Bitset保存長度為b個(gè)比特的向量，如gi,j和qi,j向量。本文中Hqc，H，G等矩陣的子矩陣階數(shù)為b，也用該結(jié)構(gòu)體來保存。該結(jié)構(gòu)體偽代碼如下。

//一個(gè)int變量長32比特

const int g_intLen←32;

//g_num為保存長為b的向量需要int的個(gè)數(shù)

const int g_num←(b-1)/g_intLen+1;

struct Bitset{

int val[g_num];

}

CUDA常量內(nèi)存是一種64 KByte的只讀緩存，它提供比全局內(nèi)存更大的訪問帶寬。編碼時(shí)生成矩陣Gqc是在第1節(jié)中已經(jīng)求解出來的常量，不會(huì)改變，因此，將Gqc的生成向量gi,j和gi,j按圖1的格式存儲(chǔ)在常量內(nèi)存的Bitset類型的二維數(shù)組generatorMtx中。

圖1中，t,q,qrow參數(shù)含義見表1。使用CUDA常量內(nèi)存存儲(chǔ)Gqc的二維數(shù)組代碼為

__constant__ __device__

Bitset generatorMtx[t-q+qrow][q];

圖1 矩陣Gqc在常量內(nèi)存中的映射Fig.1 Map of Gqc stored in constant memory

編碼算法的CUDA核函數(shù)原型為

__global__

voidcuEncode(int *mess, int *code);

其中，mess是一個(gè)數(shù)組指針，保存N個(gè)信息序列，該數(shù)組包含N×k個(gè)比特。code為編碼輸出的碼字?jǐn)?shù)組指針，數(shù)組包含N×n個(gè)比特。

當(dāng)調(diào)用CUDA核函數(shù)編碼時(shí)需要進(jìn)行線程分配，本文將分配二維線程塊numThread，其維度為(t-q+qrow)×q，numThread和存儲(chǔ)Gqc的二維數(shù)組generatorMtx具有相同的維度。在同一個(gè)線程塊中，每個(gè)線程將同時(shí)讀取gi,j和qi,j，同時(shí)進(jìn)行計(jì)算。為了充分利用GPU線程，本文還將定義一維的線程網(wǎng)格numBlock，網(wǎng)格包含N個(gè)線程塊，每個(gè)線程塊中進(jìn)行單個(gè)碼字的編碼，核函數(shù)運(yùn)行一次將同時(shí)編得N個(gè)碼字。調(diào)用該核函數(shù)的代碼為

dim3 numThread(t-q+qrow, q), numBlock(N);

cuEncode << > > (mess, code);

此外，在核函數(shù)中，本文使用共享內(nèi)存用于一個(gè)線程塊中所有線程的通信，共享內(nèi)存的變量作為中間變量，存儲(chǔ)累加和，即將同一個(gè)線程塊中所有線程的計(jì)算結(jié)果累加到該變量中。由于核函數(shù)訪問共享內(nèi)存的速度比全局內(nèi)存快10倍，因此，使用共享內(nèi)存可以減少一個(gè)線程塊中所有線程的內(nèi)存訪問時(shí)間，加快編碼。

以上是本文設(shè)計(jì)GPU編碼時(shí)，從變量存儲(chǔ)方面做的一些優(yōu)化。接下來，本文將實(shí)現(xiàn)GPU編碼核函數(shù)。設(shè)m為一個(gè)信息序列向量，相應(yīng)的碼字為c，由于QC-LDPC碼是一種線性分組碼，則有

c=mGqc

(10)

根據(jù)前面的論述及(10)式，GPU編碼核函數(shù)的偽代碼如下。

算法：GPU并行編碼。

輸入：函數(shù)傳入N個(gè)信息序列保存在CUDA全局內(nèi)存數(shù)組mess[N×k]中。

輸出：最后得到的N個(gè)碼字輸出到數(shù)組code[N×n]中。

1)定義共享內(nèi)存數(shù)組，即__shared__ Bitset codTmp[q]，并將其初始化為全零。

2) __syncthreads()，同步所有線程塊。

3)bTmp←generatorMtx[threadIdx.x][threadIdx.y]

4)if threadIdx.x>= t-q then

//計(jì)算第threadIdx.x個(gè)線程塊處理的信息

//序列對(duì)應(yīng)的起始位置bitLoc，其中qNum[i]為

//(9)式中得到的數(shù)據(jù)，k為信息位長。

for j←bitLoc to bitLoc + qNum[i]-1 do

loc ← j + k × blockIdx.x;

if mess[loc] == 1 then

//異或相加

codTmp[threadIdx.y]←

codTmp[threadIdx.y]⊕bTmp;

end if

bTmp向右循環(huán)移動(dòng)一位；

end for

else

bitLoc ← threadIdx.x × b;

//計(jì)算存放碼字的起始位置messLoc，

//用于存放信息比特。如果校驗(yàn)矩陣H滿秩，

//則得到系統(tǒng)碼，否則碼字中僅前一部分

//比特為信息比特，n為碼長。

messLoc←threadIdx.x×b+n×blockIdx.x;

for j←bitLoc to bitLoc + b-1 do

loc←j+k×blockIdx.x;

if mess[loc] == 1 then

codTmp[threadIdx.y]←

codTmp[threadIdx.y]⊕bTmp

end if

bTmp向右循環(huán)移動(dòng)一位；

if threadIdx.y == 0 then

code[messLoc]←mess[loc];

++messLoc;

end if

end for

end if

5) __syncthreads()，同步所有線程塊。

6)

if threadIdx.x == 0 then

messLoc←blockIdx.x×n+(t-q)×b

+threadIdx.y×b;

for i←b-1 to 0 do

code[messLoc]←codTmp[threadIdx.y][i];

++messLoc;

end for

end if

上述核函數(shù)中，參數(shù)的含義如表1。threadIdx，blockIdx為CUDA中索引當(dāng)前線程和當(dāng)前線程塊的變量。__syncthreads()為CUDA中同步所有線程的函數(shù)。

調(diào)用上述核函數(shù)產(chǎn)生N個(gè)碼字保存在數(shù)組code[N·n]，code數(shù)組中碼字的排列方式如圖2。

圖2 數(shù)組中N個(gè)碼字的排列方式Fig.2 Arrange of the N codes in the array

3 GPU編碼的仿真分析

本文仿真采用Nvidia TITAN Xp顯卡平臺(tái)，使用英偉達(dá)官方工具集CUDA Toolkit 9.2，在Visual Studio 2017下編譯程序，平臺(tái)其他參數(shù)如表2。

表2 仿真電腦配置

實(shí)驗(yàn)中測試了碼長從幾百到一萬多比特的幾種高碼率QC-LDPC碼。相應(yīng)校驗(yàn)矩陣大小及Tanner圖的邊數(shù)如表3。

表3 仿真用的QC-LDPC碼

為了評(píng)估編碼算法的性能，本文計(jì)算了編碼吞吐量Kraw，其粗略計(jì)算式為

(11)

(11)式中：tim為對(duì)一個(gè)碼字編碼平均所用時(shí)間，單位為s；messLen為該QC-LDPC碼的信息位長，單位為bit。

在進(jìn)行GPU編碼時(shí)，本文為核函數(shù)分配N個(gè)線程塊，N為32的倍數(shù)，運(yùn)行一次核函數(shù)編得N個(gè)碼字。本文測試了隨著N的增加，GPU編碼吞吐量的變化情況，仿真結(jié)果如圖3。

圖3 核函數(shù)分配不同線程塊后的編碼性能Fig.3 Encoding performance with kernel function assigneddifferent number of thread blocks

從仿真結(jié)果來看，隨著N的增大，這幾種QC-LDPC碼的吞吐量在前期增加較快，但在N大約增加到20×32后，編碼吞吐量不再增加，且在一個(gè)范圍中波動(dòng)。這是由于前期GPU存在很多空閑資源，隨著N的增大，分配給核函數(shù)的線程增多，核函數(shù)運(yùn)行一次同時(shí)編得的碼字增加，吞吐量增加，編碼效率明顯提高。但后期，GPU中使用的常量內(nèi)存、全局內(nèi)存、共享內(nèi)存等資源達(dá)到上限，因此，編碼吞吐量停止增加。

實(shí)驗(yàn)中也將傳統(tǒng)的LDPC碼近似下三角矩陣編碼算法[13]與該GPU加速的編碼算法進(jìn)行了對(duì)比。該傳統(tǒng)算法是一種串行的CPU編碼算法，算法復(fù)雜度為O(n)，實(shí)驗(yàn)結(jié)果如表4。

表4 與傳統(tǒng)的編碼算法對(duì)比

從表4可以看出，本文提出的GPU編碼算法對(duì)于3種不同長度的QC-LDPC碼，其編碼吞吐量均超過了10 Gbit/s，傳統(tǒng)編碼算法最高吞吐量為12.1 Mbit/s，遠(yuǎn)不及本文GPU加速的算法。

目前GPU加速的LDPC碼譯碼器很多，但編碼器不多?，F(xiàn)有的LDPC碼快速編碼方案大多是基于FPGA、集成電路的硬件編碼方案，接下來本文將分別討論。

4 同其他GPU加速的編碼器對(duì)比

在文獻(xiàn)[14]中，作者提出了針對(duì)DVB-S2系統(tǒng)中的LDPC碼的GPU編碼器。該LDPC碼的校驗(yàn)矩陣具有圖4中的形態(tài)結(jié)構(gòu)，其中，黑點(diǎn)處為非零元素，空白處為零元素。具有這種校驗(yàn)矩陣結(jié)構(gòu)的DVB-S2 LDPC碼在設(shè)計(jì)GPU編碼器時(shí)可以節(jié)省更多的存儲(chǔ)空間，因而可以同時(shí)對(duì)更多碼字進(jìn)行編碼。

在同樣的Nvidia TITAN Xp顯卡平臺(tái)上，文獻(xiàn)[14]中的編碼器對(duì)碼長16 200，碼率8/9的DVB-S2 LDPC碼一次性可以編8 000個(gè)碼字，編碼吞吐量可以達(dá)到20 Gbit/s；由圖3可知，本文提出的GPU編碼器同時(shí)對(duì)480個(gè)(12 769, 12 320)QC-LDPC碼字進(jìn)行編碼時(shí)，吞吐量為10.174 Gbit/s；最多可同時(shí)編碼1 344個(gè)碼字，此時(shí)的吞吐量約為8.0 Gbit/s。雖然本文提出的GPU編碼器編碼速度沒有文獻(xiàn)[14]快，但本文針對(duì)的是一般QC-LDPC碼的編碼方案，而文獻(xiàn)[14]針對(duì)的是DVB-S2 LDPC碼這種特定碼型LDPC碼的編碼方案。

文獻(xiàn)[8]提出了針對(duì)任意LDPC碼的GPU編碼方案。它假設(shè)LDPC碼生成矩陣為q×s維矩陣，然后直接簡單地分配q個(gè)GPU線程進(jìn)行并行編碼。該方案沒有充分考慮對(duì)單個(gè)碼字編碼時(shí)的并行處理，導(dǎo)致其編碼速度不快。該文獻(xiàn)仿真GPU為Tegra K1，在和傳統(tǒng)的編碼算法[13]進(jìn)行仿真對(duì)比時(shí)，文中給出了不同數(shù)量碼字同時(shí)編碼的運(yùn)行時(shí)間曲線圖，它的編碼器僅比傳統(tǒng)的編碼算法快2～4倍。但從表4可以看出，本文提出的編碼器比傳統(tǒng)算法[13]快900～40 180倍。

5 同硬件快速編碼器比較

在同其他硬件快速編碼器對(duì)比時(shí)，本文選擇了2種較快的硬件編碼器：①基于CMOS技術(shù)的編碼器[10]；②基于FPGA的編碼器[9]。

文獻(xiàn)[10]中提出了一種基于130 nmCMOS技術(shù)實(shí)現(xiàn)的編碼器，是目前針對(duì)802.11ac標(biāo)準(zhǔn)的最快的一種CMOS編碼器。編碼時(shí)，該編碼器主要通過在QC-LDPC碼校驗(yàn)矩陣的列方向上并行化和在電路中引入循環(huán)移位寄存器來提高編碼吞吐量。該文的仿真結(jié)果表明，對(duì)于標(biāo)準(zhǔn)中的(1 944, 1 620)QC-LDPC碼在100 MHz的頻率時(shí)可以達(dá)到最高7.7 Gbit/s的編碼吞吐量。

為了便于對(duì)比，本文同樣選擇了802.11ac標(biāo)準(zhǔn)中的碼字進(jìn)行仿真。802.11ac標(biāo)準(zhǔn)中規(guī)定了一系列具有特殊結(jié)構(gòu)的不同碼長、不同碼率的QC-LDPC碼。本文仿真選擇了該標(biāo)準(zhǔn)中1/2碼率的(648, 324)QC-LDPC碼和(1 944, 972)QC-LDPC碼，以及3/4碼率的(1 944, 1 458)QC-LDPC碼和5/6碼率的(1 944, 1 620)QC-LDPC碼，仿真結(jié)果如表5。

表5 與CMOS編碼器對(duì)比

文獻(xiàn)[10]的CMOS編碼器可以對(duì)802.11ac標(biāo)準(zhǔn)中的所有QC-LDPC碼進(jìn)行編碼，但不能對(duì)其他QC-LDPC碼編碼。文獻(xiàn)[10]對(duì)標(biāo)準(zhǔn)中不同碼率的碼進(jìn)行了仿真，但只給出了編碼最快的(1 944, 1 620)QC-LDPC碼的吞吐量數(shù)據(jù)。本文中的GPU編碼器可以對(duì)任意QC-LDPC碼編碼，本實(shí)驗(yàn)中選擇了802.11ac中3種碼率的QC-LDPC碼進(jìn)行仿真。從表5可以看出，本文的編碼器吞吐量最低為5 Gbit/s，最高為9.6 Gbit/s。且對(duì)于(1 944, 1 620)QC-LDPC碼，本文的編碼器比文獻(xiàn)[10]的CMOS編碼器提高了1.9 Gbit/s的吞吐量，效果明顯。

FPGA是編碼領(lǐng)域中常用的硬件實(shí)現(xiàn)技術(shù)，目前有很多LDPC碼的FPGA編譯碼器方案，但其中速率較快的方案同樣只針對(duì)特定的碼型或碼率的LDPC碼，不具有通用性。本文討論一種2017年提出的FPGA編碼方案[9]，該方案可以對(duì)WIMAX標(biāo)準(zhǔn)(即IEEE 802.16標(biāo)準(zhǔn))中不同碼長、不同碼率的QC-LDPC碼進(jìn)行快速編碼。

本文選擇了WIMAX標(biāo)準(zhǔn)中1/2碼率的(1 536, 768)QC-LDPC碼和(2 304, 1 152)QC-LDPC碼，以及3/4碼率的(1 536, 1 152)QC-LDPC碼和(2 304, 1 728)QC-LDPC碼，文獻(xiàn)[9]仿真平臺(tái)為117 MHz的FPGA平臺(tái)，仿真結(jié)果如表6。

表6 與FPGA編碼器對(duì)比

文獻(xiàn)[9]的FPGA編碼器是目前針對(duì)WIMAX標(biāo)準(zhǔn)中所有QC-LDPC碼通用的最快的FPGA編碼器之一，但從表6的仿真結(jié)果來看，本文提出的GPU編碼器吞吐量是該FPGA編碼器的3.94～7.73倍。而且本文提出的編碼方案可以對(duì)任意QC-LDPC碼進(jìn)行快速編碼，通用性更好。

6 結(jié) 論

本文首先引入QC-LDPC碼具有準(zhǔn)循環(huán)特性的生成矩陣，用這樣的生成矩陣編碼，其編碼過程中許多步驟可以并行化。同時(shí)，GPU具有高速并行運(yùn)算的特點(diǎn)。因此，本文基于GPU加速，給出了一種針對(duì)QC-LDPC碼通用的并行編碼方案。編碼時(shí)，GPU一個(gè)線程塊中多個(gè)線程并行編碼得到一個(gè)碼字，本文方案分配了多個(gè)線程塊同時(shí)進(jìn)行多個(gè)碼字的編碼。此外，在編碼過程中使用了共享內(nèi)存、常量內(nèi)存等高速片上存儲(chǔ)器，縮短了數(shù)據(jù)訪問時(shí)間，提高了編碼吞吐量。從仿真結(jié)果可以看出，目前無論是基于CMOS，F(xiàn)PGA等硬件技術(shù)的QC-LDPC碼編碼方案，還是已有的基于GPU或CPU的軟件編碼方案都有一定的局限性，它們局限于特殊的碼型、碼率等，通用性不好。本文提出的GPU編碼方案不僅可以針對(duì)不同碼型、不同碼率的QC-LDPC碼實(shí)現(xiàn)快速編碼，通用性好，而且編碼速度也有明顯的提升。因此，本文提出的GPU加速編碼器在實(shí)際的高速編碼系統(tǒng)中有很好的應(yīng)用前景。