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

技术特征：
1.一种基于gpu的5g多用户ldpc码高速译码器，其特征在于：所述译码器结构包括主机端和gpu设备端；主机端设置主机内存和cpu芯片，cpu芯片用于预处理和控制调度整个译码过程，控制不同码长码率ldpc码块分组，通过异步流输入gpu设备端；gpu设备端设置gpu芯片，gpu芯片由若干个流多sm处理器组成，每个sm处理器负责一组不同码型ldpc码块的译码；主机端和gpu设备端通过高速串行计算机扩展总线pci-e进行数据传输。2.一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：具体步骤如下：步骤1：主机端存储空间初始化；在主机端配置基站内待处理的小区数量、每个小区的用户数量、同时处理的最大码块数量，并为接收到的软信息分配存储空间；步骤2：gpu设备初始化；在gpu端配置内存空间，创建异步传输数据流并设置流的数量；在gpu设备端为译码判决后的码字分配全局内存，按照小区用户最大占用资源对gpu分配共享内存空间；步骤3：ldpc基矩阵信息结构体重描述；在主机端，对5g协议规定的所有ldpc基矩阵进行初始化，将基矩阵中的信息进行重新描述为基矩阵信息结构体，在开始译码前，将处理好的全部基矩阵的信息结构体传入gpu常量内存，等待译码使用；步骤4：主机端调度gpu译码；主机端设置好启动kernel函数的线程数量，主机端控制不同码长码率ldpc码块分组，采用8比特定点量化llr传输与16比特浮点运算结合方法；步骤5：llr信息拷贝；每个线程进行大粒度访问全局内存中的llr软信息，每个llr数据的处理过程为：将8比特定点量化的llr软信息转换为16bit浮点信息，作为后验概率信息拷贝到共享内存；步骤6：gpu端根据用户码块信息分配相应的线程数量，选择对应的基矩阵信息结构体，基于分层最小和算法进行迭代译码；其中，译码阶段还包括以下几个步骤：步骤6.1：采用码块分组方法进行码块划分；步骤6.2：迭代循环译码；迭代更新后验概率信息；若迭代次数达到最大迭代次数，或者所有检验行都满足校验，迭代过程结束，转到步骤7；否则继续迭代，转到步骤6.3；步骤6.3：开始逐层译码；基于分层最小和译码逐层更新后验概率信息，若译码层数超过该码字的校验层数，则停止层间循环，转到步骤6.2，准备进行下一次迭代；否则继续逐层计算，转到步骤6.4；步骤6.4：计算更新信息位置；根据码字配置，从常量内存读取对应基矩阵信息结构体，高效寻址，更新信息的具体地址，缓存到寄存器中；步骤6.5：缓存待更新的后验概率信息；根据寄存器中的地址信息，到共享内存的对应位置读取后验概率信息，每个后验概率信息为16比特half类型，每个寄存器宽度为32比特，每两个后验概率信息合并存储到一个寄存器中，每个校验行需要ceil(d/2)个寄存器，ceil(*)代表向上取整；步骤6.6：先验信息存储方式选择；存储格式采用压缩先验信息存储和非压缩先验信息存储结合的方法，先验信息采用部分更新策略，非核心阵校验行只计算前d-1个先验信息；若先验信息数量w大于阈值t，采用压缩存储，转到步骤6.7；若w≤t，采用非压缩存储，转到步骤6.8；
步骤6.7：计算压缩先验信息；将前一次迭代的压缩存储的先验信息进行解压缩，每行寻址的后验概率信息与前一次的先验信息相减后取绝对值，采用压缩先验信息两两更新方法，每个校验行用两个寄存器分别存储该行的最小值、次小值、最小值所在位置和所有先验信息符号，转到步骤6.9；步骤6.8：计算非压缩先验信息；将每行寻址的后验概率与前一次的非压缩先验信息相减后取绝对值，采用非压缩信息基于box_plus函数更新方法，直接将先验信息存到寄存器内存，每个校验行用ceil((d-1)/2)个寄存器，转到步骤6.9；步骤6.9：后验概率信息更新；后验概率信息加上更新后的先验信息，再将后验概率信息存入共享内存的相应位置，完成一层的更新，检查该层的所有信息符号是否满足校验关系，转到步骤6.3，准备进行下一层的译码；步骤7：符号判决；采用硬判决比特打包方法；步骤8：结果回传到主机端；以异步传输的方式从gpu全局内存通过pci-e传送到主机端内存。3.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤3所述的基矩阵重描述信息结构体包括：基矩阵信息在主机端进行重描述，基矩阵为校验矩阵所有循环块对应的偏移量，重新描述后的信息结构体包括循环块在每行的位置offset，以及对应循环块的偏移量shift，将全部基本图发信息结构体在译码前全部传入gpu常量内存固定位置。4.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤4所述的主机端控制不同码长码率ldpc码块分组包括：在kernel函数启动的时候分配了g个block，每个block为b个线程，整个译码过程中，线程总数将不会保持变化；对多小区多用户不同码型的ldpc码块进行分组，将前后n个码块进行合并为一组，n取满足约束关系的最大值，其中ceil(*)表示向上取整，b表示每个block中线程总数理，zc
k
表示第k个码块需要的线程数量，1≤k≤n；所述的8比特定点量化传输与16比特浮点运算结合的传输方法包括：在主机端接收到解调后的8比特定点量化软信息，通过pci-e传输到gpu全局内存，在gpu端对数据转化为16比特浮点，再进行译码。5.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤5所述的大粒度搬运方式包括：每次加载8个全局内存中的8比特定点量化数据到寄存器，将8比特定点数据转换为16比特浮点half类型，将转换后的8个16比特half数据存到共享内存，即访问和存储数据最大粒度均为128比特。6.根据权利要求3所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤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层之间不等，每次迭代进行逐层译码，每层结束后需要判断是否为该码字的最后一层，若为最后一层，直接跳出层间循环进行下一次迭代，否则继续下一层处理。7.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤6.4所述的高效寻址包括：在gpu译码过程中根据信息结构体高效计算每个校验行需要更新的llr地址，每组根据码字配置信息计算对应的基矩阵，到常量内存固定位置索引基矩阵信息结构体，根据偏移量shift计算出每个线程在循环块中的地址qc_address＝(shift zc)％zc，再根据offset和qc_address计算出最终的地址address＝offset*zc qc_address，涉及的乘加运算转移到gpu的乘加专用计算单元，利用硬件优势提高运算速度。8.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤6.6所述的压缩先验信息存储和和非压缩先验信息存储结合的方法包括：行更新采用两种方式存储先验信息，压缩存储和非压缩存储，通过将更新的先验信息从放到寄存器内存缓存，提高了运行速度；其中压缩存储引入了额外的先验信息压缩和解压缩操作，降低了速度，节省了寄存器内存空间；非压缩存储直接将更新的先验信息缓存，提高了速度，增加了寄存器压力；二者结合，通过将该行更新的先验信息数量与设置的阈值t判断比较，先验信息数量大时，选用压缩存储，先验信息数量小时选用非压缩存储，合理利用了寄存器提高译码迭代速度；步骤6.6所述的先验信息采用部分更新策略包括：基矩阵中的检验行除去核心阵校验行，其他校验行最后一列构成一个单位阵，利用单位矩阵结构，在更新过程中省略每行最后一列数据运算；即非核心阵校验行更新的先验信息数量为行重d减1。9.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤6.7所述的压缩先验信息两两更新方法包括：将两个输入信息进行大小排序，将绝对值大的放在高位，更新记录最小值的位置，记为info1；然后再输入两个信息，记为info2；先用info2的高位与info1的高位比较，选择更小的值作为info1高位；再用info2的高位与info1的低位比较，如果info2的高位更小，将info1高低位进行交换，否则info1不变，进而得到新的info1；再用info2的低位分别与info1进行相同的高位、低位比较操作，得到更新结果info1，从而完成了一轮的两两比较；步骤6.8非压缩信息基于box_plus函数更新方法包括：y
i
＝boxplus(x
j
)＝min(abs(x
j
))*xor(sign(x
j
))，i∈q，j∈q\i，其中q为向量所有位置的集合；即第i个位置计算的值y
i
的绝对值等于除去第i个位置其他所有位置x
j
的绝对值最小值，y
i
的符号为其他所有x
j
的符号的异或运算，0表示正，1表示负。10.根据权利要求2所述的一种基于gpu的5g多用户ldpc码高速译码方法，其特征在于：步骤7所述的硬判决比特打包方法包括两个阶段：第一阶段为数据判决，32个线程为1组，每组中每个线程到共享内存加载一个16比特llr数据，取符号位得到0或者1的数值，0表示正数，1表示负数；第二阶段为比特打包，每个线程将自己持有的符号传递给__ballot_sync(*)投票函数，每个线程均返回一个相同的32比特掩码mask，掩码中的比特1表示对应线程号的输入数据非0，从而得到了比特打包后的结果，再通过任意一个线程，通常为0号线程，将该组线程的结果放到全局内存。

技术总结
本发明提出一种基于GPU的5G多用户LDPC码高速译码器及其译码方法，包括：一种高速译码器架构及高速译码方法；其中，译码方法包括：1：主机端存储空间初始化；2：GPU设备初始化；3：LDPC基矩阵信息结构体重描述；4：主机端调度GPU译码；5：LLR信息拷贝；6：GPU端根据用户码块信息分配相应的线程数量，选择对应的基矩阵信息结构体，基于分层最小和算法进行迭代译码；7：符号判决；8：结果回传到主机端。本发明充分结合分层译码算法的特点和GPU的架构特点，充分利用GPU片上资源，提升访存效率和数据计算单元利用率，降低单个码块消耗资源的同时降低单个码块的译码时间，提高整体的信息吞吐量，所述的译码方式更适用于处理实际场景中多小区多用户的LDPC码块译码。区多用户的LDPC码块译码。区多用户的LDPC码块译码。


技术研发人员：刘荣科 李岩松 田铠瑞 王若诗
受保护的技术使用者：北京航空航天大学
技术研发日：2022.05.09
技术公布日：2022/8/9