一种基于GPU的5G多用户LDPC码高速译码器及其译码方法

一种基于gpu的5g多用户ldpc码高速译码器及其译码方法
技术领域
1.本发明属于通信技术领域，涉及一种基于graphics processing unit(gpu，图形处理器)的5g上行数据信道多小区多用户ldpc码高速译码器及其译码方法。


背景技术：

2.目前，通信系统主要基于两种平台进行部署：专用硬件平台和通用处理平台。专用硬件可以分为专用集成电路(asic)、现场可编程门阵列电路(fpga)等。使用基于专用硬件的通信系统，可以获得较低延时、较高的能量效率，但其开发周期长、后期运维难度高，难以满足未来通信系统对任务多样化的发展需求。新一代通信系统发展的一个重要趋势是基于通用处理器进行部署。通用处理器主要包括中央处理器(cpu)和图形处理器(gpu)，通用处理器的通信系统主要是通过软件设计进行实现，采用程序层面上的升级，灵活地进行网络部署和参数配置，比如已经提出的新一代通信系统的关键技术：软件定义无线电(sdr)、软件定义网络(sdn)和网络功能虚拟化(nfv)。由于信道译码是通信系统物理层中最为消耗资源的过程之一，因此设计高效实用的软件译码器对于新一代虚拟通信系统的发展具有非常重要的意义。
3.aronov等人分别在fpga和gpu上基于和积算法(spa)实现了面相5g的ldpc码高效译码。结果表明基于fpga的平台更容易针对高速、低延迟处理进行优化，其缺点是通常表现出较高的非重复工程成本，特别是在所需的ip核不可用或缺乏所需性能的情况下，通常需要许多设计迭代来调试；而gpu支持浮点运算，相比于fpga定点计算具有较高计算精度，误块率性能更具优势，并且具有更高的主频和更多的物理资源，能够同时并行处理更多的同种ldpc码块(cb)，吞吐量更大。(参考文献[1]：aronov a,kazakevich l,mack j,et al.5g nr ldpc decoding performance comparison between gpu&fpga platforms[c]//2019ieee long island systems,applications and technology conference(lisat).ieee,2019:1-6.)。li等人对5gldpc码的分层最小和译码算法进行了优化，实现了同种码块在gpu上的高效并行译码。通过压缩基矩阵信息、两级量化、传输数据打包等方式减少了数据传输量，降低了访存延迟。(参考文献[2]：li r,zhou x,pan h,et al.a high-throughput ldpc decoder based on gpus for 5g new radio[c]//2020ieee symposium on computers and communications(iscc).ieee,2020:1-7)。tarver等人基于洪泛最小和算法在gpu上实现了低时延、高吞吐5gnrldpc译码器，并在oai软件栈完成了集成测试，最低时延可达到87us，最高吞吐可达到4gbps。采用三个内核(kernel)函数进行译码，分别为校验节点运算内核、变量节点运算内核和硬判决内核。tarver等人提出宏码字(macro-codeword)技术，通过灵活调整宏码字数量以及其中承载的子码字数量可以达到灵活重配置的效果。通过异步传输流来充分利用gpu调度机制掩盖数据拷贝延迟，减少计算等待时间。(参考文献[3]：c.tarver,m.tonnemacher,h.chen,j.zhang and j.r.cavallaro,"gpu-based,ldpc decoding for 5g and beyond,"in ieee open journal of circuits and systems,vol.2,pp.278-290,2021.)
[0004]
近期，相关专家和研究人员提出的各种基于gpu的5g上行数据信道的ldpc码高速译码器，启动一次gpu都仅仅对应一种码长码率的ldpc译码，不具备kernel内部码率兼容的能力；现有的大多数高吞吐译码器靠增加同种码块间并行度的方式提高译码器的吞吐量，单次译码时间较长。这样的设计方案会带来两大问题：1)上行信道基站中对应多种用户，由于使用场景不同，一个批次中译码器需要处理的码长码率不同。现有的基于gpu的译码器缺少不同码长码率兼容的ldpc译码方法的研究，不能有效地同时处理不同用户的码块，只能通过频繁地切换启动kernel函数调整不同的线程映射机制来完成译码，因此主机与gpu之间频繁的控制、通信、同步等操作将带来极大的延迟开销，不能满足实际延迟和吞吐需求。2)上行信道基站中需要处理大量码块(单个基站最多要处理1000个以上的码块)，但是基站中每个用户需要处理的码块数量不多，通过不断增加码块间并行度提高吞吐量的这种方式有很大局限性，若不能在一次kernel调度过程中同时传入足够多的码块将无法获得高吞吐译码性能。因此实际的多小区多用户场景中少量码块数量的并行计算将大大降低gpu的处理能力和运行效率，并且不利于异步流的灵活分配。


技术实现要素：

[0005]
本发明提出一种基于gpu的5g上行数据信道多小区多用户ldpc码高速译码器及其译码方法，解决现有译码器在并行处理不同码长码长码率ldpc译码兼容问题，进一步降低处理单个码块的时延，提高译码器吞吐量。
[0006]
本发明首先提供一种基于gpu的5g上行数据信道多小区多用户的ldpc码高速译码器，所述译码器结构主要包括主机端和gpu设备端；主机端设置主机内存和cpu芯片，cpu芯片用于预处理和控制调度整个译码过程，控制不同码长码率ldpc码块分组，通过异步流输入gpu设备端；gpu设备端设置gpu芯片，gpu芯片由若干个流多(sm)处理器组成，每个sm处理器负责一组不同码型ldpc码块的译码；主机端和gpu设备端通过高速串行计算机扩展总线(pci-e)进行数据传输。
[0007]
本发明还提供一种基于gpu的5g上行数据信道多小区多用户的ldpc码高速译码方法，整个译码过程可以分为两个阶段：初始化阶段和译码判决阶段。初始化阶段包括步骤1～步骤3，译码判决阶段包括步骤4～步骤8是判决阶段，具体译码流程步骤如下：
[0008]
步骤1：主机端存储空间初始化。
[0009]
在主机端配置基站内待处理的小区数量、每个小区的用户数量、同时处理的最大码块数量，并为接收到的软信息分配存储空间；
[0010]
步骤2：gpu设备初始化。
[0011]
在gpu端配置内存空间，创建异步传输数据流并设置流的数量；在gpu设备端为译码判决后的码字分配全局内存，按照小区用户最大占用资源对gpu分配共享内存空间；
[0012]
步骤3：ldpc基矩阵信息结构体重描述。
[0013]
在主机端，对5g协议规定的所有ldpc基矩阵进行初始化，将基矩阵中的信息进行重新描述为基矩阵信息结构体，在开始译码前，将处理好的全部基矩阵的信息结构体传入gpu常量内存，等待译码使用；
[0014]
步骤4：主机端调度gpu译码。主机端设置好启动kernel函数的线程数量，主机端控制不同码长码率ldpc码块分组，采用8比特定点量化llr传输与16比特浮点运算结合方法；
[0015]
步骤5：llr信息拷贝。
[0016]
每个线程进行大粒度访问全局内存中的llr软信息，每个llr数据的处理过程为：将8比特定点量化的llr软信息转换为16bit浮点信息(half类型)，作为后验概率信息拷贝到共享内存；
[0017]
步骤6：gpu端根据用户码块信息分配相应的线程数量，选择对应的基矩阵信息结构体，基于分层最小和算法进行迭代译码。译码阶段共包括以下9个步骤：
[0018]
步骤6.1：采用码块分组方法进行码块划分。
[0019]
步骤6.2：迭代循环译码。迭代更新后验概率信息。若迭代次数达到最大迭代次数，或者所有检验行都满足校验，迭代过程结束，转到步骤7；否则继续迭代，转到步骤6.3。
[0020]
步骤6.3：开始逐层译码。基于分层最小和译码逐层更新后验概率信息，若译码层数超过该码字的校验层数，则停止层间循环，转到步骤6.2，准备进行下一次迭代；否则继续逐层计算，转到步骤6.4。
[0021]
步骤6.4：计算更新信息位置。根据码字配置，从常量内存读取对应基矩阵信息结构体，高效寻址，更新信息的具体地址，缓存到寄存器中。
[0022]
步骤6.5：缓存待更新的后验概率信息。根据寄存器中的地址信息，到共享内存的对应位置读取后验概率信息，每个后验概率信息为16比特half类型，每个寄存器宽度为32比特，每两个后验概率信息合并存储到一个寄存器中，每个校验行需要ceil(d/2)个寄存器，ceil(*)代表向上取整。
[0023]
步骤6.6：先验信息存储方式选择。
[0024]
存储格式采用压缩先验信息存储和和非压缩先验信息存储结合的方法，先验信息采用部分更新策略，非核心阵校验行只计算前(d-1)个先验信息。若先验信息数量w大于阈值t，采用压缩存储，转到步骤6.7；若w≤t，采用非压缩存储，转到步骤6.8。
[0025]
步骤6.7：计算压缩先验信息。
[0026]
将前一次迭代的压缩存储的先验信息进行解压缩，每行寻址的后验概率信息与前一次的先验信息相减后取绝对值，采用压缩先验信息两两更新方法，每个校验行用两个寄存器分别存储该行的最小值、次小值、最小值所在位置和所有先验信息符号，转到步骤6.9。
[0027]
步骤6.8：计算非压缩先验信息。
[0028]
将每行寻址的后验概率与前一次的非压缩先验信息相减后取绝对值，采用非压缩信息基于box_plus函数更新方法，直接将先验信息存到寄存器内存，每个校验行用ceil((d-1)/2)个寄存器，转到步骤6.9。
[0029]
步骤6.9：后验概率信息更新。
[0030]
后验概率信息加上更新后的先验信息，再将后验概率信息存入共享内存的相应位置，完成一层的更新，检查该层的所有信息符号是否满足校验关系，转到步骤6.3，准备进行下一层的译码。
[0031]
步骤7：符号判决。采用硬判决比特打包方法；
[0032]
步骤8：结果回传到主机端。以异步传输的方式从gpu全局内存通过pci-e传送到主机端内存。
[0033]
可选的，步骤3所述的基矩阵重描述信息结构体包括：基矩阵信息在主机端进行重描述，基矩阵为校验矩阵所有循环块对应的偏移量，重新描述后的信息结构体包括循环快
在每行的位置offset，以及对应循环块的偏移量shift，将全部基本图发信息结构体在译码前全部传入gpu常量内存固定位置；
[0034]
可选的，步骤4所述的主机端控制不同码长码率ldpc码块分组包括：在kernel函数启动的时候分配了g个block，每个block为b个线程，整个译码过程中，线程总数将不会保持变化。对多小区多用户不同码型的ldpc码块进行分组，将前后n个码块进行合并为一组，n取满足约束关系的最大值，其中ceil(*)表示向上取整，b表示每个block中线程总数理，zck表示第k个码块需要的线程数量，1≤k≤n。
[0035]
可选的，步骤4所述的8比特定点量化传输与16比特浮点运算结合的传输方法包括：在主机端接收到解调后的8比特定点量化软信息，通过pci-e传输到gpu全局内存，在gpu端对数据转化为16比特浮点，再进行译码。
[0036]
可选的，步骤5所述的大粒度数据搬运方式包括：每次加载8个全局内存中的8比特定点量化数据到寄存器，将8比特定点数据转换为16比特浮点half类型，将转换后的8个16比特half数据存到共享内存，即访问和存储数据最大粒度均为128比特。
[0037]
可选的，步骤6.1所述的码块分组方法包括：处理不同zc的ldpc码字，block中的b个线程为一维排列，每个线程标识号为threadidx.x，根据不同码块的zc计算每个码块需要分配的线程数量，计算方式为ceil(zc/32)*32，将线程的标识号重新赋值为cnidx，每个码块对应的线程标识号cnidx为[0,ceil(zc/32)*32-1]，译码过程中限制每个码块只有前zc个线程工作，即cnidx《zc；处理不同码率的ldpc码字，5g协议规定的ldpc码字由于码率不同，处理的层数在4层到46层之间不等，每次迭代进行逐层译码，每层结束后需要判断是否为该码字的最后一层，若为最后一层，直接跳出层间循环进行下一次迭代，否则继续下一层处理。
[0038]
可选的，步骤6.4所述的高效寻址包括：在gpu译码过程中根据信息结构体高效计算每个校验行需要更新的llr地址，每组可以根据码字配置信息计算对应的基矩阵，到常量内存固定位置索引基矩阵信息结构体，根据偏移量shift计算出每个线程在该循环块中的地址qc_address＝(shift zc)％zc，再根据offset和qc_address计算出最终的地址address＝offset*zc qc_address，涉及的乘加运算转移到gpu的乘加专用计算单元，利用硬件优势提高运算速度。
[0039]
可选的，步骤6.6所述的压缩先验信息与非压缩先验信息结合存储方法包括：行更新采用两种方式存储先验信息，压缩存储和非压缩存储，通过将更新的先验信息从放到寄存器内存缓存，极大地提高了运行速度。其中压缩存储引入了额外的先验信息压缩和解压缩操作，降低了速度，但是节省了寄存器内存空间。非压缩存储直接将更新的先验信息缓存，提高了速度，增加了寄存器压力。二者结合，通过将该行更新的先验信息数量与设置的阈值t判断比较，先验信息数量较大时，选用压缩存储，先验信息数量较小时选用非压缩存储，合理利用了寄存器提高译码迭代速度。
[0040]
可选的，步骤6.6所述的部分先验信息更新包括：基矩阵中的检验行除去核心阵校验行，其他校验行最后一列构成一个单位阵，利用单位矩阵的结构的特点，在更新过程中省略每行最后一列数据运算。即非核心阵校验行更新的先验信息数量为行重d减1。
[0041]
可选的，步骤6.7所述的压缩先验信息两两更新方法包括：将两个输入信息进行大
小排序，将绝对值大的放在高位，更新记录最小值的位置，记为info1；然后再输入两个信息，记为info2；先用info2的高位与info1的高位比较，选择更小的值作为info1高位；再用info2的高位与info1的低位比较，如果info2的高位更小，将info1高低位进行交换，否则info1不变，进而得到新的info1；再用info2的低位分别与info1进行相同的高位、低位比较操作，得到更新结果info1，从而完成了一轮的两两比较。
[0042]
可选的，步骤6.8非压缩信息基于box_plus函数更新方法包括：yi＝boxplus(xj)＝min(abs(xj))*xor(sign(xj))，i∈q，j∈q\i，其中q为向量所有位置的集合。即第i个位置计算的值yi的绝对值等于除去第i个位置其他所有位置xj的绝对值最小值，yi的符号为其他所有xj的符号的异或运算，0表示正，1表示负。
[0043]
可选的，步骤7所述的硬判决比特打包方法包括两个阶段：第一阶段为数据判决，32个线程为1组，每组中每个线程到共享内存加载一个16比特llr数据，取符号位得到0或者1的数值，0表示正数，1表示负数；第二阶段为比特打包，每个线程将自己持有的符号传递给__ballot_sync(*)投票函数，每个线程均返回一个相同的32比特掩码mask，掩码中的比特1表示对应线程号的输入数据非0，从而得到了比特打包后的结果，再通过任意一个线程，通常为0号线程，将该组线程的结果放到全局内存。
[0044]
与现有技术相比，本发明的优点与积极效果在于：
[0045]
本发明所述的基于gpu的5g上行数据信道多小区多用户高速译码器可以在kernel函数内从软件逻辑层面为每个码块灵活划分相应的线程数量，支持同时对不同码长码率的码块进行译码，相比于传统的仅支持同种码型并行处理的译码器，在处理多小区多用户的过程中，本发明所述译码器避免了内核函数的频繁调度和切换，有效地减小了主机与设备之间的同步开销和内核函数的启动开销，提高了译码器的灵活性和实用价值。
[0046]
本发明充分结合分层译码算法的特点和gpu的架构特点，充分利用gpu片上资源，提升访存效率和数据计算单元利用率，降低单个码块消耗资源的同时降低单个码块的译码时间，提高整体的信息吞吐量，所述的译码方式更适用于处理实际场景中多小区多用户的ldpc码块译码。
附图说明
[0047]
图1为本发明所述的基于gpu的ldpc高速译码器结构示意图。
[0048]
图2为本发明所述的基于gpu的ldpc高速译码器的译码方法流程图。
[0049]
图3为本发明所述的基矩阵信息高效寻址示意图。
[0050]
图4为本发明所述的多用户不同码型码块划分示意图。
[0051]
图5为本发明所述的线程大粒度访存示意图。
[0052]
图6为本发明所述的压缩先验信息与非压缩先验信息存储模型示意图。
[0053]
图7为本发明所述的计算压缩先验信息的示意图。
[0054]
图8为本发明所述的计算非压缩先验信息的更新示意图。
[0055]
图9为本发明所述的硬判决与比特打包示意图。
具体实施方式
[0056]
下面结合附图和实施示例对本发明进行详细说明。
[0057]
本发明首先提供一种基于gpu的5g上行数据信道多小区多用户的ldpc码高速译码器，如图1所示，其中n
sm
表示gpu芯片上流多处理器的个数。所述译码器结构主要包括主机端和gpu设备端(简称gpu端)，主机端设置主机内存和cpu芯片，cpu芯片用于预处理和控制调度整个译码过程。gpu端设置gpu芯片，gpu芯片由若干个流多(sm)处理器组成，每个sm处理器负责一个批次不同码型ldpc码块的译码。sm处理器的逻辑单元包括全局内存、常量内存、共享内存、寄存器内存等。其中，全局内存与主机内存之间通过pci-e总线连接，通信接收端对数似然比软信息、ldpc码块信息以及判决的码字在主机内存和gpu端的全局内存之间以异步传输的方式进行传输；常量内存存储基矩阵的描述信息用于生成译码过程中的寻址信息；共享内存存储译码迭代过程中的后验概率信息，从而降低逐层更新访问存储信息时延；寄存器内存用于存放所有计算访存过程产生的中间量等临时变量。
[0058]
本发明还提供一种基于gpu的5g上行数据信道多小区多用户的ldpc码高速并行译码方法，译码流程如图2所示，整个译码过程可以分为两个阶段：初始化阶段和译码判决阶段。初始化阶段包括步骤1～步骤3，译码判决阶段包括步骤4～步骤8是判决阶段，具体译码流程步骤如下：
[0059]
步骤1：主机端存储空间初始化。
[0060]
在主机端配置基站内待处理的小区数量、每个小区的用户数量、同时处理的最大码块数量，并为接收到的软信息分配存储空间；
[0061]
步骤2：gpu设备初始化。
[0062]
在gpu端配置内存空间，创建异步传输数据流并设置流的数量；在gpu设备端为译码判决后的码字分配全局内存，按照小区用户最大占用资源对gpu分配共享内存空间；
[0063]
步骤3：ldpc基矩阵信息结构体重描述。
[0064]
在主机端，对5g协议规定的所有ldpc基矩阵进行初始化，将基矩阵中的信息进行重新描述为基矩阵信息结构体，在开始译码前，将处理好的全部基矩阵的信息结构体传入gpu常量内存，等待译码使用；
[0065]
步骤4：主机端调度gpu译码。
[0066]
主机端设置好启动kernel函数的线程数量，将多小区多用户的码块进行分组，采用8比特定点量化llr传输与16比特浮点运算相结合的方式，通过pci-e总线将llr和码块信息异步流依次传入gpu设备；
[0067]
步骤5：llr信息拷贝。
[0068]
每个线程进行大粒度访问全局内存中的llr软信息，每个llr数据的处理过程为：将8比特定点量化的llr软信息转换为16bit浮点信息(half类型)，作为后验概率信息拷贝到共享内存；
[0069]
步骤6：gpu端根据用户码块信息分配相应的线程数量，选择对应的基矩阵信息结构体，基于分层最小和算法进行迭代译码。译码阶段共包括以下9个步骤：
[0070]
步骤6.1：码块分组。kernel函数在逻辑层面根据每个ldpc码块的提升值zc为每个码块分配ceil(zc/32)*32个线程，ceil(*)代表向上取整，同时对不同码型的ldpc码字进行译码。
[0071]
步骤6.2：迭代循环译码。迭代更新后验概率信息。若迭代次数达到最大迭代次数，或者所有检验行都满足校验，迭代过程结束，转到步骤7；否则继续迭代，转到步骤6.3。
[0072]
步骤6.3：开始逐层译码。基于分层最小和译码逐层更新后验概率信息，若译码层数超过该码字的校验层数，则停止层间循环，转到步骤6.2，准备进行下一次迭代；否则继续逐层计算，转到步骤6.4。
[0073]
步骤6.4：计算更新信息位置。根据码字配置，从常量内存读取对应基矩阵信息结构体，更新信息的具体地址，缓存到寄存器中。
[0074]
步骤6.5：缓存待更新的后验概率信息。根据寄存器中的地址信息，到共享内存的对应位置读取后验概率信息，每个后验概率信息为16比特half类型，每个寄存器宽度为32比特，每两个后验概率信息合并存储到一个寄存器中，每个校验行需要ceil(d/2)个寄存器，ceil(*)代表向上取整。
[0075]
步骤6.6：先验信息存储方式选择。
[0076]
先验信息存储到寄存器，首次迭代初始值为0，存储格式分为压缩存储和和非压缩存储，先验信息采用部分更新策略，非核心阵校验行只计算前(d-1)个先验信息。若先验信息数量w大于阈值t，采用压缩存储，转到步骤6.7；若w≤t，采用非压缩存储，转到步骤6.8。
[0077]
步骤6.7：计算压缩先验信息。
[0078]
将前一次迭代的压缩存储的先验信息进行解压缩，每行寻址的后验概率信息与前一次的先验信息相减后取绝对值，再从中两两比较找到所有信息中的最小值和次小值，用常数α进行归一化处理，每个校验行用两个寄存器分别存储该行的最小值、次小值、最小值所在位置和所有先验信息符号，转到步骤6.9。
[0079]
步骤6.8：计算非压缩先验信息。
[0080]
将每行寻址的后验概率与前一次的非压缩先验信息相减后取绝对值，通过boxplus函数分别计算出每个位置的先验信息，用常数α进行归一化处理，直接将先验信息存到寄存器内存，每个校验行用ceil((d-1)/2)个寄存器，转到步骤6.9。
[0081]
步骤6.9：后验概率信息更新。
[0082]
后验概率信息加上更新后的先验信息，再将后验概率信息存入共享内存的相应位置，完成一层的更新，检查该层的所有信息符号是否满足校验关系，转到步骤6.3，准备进行下一层的译码。
[0083]
步骤7：符号判决。当译码迭代次数达到设定的最大值时，利用译码迭代得到的更新后的软信息进行判决，判决后的比特每32个进行一次打包，存到全局内存。
[0084]
步骤8：结果回传到主机端。以异步传输的方式从gpu全局内存通过pci-e传送到主机端内存。
[0085]
其中，步骤3所述的基矩阵信息结构体重描述具体操作为：将基矩阵中记录的循环块偏移量重新描述为信息结构体，每个信息结构体包含2个描述符，分别为循环块在每行的位置offset(占用16比特)，以及对应循环块的偏移量shift(占用16比特)，为了节省存储空间，每两个信息节点中offset/shift的值合并在一个32比特空间。将信息结构体在译码前全部传入gpu常量内存固定位置，提高gpu译码索引信息的效率。
[0086]
其中，步骤4～步骤7为译码阶段，将该阶段合并到一个内核函数去执行，减少线程块间的同步开销、内核函数启动开销以及全局内存的访问量，后验概率信息放入共享内存，减少全局内存的访问，提高译码迭代速度。
[0087]
其中，步骤4中的调度gpu译码具体指：在kernel函数启动时分配g个block，每个
block为b个线程，整个译码过程中，线程总数将不会保持变化。对多小区多用户不同码型的ldpc码块进行分组，将前后n个码块进行合并为一组，n取满足约束关系的最大值，其中ceil(*)表示向上取整，b表示每个block中线程总数理，zck表示第k个码块需要的线程数量，1≤k≤n。
[0088]
其中，步骤4中所述的8比特定点量化llr传输与16比特浮点运算相结合具体为：主机端与gpu端通过pci-e传输信息，相比于gpu强大的计算吞吐能力，pci-e的传输存在较大瓶颈，数据需要花费大量延时在主机和设备之间传输，本发明采用8比特定点量化llr传输，节省传输时延；现有基于fpga的ldpc译码器在进行分层最小和译码时，常常采用8比特定点量化计算中间值，能够满足基本的性能需求，但对于高阶调制后的ldpc译码，llr幅度的范围较大，8比特定点量化计算中间值性能损失较大，并且gpu大部分运算资源为浮点运算，浮点计算单元相比于定点计算单元具有更高的吞吐量，32比特浮点计算对于ldpc码的计算又存在较大计算冗余，性能提升不明显，因此本发明采用16比特浮点计算中间值，最大10次迭代，能够得到与32比特浮点运算几乎相同的误码性能。
[0089]
其中，步骤5所述的大粒度访存具体是指：在迭代译码前，需要将全局内存中的软信息加载到寄存器内存，进行8比特定点到16比特浮点的转换，再存入共享内存。为了提高gpu的访存效率，每个线程每次搬运多个数据，达到在相同的指令操作数量下可以搬运更多数据的效果，整体上减少了循环次数，加快gpu的搬运过程。图5展示了本发明所述的线程大粒度访存示意图，每个元素经过的操作流程为：读取全局内存中的8比特定点数据到寄存器，将8比特定点数据转换为16比特浮点half类型，将转换后的half数据存到共享内存。为了最大限度的加速，需要增加每个线程的访存粒度，访存粒度越大，访问效率越高，在相同的指令周期中可以搬运更多的数据。现有gpu访存粒度最大为128比特，即访问和存储数据最大粒度均为128比特。但是由于存在数据转换，若加载128比特llr数据，转换后将变为256比特，造成共享内存存储时产生访存冲突，因此采用的最大访问粒度为64比特，在寄存器中将其转换扩展，最后按照128比特存入共享内存。图5中每个线程需要在全局内存加载8个llr数据，每个llr占8bit，再将其转换为half类型，每个llr占16比特。每个线程最终形成128比特的数据，再将其存入共享内存。
[0090]
其中，步骤6所述的译码阶段采用的分层最小和译码具有层内校验行并行的特点，即属于同一层的zc个校验行可以采用zc个线程并行更新，每个线程对应一个校验行，不同层的信息进行顺序更新。因此基于分层最小和算法处理一个码块需要zc个线程。图4展示了本发明提出的多用户不同码型码块划分示例，其中每个block开启的线程数量为384，某个block处理了四种码字，分别为基本图1、基本图1、基本图1、基本图2，zc分别为64、80、56、208，每个码块逻辑上分别分配ceil(zc/32)*32个线程，每个线程对应一个校验行，一层内的校验行大小等于提升值zc，cnidx指示了每个码块内部校验行的索引，为[0,63]∪[0,95]∪[0,63]∪[0,223]，译码时每个码块对应的前zc个线程工作。
[0091]
其中，步骤6.4所述的计算更新信息位置具体操作为：在gpu译码过程中根据信息结构体高效计算每个校验行需要更新llr的地址，每组可以根据码字配置信息计算对应的基矩阵，到常量内存固定位置索引基矩阵信息结构体，根据偏移量shift计算出每个线程在该块中的地址qc_address＝(shift threadidx.x)％zc，再根据offset和qc_address计算
出最终的地址address＝offset*zc qc_address，涉及的乘加运算转移到gpu的乘加专用运算单元，进行高速计算。
[0092]
图3展示了本发明提出的基矩阵信息高效计算更新位置示意图，其中展示了zc为8，shift为3，offset为2的示例。先根据偏移量shift计算出每个线程在该块中的地址qc_address＝(3 threadidx.x)％8，0～4号线程对应索引qc_address为3～7，5～7号线程对应索引qc_address为0～2，再根据offset＝2和qc_address计算出最终的地址address＝offset*zc qc_address＝2*8 qc_address。
[0093]
其中，步骤6.6所述的行更新采用两种方式存储先验信息，压缩存储和非压缩存储，通过将更新的先验信息从放到寄存器内存缓存，极大地提高了运行速度。其中压缩存储引入了额外的先验信息压缩和解压缩操作，降低了速度，但是节省了寄存器内存空间。非压缩存储直接将更新的先验信息缓存，提高了速度，增加了寄存器压力。二者结合，通过将该行更新的先验信息数量与设置的阈值t判断比较，先验信息数量较大时，选用压缩存储，先验信息数量较小时选用非压缩存储，合理利用了寄存器提高译码迭代速度。
[0094]
图6展示了本发明所述的压缩传递信息与非压缩传递信息存储模型示例，其中阈值t＝4，因此当先验信息数量w》t时，采用压缩存储方式，min1和min0分别表示每行传递信息的绝对值次小值和绝对值最小值，signs表示每行传递信息的符号，index表示绝对值最小值位置。由于llr均为16bit的half类型，只需要两个32比特寄存器c_c2v[0]和c_c2v[1]可以对min1、min0、signs和index进行缓存。具体来说，由于所有基本图矩阵中的最大行重(度)d为19，因此index占用5比特即可以表示19个数中的任意一个数字，并且signs最多需要19比特，两者共同存于一个32比特的寄存器空间，因此第一个寄存器存储min1和min0，第二个寄存器存储signs和index，压缩存储引入了额外的操作，但是最大程度减少了度很大的情况下的寄存器开销。当w≤t＝4时无需进行压缩，非压缩寄存器数组n_c2v共有ceil(t/2)＝2个元素，每个元素可以存储两个half数据，因此t＝4时数组n_c2v最多存储4个先验信息数据。t的值的可以自定义传入，在寄存器够用的情况下，t值越大，非压缩存储的校验行所占比例越大，处理速度越快，寄存器压力也会越大，应根据gpu的具体情况调整。
[0095]
其中，步骤6.6所述的部分更新策略具体解释为：5g数据信道ldpc码字采用了“raptor-like”结构，其基矩阵可以通过一个高码率的核心矩阵逐步扩展到低码率，基矩阵中的检验行除去核心阵的校验行，其他校验行最后一列构成一个单位阵，利用单位矩阵的结构的特点可以采用部分更新策略简化分层最小和算法的更新过程。分层最小译码在每次迭代过程中需要逐层更新，执行的操作分别为：初始化(见式(1))、变量节点传递信息更新(见式(2))、校验节点传递信息更新(见式(3))、后验概率更新(见式(4))。
[0096]
l
(l)
(v
i,j
)＝llr
in
，l＝0，
ꢀꢀꢀꢀꢀꢀꢀꢀ
(1)
[0097][0098][0099][0100]
其中，n(j)表示第j个校验节点相连的所有变量节点的位置的集合，n(j)\i表示在n(j)中除去第i个变量节点位置组成的集合；k
t
表示第t层的校验节点位置索引的集合；
llr
in
表示信道输入的对数似然比值(llr)；l
(l)
(vi)表示在第l次迭代第i个变量节点的llr，分别表示在第l次迭代、第l 1次迭代时，第i个变量节点的后验概率的值；l
(l-1)
(u
j,i
)、l
(l)
(u
j,i
)分别表示在第l-1次迭代、第l次迭代时，与第i个变量节点相连的第j个校验节点的llr的值；l
(l)
(v
i,j
)表示在第l次迭代时，与第j个校验节点相连的第i个变量节点的llr的值；l
(l)
(vi′
,j
)，i’∈n(j)\i，表示取第l次迭代，第j个校验节点相连的去掉第i个变量节点的llr的值；α表示归一化因子。校验层为单位矩阵，层间更新信息互不干扰，即第t层与第t 1层第i个变量节点的后验概率信息没有变化，因此在下一次迭代变量节点传递信息更新时，l
(l)
(v
i,j
)将与l
(l-1)
(v
i,j
)一样，单位矩阵的变量节点传递的值在迭代更新过程中无变化。因此在gpu译码过程中，对于非核心阵的最后一列数据不进行更新，对译码性能没有影响，减少了处理步骤，降低了存储空间和译码延迟。
[0101]
其中，步骤6.7所述的压缩存储先验信息具体操作为：由于分层最小和算法先验信息存储的数据有很大冗余性，对于度较大的行可以进行压缩存储。分层最小和算法中的先验信息绝对值只存在两种可能取值，即每行计算产生的最小值和次小值，由于数据均为16比特的half类型，因此每行的先验信息可以存储为采用两个32比特的寄存器，第一个寄存器存储最小值和次小值的绝对值，第二个寄存器存储每行的符号以及最小值的位置，由于所有基本图中的最大度为19，因此采用最多5比特表示19个数中的任意一个数字，并且采用最多19比特表示所有符号，因此前后最多需要两个寄存器将整个检验行的先验信息进行缓存。压缩存储引入了额外的操作，但是最大程度减少了行重d很大的情况下的寄存器开销。
[0102]
其中，步骤6.7所述的两两比较具体操作为：将两个输入信息进行大小排序，将绝对值大的放在高位，更新记录最小值的位置，记为input[0]；然后再输入两个信息，记为input[1]；先用input[1]的高位与input[0]的高位比较，选择更小的值作为高位；再用input[1]的高位与input[0]的低位比较，如果input[1]的高位更小，将input[0]高低位进行交换，否则不变，进而得到新的input[0]；再用input[1]的低位分别与input[0]进行相同的高位、低位比较操作，得到更新结果，从而完成了一轮的两两比较。
[0103]
图7展示了计算压缩先验信息的两两更新示意图，以行重d＝5，待计算的先验信息个数w＝4为例。先初始化更新input[0]，将两个数据进行排序，将绝对值大的放在高位；然后将input[1]的高位与input[0]高位比较，将更小的值放入input[0]高位；再将input[1]高位与input[0]中低位比较，如果input[1]高位更小，将input[0]高位低位互换，完成了input[1]高位的比较；input[1]低位重复上述过程，完成input[1]与input[0]比较，最小值和次小值分别记录在input[0]的低位和高位；再将input[2]与input[0]按照上述方式比较，最终将input[0]计算结果赋值给output，其中input[2]高位为无效位置，设置最大值。
[0104]
其中，步骤6.8所述的非压缩存储先验信息具体操作为：对于度较小的行，直接用寄存器数组存储先验信息，寄存器数组每个元素存储两个half数据，存储t个half数据，数组大小为ceil(t/32)。t的值作为选择压缩存储和非压缩存储的判断阈值，其大小可以自定义传入，在寄存器够用的情况下，t的值越大，处理速度越快，寄存器压力也会越大，根据gpu的具体情况调整。
[0105]
其中，步骤6.8所述的boxplus具体操作为：yi＝boxplus(xj)＝min(abs(xj))*xor(sign(xj))，i∈q，j∈q\i，其中q为向量所有位置的集合。即第i个位置计算的值yi的绝对值
等于除去第i个位置其他所有位置xj的绝对值最小值，yi的符号为其他所有xj的符号的异或运算，0表示正，1表示负。
[0106]
图8展示了计算非压缩先验信息的示意图，在进行比较大小的过程中，运用box_plus函数进行运算，以行重d＝5，待计算的先验信息个数w＝4为例。input[0]、input[1]和input[2]所有值共同参与运算，[2]所有值共同参与运算，以此类推。
[0107]
其中，步骤7所述的硬判决与比特打包具体操作如图9所示，图9展示了对于32个迭代更新后llr数据进行判决、比特打包的示例，其中，tn表示第n号线程，第一阶段为数据判决，每个线程到共享内存加载一个16比特llr数据，记为fn，取符号位得到0或者1的数值，记为sn，0表示正数，1表示负数；第二阶段为比特打包，每个线程将自己持有的符号传递给__ballot_sync(*)函数，每个线程均返回一个相同的32比特掩码m，掩码中的比特1表示对应线程号的输入数据非0，从而得到了比特打包后的结果，再通过任意一个线程，通常为0号线程，将该组线程的结果放到全局内存。