基于CUDA架構的LDPC碼并行譯碼設計與實現

魯鄒晨

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

0 引 言

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

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

1 系統開發環境

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

表1 CUDA架構的硬件環境

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

表2 Geforce 750Ti顯卡參數表

2 并行歸一化MSA譯碼的CUDA實現

2.1 歸一化 MSA譯碼算法

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

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

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

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

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

Lqnm為VN節點外信息；

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

Lrmn為VN節點外信息；

lmax為最大迭代次數；

LQn為VN節點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) 對校驗節點傳送到變量節點的信息進行更新，k為歸一化因子。

(2)

αn′m=sign(Lqn′m)

(3)

Φ(x)=tanh(x/2)

(4)

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

(5)

(6)

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

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

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

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

(4) 嘗試譯碼判決。

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

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

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

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

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

調用GPU的Grid和Block前先進行配置和聲明：

dim3 dimBlock(x1，y1，1)

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

接下來是內核函數的執行：

CNP《(，，，)

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

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

2.2 仿真模型

圖1 LDPC碼編譯碼的仿真系統模型

2.3 譯碼器的CUDA實現

基于CUDA架構實現的譯碼器如圖2所示，共包含以下處理流程：

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

譯碼過程中， CNP和VNP 2個核函數采用各自線程網格并行處理：

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

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

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

圖2 基于CUDA的LDPC并行譯碼器的實現框圖

2.4 并行譯碼的速度分析

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

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

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

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

2.5 采用不同線程數量的加速分析

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

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

圖4 采用不同線程數量的加速分析

3 結束語

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

[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.