魯鄒晨

(中國電子科技集團公司第二十研究所，陜西 西安 710068)

0 引 言

信息在信道中傳輸時，難免受到噪聲和衰落的干擾而出錯。隨著Shannon在信道編碼定理[1]中證明采用信道編碼技術(shù)能夠在噪聲干擾的環(huán)境中保持信息的高可靠傳輸，學者們陸續(xù)研究了多種有效的信道編碼設計方法，其中包括了著名的Turbo碼[2]和低密度奇偶校驗碼(LDPC)碼。相較于Turbo碼，LDPC碼分組誤碼性能更優(yōu)，譯碼算法簡單。

隨著半導體工藝的發(fā)展，高速計算研究理論逐漸興起，在圖形處理器(GPU)強大的并行運算架構(gòu)下來實現(xiàn)算法加速具有的優(yōu)越性使其成為可能。LDPC碼的Scaled-MSA譯碼算法中，校驗陣H各行(列)并行處理信息，符合采用全并行架構(gòu)加速的特征。本文采用VS2010標準C編譯器和CUDA6.0的集成開發(fā)環(huán)境，在GPU架構(gòu)上實現(xiàn)了LDPC碼的低時延、高吞吐量譯碼器，并與CPU譯碼以及不同線程數(shù)、并行度間的譯碼速度進行了比較。

1 系統(tǒng)開發(fā)環(huán)境

本文基于CUDA的集成開發(fā)環(huán)境研究并設計了高速率LDPC碼編譯碼器，其中集成系統(tǒng)的開發(fā)環(huán)境需要以下支持：可以運行 CUDA平臺的顯卡和匹配驅(qū)動程序，CUDA工具包和C編譯器。本文GPU采用GeForce 750Ti處理器，CPU采用Intel i3處理器，操作系統(tǒng)為WinXP。整個編譯碼系統(tǒng)采用CPU+GPU混合編譯，GPU開發(fā)環(huán)境由NVIDIA公司開發(fā)的CUDA6.0提供；而CPU開發(fā)環(huán)境采用了VS2010中的標準C編譯器。

表1 CUDA架構(gòu)的硬件環(huán)境

此款顯卡的規(guī)格參數(shù)如表2所示，GPU并行計算加速比很大程度上依賴于流多處理器的數(shù)目和CUDA核心數(shù)。設備的顯存決定了能存儲的資源大小，各Block能夠同時運行的最大線程數(shù)取決于圖像處理器的硬件性能，在研究采用線程數(shù)對譯碼算法的加速能力時，調(diào)用線程數(shù)不能超過這個值。譯碼過程中只讀不寫的常量存儲在常量存儲器中，共享存儲器的存儲容量較小但可以顯著提高數(shù)據(jù)的讀寫速度[3]。

表2 Geforce 750Ti顯卡參數(shù)表

2 并行歸一化MSA譯碼的CUDA實現(xiàn)

2.1 歸一化 MSA譯碼算法

LDPC碼的歸一化最小和譯碼算法，其譯碼步驟可以劃分為以下4項，其中：

M(n)為和第n個變量節(jié)點(VN)連接的檢驗節(jié)點(CN)的集合；

M(n)m為集合M(n)中除去m的子集；

N(m)為第m個校驗方程中的VN節(jié)點集合；

N(m) 為從集合N(m)中去掉n之后的子集；

Lqnm為VN節(jié)點外信息；

LLRn為比特n的信道初始值(對數(shù)似然比)；

Lrmn為VN節(jié)點外信息；

lmax為最大迭代次數(shù)；

LQn為VN節(jié)點n的后驗概率；

(1) 對LDPC碼的校驗矩陣H中的各個非零元素進行初始化。

Lqmn=LLRn=L(xn|yn)=lg(P(xn=0|yn)/

P(xn=1/yn))=2yn/σ2

(1)

式中：0≤m≤M；0≤n≤N。

(2) 對校驗節(jié)點傳送到變量節(jié)點的信息進行更新，k為歸一化因子。

(2)

αn′m=sign(Lqn′m)

(3)

Φ(x)=tanh(x/2)

(4)

(3) 對變量節(jié)點傳送到校驗節(jié)點的信息進行更新。

(5)

(6)

Scaled-MSA譯碼算法中，消息值在CN和VN之間傳遞、每次迭代計算譯碼結(jié)果并根據(jù)校驗方程完成校驗的步驟都彼此獨立。將這些流程映射為幾個獨立CUDA核函數(shù)，再運行到GPU上利用劃分的線程網(wǎng)格就能完成并行加速。編譯碼系統(tǒng)可簡述為：

(1) 在CPU上完成信道初始化；

(2) 將CN節(jié)點的更新映射為設備上的一個Kernel函數(shù)(即CNP核)，為該核在設備上分配一個對應GPU網(wǎng)格Grid1；

(3) 同樣將VN節(jié)點的更新映射為VNP核，再為其分配一個對應網(wǎng)格Grid2；

(4) 嘗試譯碼判決。

GPU并行譯碼包括GPU初始化、線程資源聲明、核函數(shù)定義及運行。為了優(yōu)化編譯碼系統(tǒng)的處理效率，應該減少PCIE總線非必需的信息傳輸，主機CPU和設備GPU間只傳輸初始化后的對數(shù)似然比值和硬判結(jié)果。其它變量由GPU核函數(shù)訪問，無需在CPU和GPU間傳輸，例如Lrmn。

GPU初始化時完成內(nèi)存分配和參數(shù)傳遞。QC-LDPC碼的H矩陣高度結(jié)構(gòu)化，可分為多個Z×Z子陣，有3種類型：全零陣、單位陣和單位移位陣。在GPU的Constant memory中為校驗陣分配內(nèi)存并以一種壓縮形態(tài)存儲，僅用4個字節(jié)實現(xiàn)矩陣元素的存儲，有效節(jié)省了編譯碼系統(tǒng)的存儲資源和訪問效率。前兩字節(jié)分別表示行標值和列標值，第3字節(jié)代表矩陣相對單位陣的移位值，末尾字節(jié)表征當前元素是否為0。

再調(diào)用函數(shù)cudaMemcpy()把信道初始化對數(shù)似然比(LLR)值傳到GPU中：

cudaMemcpy(dev_lratio,lratio,sizeof(double),cudaMemcpyHostToDevice)；

核函數(shù)CNP(VNP)中的線程和線程塊的索引ID不需要初始化，值按自然數(shù)順序遞增。線程網(wǎng)格的Block數(shù)和CN節(jié)點(VN節(jié)點)數(shù)目有關(guān)；線程塊的線程數(shù)和校驗陣的行列重有關(guān)。各線程根據(jù)blockId和threadId計算對應H陣中元素的地址。

調(diào)用GPU的Grid和Block前先進行配置和聲明：

dim3 dimBlock(x1，y1，1)

聲明Block大小為x1×y1， CN節(jié)點處理的CNP中Block的大小設定為校驗矩陣H的行重，滿足各線程并行處理矩陣中的信息。

接下來是內(nèi)核函數(shù)的執(zhí)行：

CNP《(，，，)

《<》>為Kernel函數(shù)執(zhí)行符，尖括號內(nèi)依次為Grid內(nèi)的Block數(shù)目、Block內(nèi)的Thread數(shù)量，()里面包含了CNP函數(shù)的參數(shù)。

LDPC碼的編譯碼耗時和譯碼算法復雜度有較強關(guān)聯(lián)，完成迭代譯碼步驟耗費的時延在系統(tǒng)總耗時中占了較大比重。傳統(tǒng)CPU平臺只能順序、串行譯碼，而配置足夠的GPU線程資源完成的LDPC碼高速并行譯碼，在不改變Scaled-MSA譯碼算法的工作原理的前提下，卻能以協(xié)同處理的方式利用設備豐富的計算資源，并行、高效完成譯碼，能夠帶來譯碼時延的顯著減少。

2.2 仿真模型

圖1 LDPC碼編譯碼的仿真系統(tǒng)模型

2.3 譯碼器的CUDA實現(xiàn)

基于CUDA架構(gòu)實現(xiàn)的譯碼器如圖2所示，共包含以下處理流程：

初始化：在GPU上開辟內(nèi)存并對校驗(變量)節(jié)點賦初值，另開辟內(nèi)存完成校驗矩陣的存儲。H矩陣的元素都是常量，利用constantmemory能夠顯著優(yōu)化訪問時延。信息比特經(jīng)編碼和調(diào)制后進入信道，加噪后的接收值從CPU上傳遞到GPU上。

譯碼過程中， CNP和VNP 2個核函數(shù)采用各自線程網(wǎng)格并行處理：

CNP核：每次迭代時，根據(jù)校驗矩陣H每一行關(guān)聯(lián)的VN節(jié)點對CN節(jié)點更新，各GPU線程選擇相關(guān)VN節(jié)點信息并獨立計算，再將結(jié)果回傳CN節(jié)點。本文研究的碼率為二分之一的(1 024，512)LDPC碼，其H矩陣行重為6，即每個CN節(jié)點將與6個VN節(jié)點連接。那么對于各個CN節(jié)點，都可以分配設備的一個線程塊進行運算處理，這個線程塊內(nèi)應至少包含6個thread。

VNP核： 在CN節(jié)點消息值更新后，校驗陣H每列關(guān)聯(lián)的CN節(jié)點對VN節(jié)點更新，譯碼迭代完成后進行判決。

譯碼結(jié)果回傳：判決后的結(jié)果從設備GPU返回主機CPU，并釋放在設備上開辟的內(nèi)存資源。

圖2 基于CUDA的LDPC并行譯碼器的實現(xiàn)框圖

2.4 并行譯碼的速度分析

統(tǒng)計CPU平臺譯碼耗時使用clock()函數(shù)， GPU平臺統(tǒng)計譯碼耗時能利用CUDA API中的事件管理函數(shù)來測量。

本文測試并行度對譯碼時間影響時分為4種方案。方案一在GPU上執(zhí)行CN節(jié)點的處理，VN的處理在CPU上完成；方案二在CPU上對校驗節(jié)點的更新進行處理， 變量節(jié)點運算在設備上進行；方案三將CN和VN的處理都放到了GPU來并行加速，方案四是原始的CPU串行譯碼方案。仿真使用的(1 024，512)LDPC碼測試總幀數(shù)為10 000幀，采用BPSK的調(diào)制方式， 模擬信道的信噪比設置為3.0 dB，GPU分配的總線程數(shù)為256個。

采用CUDA架構(gòu)并行譯碼方式的幾種方案，雖然譯碼速率加速比的大小不同，但對比CPU平臺的串行譯碼均取得明顯加速，如圖3所示。僅采用GPU線程資源支持CN節(jié)點并行處理或者僅對VN并行處理，加速比提升均不顯著，為2.5倍左右；若采用行列全并行的譯碼方案，則獲得大約7倍的譯碼速度提升。

圖3 不同并行度方案的加速分析

2.5 采用不同線程數(shù)量的加速分析

在表1所示的實驗平臺上，對經(jīng)過模擬的加性高斯白噪聲信道傳輸，采用CUDA并行譯碼的LDPC碼譯碼耗時進行統(tǒng)計。LDPC碼測試幀數(shù)目為10 000幀，采用BPSK調(diào)制方式，模擬噪聲信道的信噪比設置為3.0 dB。

并行譯碼時，采用不同數(shù)量的線程來比較譯碼加速效果的區(qū)別。從圖4不難看出分配的GPU線程數(shù)對譯碼速度有明顯的影響。當分配的設備線程較少時，譯碼速度可能要慢于CPU平臺的串行譯碼，可見調(diào)用的GPU線程不足時，譯碼時帶來的加速十分有限。隨著為LDPC碼譯碼分配的GPU線程資源的增加，CUDA高速率并行譯碼的加速特性逐漸明顯。分配的線程數(shù)目在一定范圍內(nèi)變化時，譯碼速度與調(diào)用GPU線程數(shù)近似成線性增長。

圖4 采用不同線程數(shù)量的加速分析

3 結(jié)束語

結(jié)合本文對編譯碼系統(tǒng)譯碼速度性能決定因素的研究，不難發(fā)現(xiàn)：在設計基于GPU平臺的LDPC碼的高速編譯碼器系統(tǒng)時，合理利用CUDA架構(gòu)下各存儲器的訪存特性，增加譯碼算法中各環(huán)節(jié)對圖像處理器并行運算單元的利用度，盡可能利用更多數(shù)量的GPU線程對碼字并行譯碼，就能最大限度減少編譯碼的訪存時延，進而提升系統(tǒng)吞吐量。

[1] SHANNON C E.A mathematical theory of communication [J].Bell System Technical Journal,1948，27(3):379-423.

[2] BERROU C,GLAVIEUX A,THITIMAJSHIMA P.Near Shannon limit error-correcting coding and decoding:Turbo-Codes[C]//Proceedings of ICC 1993,Geneva,Switzerland,1993:1064-1070.

[3] WANG S,CHENG S,WU Q.A parallel decoding algorithm of LDPC codes using CUDA[C]//Signals,Systems and Computers,2008 42nd Asilomar Conference，2008:172-174.