更全的杂志信息网

基于CUDA架构的LDPC码并行译码设计与实现

更新时间:2009-03-28

0 引 言

信息在信道中传输时,难免受到噪声和衰落的干扰而出错。随着Shannon在信道编码定理[1]中证明采用信道编码技术能够在噪声干扰的环境中保持信息的高可靠传输,学者们陆续研究了多种有效的信道编码设计方法,其中包括了著名的Turbo码[2]和低密度奇偶校验码(LDPC)码。相较于Turbo码,LDPC码分组误码性能更优,译码算法简单。

通过协商,甲和乙的策略选择受对方提供的新信息所影响,在新的博弈情境中,甲对于策略集合A中的每一个行动选择,都存在一个新的相对于策略集B的条件概率q;同理,乙对于策略集合B中的每一个行动选择,都存在一个新的相对于策略集A的条件概率r。据此,甲和乙的期望效用演变为条件期望效用。对于彼此独立的环境或事态,主体之间有非条件概率。在纳什均衡中,甲和乙的主观概率没有被任何实质性的条件限制;在纯粹策略中,两者按照给定的占优策略行动;在混合策略中,彼此行动的概率选择相互保密,不为对方所知。然而,通过理性协商,主体信念和行动的概率选择成为公共知识,这就决定了协商机制下的行动博弈超越于纳什均衡。

随着半导体工艺的发展,高速计算研究理论逐渐兴起,在图形处理器(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架构的硬件环境

  

平台名称件环境型号平台核心内存存储容量CPUIntel i344 GBGPUGeForce 750Ti6402 GB

(1) 对LDPC码的校验矩阵H中的各个非零元素进行初始化。

 

2 Geforce 750Ti显卡参数表

  

显存大小2 048 000 bytes流多处理器数目5CUDA核心数640常量存储器大小 65 536 bytes线程块中的共享存储器大小49 152 bytes线程块中的最大线程数1 024

2 并行归一化MSA译码的CUDA实现

2.1 归一化 MSA译码算法

LDPC码的归一化最小和译码算法,其译码步骤可以划分为以下4项,其中:

CNP《(,,,)

M(n)m为集合M(n)中除去m的子集;

N(m)为第m个校验方程中的VN节点集合;

N(m)
为从集合N(m)中去掉n之后的子集;

调用GPU的Grid和Block前先进行配置和声明:

LLRn为比特n的信道初始值(对数似然比);

Lrmn为VN节点外信息;

lmax为最大迭代次数;

(2) 将CN节点的更新映射为设备上的一个Kernel函数(即CNP核),为该核在设备上分配一个对应GPU网格Grid1;

此款显卡的规格参数如表2所示,GPU并行计算加速比很大程度上依赖于流多处理器的数目和CUDA核心数。设备的显存决定了能存储的资源大小,各Block能够同时运行的最大线程数取决于图像处理器的硬件性能,在研究采用线程数对译码算法的加速能力时,调用线程数不能超过这个值。译码过程中只读不写的常量存储在常量存储器中,共享存储器的存储容量较小但可以显著提高数据的读写速度[3]

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

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

(1)

式中:0≤mM;0≤nN

(2) 对校验节点传送到变量节点的信息进行更新,k为归一化因子。

顺天者生逆天死,贼尸未僵魂已禠。当风拉杂扬其灰,到头倔强将何为?告尔促浸,尔果悔罪,何不早生致?天兵已临,安敢贳!不见梭磨土妇贤且智,断割私姻名大义,子戴翎枝母受赐。[3]65

 

(2)

αnm=sign(Lqnm)

(3)

Φ(x)=tanh(x/2)

(4)

(3) 对变量节点传送到校验节点的信息进行更新。

 

(5)

 

(6)

(4) 尝试译码。如果LQn≥0,则否则为1。对于码字,如果则停止当前迭代,将x作为译码结果而输出,否则进入第2步继续执行。如迭代次数等于lmax时译码尚未成功,则宣告译码失败,停止迭代。

Scaled-MSA译码算法中,消息值在CN和VN之间传递、每次迭代计算译码结果并根据校验方程完成校验的步骤都彼此独立。将这些流程映射为几个独立CUDA核函数,再运行到GPU上利用划分的线程网格就能完成并行加速。编译码系统可简述为:

首钢水厂铁选厂采用磁选—磨矿—磁选工艺处理TFe品位10%左右的尾砂,每年处理尾矿787万t,生产TFe品位为66.95%的铁精矿28.8万t,回收金属量19.28万t,直接经济价值达2.3亿元;年减少尾矿排放量28.8万t,减少占用尾矿库库容9万m3左右,既产生了经济效益,又缓解尾矿排放对环境的影响[28]。

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

LQn为VN节点n的后验概率;

随着科技的发达,纸质文本向电子书的发展是一种必然,手机阅读相对纸本阅读更加方便,于全民阅读的普及来说,是很有益处的。还有各类专门的电子书阅读器的出现,譬如亚马逊的kindle阅读器,因其携带方便,也日益被人们喜欢和接受。

(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阵中元素的地址。

Lqnm为VN节点外信息;

dim3 dimBlock(x1y1,1)

声明Block大小为x1 × y1, CN节点处理的CNP中Block的大小设定为校验矩阵H的行重,满足各线程并行处理矩阵中的信息。

接下来是内核函数的执行:

M(n)为和第n个变量节点(VN)连接的检验节点(CN)的集合;

因此,首先无论作为反补贴调查国或是被调查国,当出现相关问题的首例案件时,调查机构和立法机构应及时完善相关的法律体系。对缺乏解释的法规进行修订及补充,做到有法可依,防止一个单一的问题因为前期缺乏重视而不断演化,在后期变成复杂棘手的问题。并且,完善相关的监督机制,防止自由裁量权无限制的扩大,维护司法公平。同时针对不同国家不同企业,或者同一国家不同企业的情况,应该做到实事求是。立足于不同国家的市场和企业情况进行分析,防止一概而论而造成的失实判断。作为被调查国,当调查结论严重失实时,应及时与调查机构进行沟通,并依据现实情况上报情况说明,以维护合法权利。

《<》>为Kernel函数执行符,尖括号内依次为Grid内的Block数目、Block内的Thread数量,()里面包含了CNP函数的参数。

LDPC码的编译码耗时和译码算法复杂度有较强关联,完成迭代译码步骤耗费的时延在系统总耗时中占了较大比重。传统CPU平台只能顺序、串行译码,而配置足够的GPU线程资源完成的LDPC码高速并行译码,在不改变Scaled-MSA译码算法的工作原理的前提下,却能以协同处理的方式利用设备丰富的计算资源,并行、高效完成译码,能够带来译码时延的显著减少。

(2)进行电气设备配管下料时,必须要由专业的监理工程师进行施工监督,且施工的整个过程,都需要按照施工规范与设计规范进行相应施工,施工时若发现下料的材料不符合质量要求,必须进行更换,以确保电气设备的质量与安全性。

(1)矿石矿物成分:原生矿物有黄铁矿、黄铜矿、磁铁矿、辉钼矿等,次生矿物主要有孔雀石、褐铁矿、钼华等;脉石矿物主要为斜长石、钾长石、石英、少量黑云母等。

①侵袭组患者的CT表现与病理学表现比较:统计CT诊断21例侵袭性胸腺瘤患者心血管结构侵犯、心膈角和胸腔侵犯、肺侵犯、胸膜种植、纵隔胸膜侵犯和纵隔肿块的灵敏度、特异度、阳性预测值和阴性预测值。②侵袭组和非侵袭组患者的CT表现比较:对比2组患者的胸腺瘤形态、钙化、囊性变或中心坏死、纵隔脂肪线、均匀增强程度。

2.2 仿真模型

基于混合编译环境的LDPC码编译码系统如图1所示,方框表示CPU平台完成的流程,圆框表示GPU平台上的流程。首先在CPU上生成信息比特u,编码后生成码字c∈{0,1},经过二进制相移键控(BPSK)调制后(x=2c-1)得到比特x∈{-1,1}并进入一定信噪比的加性高斯白噪声信道传输。信道传输完成后,接收端收到干扰信号y=x+n,进入设备GPU进行译码,得到译码结果并传回主机CPU。

  

图1 LDPC码编译码的仿真系统模型

2.3 译码器的CUDA实现

基于CUDA架构实现的译码器如图2所示,共包含以下处理流程:

初始化:在GPU上开辟内存并对校验(变量)节点赋初值,另开辟内存完成校验矩阵的存储。H矩阵的元素都是常量,利用constantmemory能够显著优化访问时延。信息比特经编码和调制后进入信道,加噪后的接收值从CPU上传递到GPU上。

译码过程中, CNP和VNP 2个核函数采用各自线程网格并行处理:

3.加速产业结构调整的步伐。加速深圳产业结构的调整步伐,增加商业性融资,促进深圳金融企业发展。目前深圳金融企业的重要资金来源是商业性融资,政府应该继续简政放权,大力推动供给侧改革,以市场为主导,提高资源配置效率,使资源从一些僵尸企业和产能过剩产业流向新产业、新技术等,培育新动力,发展新产业,从而加速深圳产业结构的调整步伐。

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倍的译码速度提升。

一是大幅增加公共财政对水利的投入。2013年财政预算安排水利投入的增幅超过同期财政经常性收入增幅,市县水利投入达134亿元,增长20%以上;切实加大中央投资水利项目的地方配套资金落实力度,省级配套资金比2012年增长了24%。

  

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

 
鲁邹晨
《舰船电子对抗》 2018年第02期
《舰船电子对抗》2018年第02期文献

服务严谨可靠 7×14小时在线支持 支持宝特邀商家 不满意退款

本站非杂志社官网,上千家国家级期刊、省级期刊、北大核心、南大核心、专业的职称论文发表网站。
职称论文发表、杂志论文发表、期刊征稿、期刊投稿,论文发表指导正规机构。是您首选最可靠,最快速的期刊论文发表网站。
免责声明:本网站部分资源、信息来源于网络,完全免费共享,仅供学习和研究使用,版权和著作权归原作者所有
如有不愿意被转载的情况,请通知我们删除已转载的信息 粤ICP备2023046998号