[go: up one dir, main page]

CN111966405B - 一种基于GPU的Polar码高速并行译码方法 - Google Patents

一种基于GPU的Polar码高速并行译码方法 Download PDF

Info

Publication number
CN111966405B
CN111966405B CN202010629868.3A CN202010629868A CN111966405B CN 111966405 B CN111966405 B CN 111966405B CN 202010629868 A CN202010629868 A CN 202010629868A CN 111966405 B CN111966405 B CN 111966405B
Authority
CN
China
Prior art keywords
thread
stage
local
thread block
parallel
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
CN202010629868.3A
Other languages
English (en)
Other versions
CN111966405A (zh
Inventor
李舒
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Beihang University
Original Assignee
Beihang University
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Beihang University filed Critical Beihang University
Priority to CN202010629868.3A priority Critical patent/CN111966405B/zh
Publication of CN111966405A publication Critical patent/CN111966405A/zh
Application granted granted Critical
Publication of CN111966405B publication Critical patent/CN111966405B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005—Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5011—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals
    • G06F9/5016—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals the resource being the memory
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/54—Interprogram communication
    • G06F9/544—Buffers; Shared memory; Pipes
    • H—ELECTRICITY
    • H03—ELECTRONIC CIRCUITRY
    • H03M—CODING; DECODING; CODE CONVERSION IN GENERAL
    • H03M13/00—Coding, decoding or code conversion, for error detection or error correction; Coding theory basic assumptions; Coding bounds; Error probability evaluation methods; Channel models; Simulation or testing of codes
    • H03M13/03—Error detection or forward error correction by redundancy in data representation, i.e. code words containing more digits than the source words
    • H03M13/05—Error detection or forward error correction by redundancy in data representation, i.e. code words containing more digits than the source words using block codes, i.e. a predetermined number of check bits joined to a predetermined number of information bits
    • H03M13/13—Linear codes
    • H—ELECTRICITY
    • H04—ELECTRIC COMMUNICATION TECHNIQUE
    • H04L—TRANSMISSION OF DIGITAL INFORMATION, e.g. TELEGRAPHIC COMMUNICATION
    • H04L1/00—Arrangements for detecting or preventing errors in the information received
    • H04L1/004—Arrangements for detecting or preventing errors in the information received by using forward error control
    • H04L1/0056—Systems characterized by the type of code used
    • H04L1/0057—Block codes

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • General Physics & Mathematics (AREA)
  • Computer Networks & Wireless Communication (AREA)
  • Signal Processing (AREA)
  • Probability & Statistics with Applications (AREA)
  • Compression Or Coding Systems Of Tv Signals (AREA)

Abstract

本发明公开了一种基于GPU的Polar码高速并行译码方法，整个译码过程可以分为三个阶段：初始化阶段、译码阶段、结果回传阶段，具体包括：步骤1：主机初始化；步骤2：GPU初始化；步骤3：译码内核函数进行若干次循环迭代，最大循环次数由程序预先设定；步骤4：对于因子图p_good的所有线程块的0号线程，将其共享内存中的Local_L[][0]+Local_R[][0]经过逆置换后，作为译码结果；步骤5：主机将译码结果从GPU传回到主机。本发明方法包括了三个层次的并行，即多子图之间、多线程块之间和多线程之间的并行。此外，本发明方法最大限度地降低了内核函数的启动开销；提高了访存效率和运行速度。

Description

一种基于GPU的Polar码高速并行译码方法
技术领域
本发明属于通信技术领域，涉及一种基于GPU(Graphics Processing Unit，图形处理器)的Polar码高速并行译码方法。
背景技术
Polar码由Erdal Arikan于2008年提出(参考文献[1]:Erdal Arikan,“ChannelPolarization:A Method for Constructing Capacity-Achieving Codes”,IEEEISIT2008)，是目前唯一能够被严格证明可以达到香农极限的信道编码方法。Polar码已经被5G标准化组织正式采纳。Polar码的译码方法可以分为两类：基于串行抵消的方法和基于置信传播的方法。基于串行抵消的方法运算量较小，但是算法本质上是串行的，因此译码延迟较大；对于基于置信传播的方法来说，为了保证Polar码译码的纠错性能，通常采用置信传播列表算法，即基于多个置换因子图的迭代算法，因此这种译码方法运算量很大，但是置信传播列表算法具有并行实现的潜力。
另一方面，近年来GPU技术得到了迅猛的发展，一张商业级的GPU卡上可以拥有超过4000个并行处理的核心，这为并行计算提供了高性价比的硬件基础。
发明内容
本发明的目的在于提供一种基于GPU的Polar码高速并行译码方法，以实现低延迟、高吞吐量的译码。
本发明提出了一种基于GPU的Polar码高速并行译码方法。本方法包括三个层次的并行，可以充分利用GPU上的核资源。本发明还设计了高效的分布式存储方法，提高了访存效率和运行速度。整个译码过程可以分为三个阶段：初始化阶段、译码阶段、结果回传阶段。初始化阶段包括以下步骤1和步骤2，译码阶段包括以下步骤3和步骤4，以下步骤5是结果回传阶段。
步骤1：主机初始化。依次包括：为信息比特标志、因子图置换和逆置换信息、接收机接收到的信号、译码的结果即源比特的对数似然比分配内存空间，信息和变量的初始化，存储接收到的信号并计算编码比特对数似然比。
步骤2：GPU初始化。依次包括：GPU全局内存分配，主机将数据发送给GPU，启动GPU的并行译码线程，GPU分配共享内存，初始化共享内存，根据全局内存给共享内存的数组赋值。
步骤3：译码内核函数进行若干次循环迭代，最大循环次数由程序预先设定。每次循环依次包括：L1阶段、L1-L2阶段间交换线程块共享内存、L2阶段、R1阶段、R1-R2阶段间交换线程块共享内存、R2阶段和循环终止条件判断。如果在循环过程中有因子图满足早期终止条件，或者已达到最大循环次数，则设置变量p_good，并终止循环，跳转到步骤4。
步骤4：对于因子图p_good的所有线程块的0号线程，即线程((p_good,b),0)，其中b＝0,1,...,N1-1，
Figure BDA0002568040640000021
N为Polar码的码长，将其共享内存中的Local_L[][0]+Local_R[][0]经过逆置换后，作为译码结果。
步骤5：主机将译码结果从GPU传回到主机。
其中，步骤3循环迭代的每次循环包括以下步骤：
步骤3.1：向左迭代的第一阶段，即L1阶段，包括第n-1,...,n-n1级迭代，其中n＝log2N，n1＝log2N1；
步骤3.2：每个因子图的线程块之间通过全局内存交换共享内存中的Local_L[][n-n1+1]；
步骤3.3：向左迭代的第二阶段，即L2阶段，包括第n-n1-1,...,0级迭代；
步骤3.4：向右迭代的第一阶段，即R1阶段，包括第0,...,n-n1-1级迭代；
步骤3.5：每个因子图的线程块之间通过全局内存交换共享内存中的Local_R[][n-n1]；
步骤3.6：向右迭代的第二阶段，即R2阶段，包括第n-n1,...,n-1级迭代；
步骤3.7：判断是否有因子图满足早期终止条件或者已达到最大循环次数，并设置变量p_good。
其中，步骤3所述的L1阶段、L2阶段、R1阶段和R2阶段，每个阶段都包括三个层次的并行具体如下：
第一个层次是多因子图之间的并行，每个因子图由N1个线程块负责。因为不同因子图在迭代过程没有数据依赖关系，所以不同因子图的线程块可以自然地并行运行。
第二个层次是同一因子图多线程块的并行。当Polar码的码长N较大时，GPU上一个流多处理器的核资源和共享内存资源不能支持一个因子图充分并行化的需求。为此，在本发明中，由N1个线程块负责一个因子图的迭代。本发明将每次迭代的向左传播和向右传播各分为两个阶段，共四个阶段。向左传播包括第n-1,..,1,0级的迭代，分为两个阶段：第一阶段为第n-1级到第n-n1级，称为L1阶段；第二阶段为第n-n1-1级到第0级，称为L2阶段。向右传播和向左传播的方向相反，包括第0,1,...,n-1级的迭代，分为两个阶段：第一阶段为第0级到第n-n1-1级，称为R1阶段；第二阶段为第n-n1级到第n-1级，称为R2阶段。
第三个层次是同一线程块内的多线程并行。每个线程块在每一级的计算可以分为N/N1/2＝2n-n1-1个子任务，各子任务之间没有数据依赖关系，各子任务被分成min(T,2n-n1-1)组，每组子任务由线程块内的一个线程负责，各线程可以并行执行；每个线程完成所负责的这组子任务后，进行线程块内的线程同步。
其中，第二个层次所述的同一因子图多线程块的并行和第三个层次所述的同一线程块内的多线程并行(即多个子任务)的分工细节如下：
(1)在L1阶段，不同线程块的数据之间没有依赖关系，各线程块可以并行运行，即L1阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即L1阶段同一线程块内多线程并行；
(2)在L2阶段，不同线程块的数据之间没有依赖关系，各线程块可以并行运行，即L2阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即L2阶段同一线程块内多线程并行；
(3)在R1阶段，不同线程块的数据之间没有依赖关系，各线程块可以并行运行，即R1阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即R1阶段同一线程块内多线程并行；
(4)在R2阶段，不同线程块的数据之间没有依赖关系，各线程块可以并行运行，即R2阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即R2阶段同一线程块内多线程并行；
其中，步骤2在GPU初始化时，译码过程中使用的L和R数组分布式存储在各线程块共享内存，即在一次完整的循环过程中，只需要通过全局内存在线程块之间交换共享内存2次，其他所有操作都可以使用线程块内的共享内存。具体如下：
GPU上的存储空间主要包含全局内存和线程块内共享内存(简称为块内共享内存或共享内存)。共享内存访问速度快，但是容量相对较小；全局内存空间大，访问速度较慢。
Polar码置信传播译码的主要数据为矩阵L和R，这两个矩阵的大小都是N*(n+1)。当码长N较大时，L和R所需的存储空间较大。为了提高访存速度，本发明将计算过程中的L和R进行分割后，存放在N个线程块的块内共享内存中的二维数组Local_L和Local_R中，Local_L的第一维和第二维大小分别为N/N1和n+1，Local_R的大小和Local_L相同。矩阵L和R分布式存储在各线程块内的共享内存方法如下：
(1)对于0<＝j<＝n-n1，Lb*(N/N1)+d2*N1+d1,j和Rb*(N/N1)+d2*N1+d1,j分别存放在第b个线程块的共享内存中的Local_L[d2*N1+d1][j]和Local_R[d2*N1+d1][j]，其中b＝0,1,...,N1-1；d2＝0,1,...,N/(N1*N1)-1；d1＝0,1,...,N1-1；
(2)对于n-n1<＝j<＝n，Lb*(N/N1)+d2*N1+d1,j和Rb*(N/N1)+d2*N1+d1,j存放在第d1个线程块的共享内存中的Local_L[d2*N1+b][j+1]和Local_R[d2*N1+b][j+1]，其中b＝0,1,...,N1-1；d2＝0,1,...,N/(N1*N1)-1；d1＝0,1,...,N1-1。
根据上述的分布式存储方法，在L1阶段的各级(s＝n-1,...,n-n1-1)，L和R分布式存储在各线程块的方案相同，所以各线程块在L1阶段内的各级迭代中使用块内共享内存即可，不需要在线程块之间交换数据。同理，在L2、R1和R2各阶段内的迭代，也不需要在线程块之间交换数据。L1和L2阶段的分布式存储方案不同，所以在完成L1阶段和开始L2阶段之间，需要通过全局内存在线程块之间交换数据；L2和R1阶段的分布式存储方案相同，所以在完成L2阶段和开始R1阶段之间，不需要通过全局内存在线程块之间交换数据。同理，在完成R1阶段和开始R2阶段之间，需要通过全局内存在线程块之间交换数据；在完成R2阶段和开始下一次迭代的L1阶段之间，不需要通过全局内存在线程块之间交换数据。由此可见，本发明在一次完整的循环(即包括4个阶段，共2n级迭代)过程中，只需要通过全局内存在线程块之间交换共享内存2次，其他所有操作都可以使用线程块内的共享内存。
其中，步骤2所述的GPU全局内存分配，为了提高访存的局部性，本发明做了如下优化：
(1)全局内存包括所有因子图的公共数据和每个因子图的专有数据。每个因子图使用的全局内存是连续存放的。第p个线程块使用的全局内存为结构体graph_info[p]，其成员包括：因子图置换数组、逆置换数组和用于线程块交换共享内存的全局内存空间。
(2)用于线程块之间交换共享内存的graph_info[p].swap按照线程块读取的顺序连续存放，即每个线程块从graph_info[p].swap读取时，其地址空间是连续的。
本发明的优点与积极效果在于：
(1)本发明的并行译码包括了三个层次的并行：第一个层次是多因子图之间的并行，每个因子图由若干个线程块负责。因为不同因子图在迭代过程没有数据依赖关系，所以不同因子图的线程块可以自然地并行运行；第二个层次是同一因子图多线程块的并行，在本发明中，一个因子图的迭代由N1个线程块负责，本发明设计的线程块分工保证了同一阶段内(即L1、L2、R1、R2)，线程块之间没有数据依赖，从而这N1个线程块可以并行执行；第三个层次是同一线程块内的多线程并行，本发明将同一线程块在每一级的迭代分为N/2N1个相互之间没有依赖的子任务，然后将这些子任务分配给线程块内的多个线程并行执行，每个线程执行完本级所负责的所有子任务后，进行线程块内的同步。这种方法可以充分利用GPU的并行核资源，计算效率高。而且本发明的同步开销小：在一次完整的循环(即包括4个阶段，共2n级迭代)过程中，本发明只需要使用2次线程块之间的同步，其余2n-2级之间都使用线程块内的同步机制，同步开销小。此外，GPU上整个译码过程使用一个内核函数，最大限度地降低了内核函数的启动开销。
(2)本发明将迭代中的主要数据分布式存储在各线程块内的共享内存，从而提高了访存效率和运行速度。本发明在一次完整的循环(即包括4个阶段，共2n级迭代)过程中，只需要通过全局内存在线程块之间交换共享内存2次，其他所有操作都可以使用线程块内的共享内存。此外，在本发明使用的全局空间中，同一个因子图的专有数据连续存放，且用于线程块之间交换共享内存的全局内存空间按照线程块读取的顺序连续存放，从而优化了存储的局部性，提高了访存效率和运行速度。
附图说明
图1为本发明所述的基于GPU的Polar码高速并行译码方法的流程图。
具体实施方式
下面结合附图对本发明进行详细说明。
本发明提供了一种基于GPU的Polar码高速并行译码方法，采用的是基于置信传播列表的译码算法。所述的译码方法包含三个阶段：初始化阶段、译码阶段、结果回传阶段。初始化阶段包括下述步骤1和步骤2，译码阶段包括下述步骤3和步骤4，下述步骤5是结果回传阶段，整个译码流程如图1所示，具体如下：
步骤1：主机初始化。依次包括：为信息比特标志、因子图置换和逆置换信息、接收机接收到的信号、译码的结果即源比特的对数似然比分配内存空间(步骤1.1)，信息和变量的初始化(步骤1.2)，存储接收到的信号并计算编码比特对数似然比(步骤1.3)；具体如下：
步骤1.1：分配主机内存空间，依次包括：信息比特标志数组InfoBitFlags，接收机接收到的信号数组y，编码比特的对数似然比数组cLLR。
步骤1.2：主机初始化信息比特标志数组InfoBitFlags，因子图置换数组Perm，根据因子图置换数组计算逆置换数组InvPerm。
步骤1.3：主机将接收机接收到的信号存放在数组y，信道的信噪比存放在变量SNR中。主机根据接收信号和信噪比计算编码比特对数似然比数组cLLR＝y*SNR。
步骤2：GPU初始化。依次包括：GPU全局内存分配，主机将数据发送给GPU，启动GPU的并行译码线程，GPU分配共享内存，初始化共享内存，根据全局内存给共享内存的数组赋值；具体过程如下：
步骤2.1：分配GPU上的全局内存。包括所有因子图的公共数据和每个因子图的专有数据。所有因子图的公共数据包括编码比特对数似然比数组cLLR、信息比特标志数组InfoBitFlags、译码结果即源比特对数似然比数组uLLR。每个因子图的专有数据连续存放，第p个因子图的专有数据存放在结构体graph_info[p]中，其成员包括：因子图置换数组Perm、逆置换数组InvPerm和用于线程块交换共享内存的全局内存空间swap。
步骤2.2：将主机内存的信息比特标志数组、因子图置换数组和逆置换数组、编码比特对数似然比数组发送给GPU，其中因子图置换数组和逆置换数组是因子图专有数据，存放在结构体graph_info[p]中。
步骤2.3：主机启动GPU的并行译码线程，共P*N1个线程块，每个线程块包含T个线程，其中T等于每个流多处理器包含的核心数。所有的线程执行相同的译码内核函数，线程之间通过线程索引来区分，每个线程的索引为((p,b),t)，其中(p,b)为线程块索引，t为线程块内的线程索引。
步骤2.4：GPU的线程分配线程块内的共享内存，包括二维数组Local_L[N/N1][n+2]和Local_R[N/N1][n+2]，所有元素初始化为0。
步骤2.5：GPU上每个线程块的0号线程，即线程((p,b),0)(p＝0,1,...,P-1；b＝0,1,...,2n1-1)根据信息比特标志和编码比特对数似然比给Local_L[][n+2]和Local_R[][0]赋值。
在程序中，可以用一重循环来实现本步骤，具体流程如下：
(1)循环索引变量为d，d＝0,1,...,N/N1-1
(2)计算dd＝(d％N1)*N1+(d/N1)
(3)局部内存中的Local_L[d][n+2]对应的全局内存地址为dd*N1+b，在原因子图(即置换前的因子图)的索引为graph_info[p].InvPerm[dd*N1+b]，因此将cLLR[graph_info[p].InvPerm[dd*N1+b]]的值赋给Local_L[d][n+2]。
(4)局部内存中的Local_R[d][0]对应的全局内存地址为b*N1+d，在原因子图(即置换前的因子图)的索引为graph_info[p].InvPerm[b*N1+d]。如果信息比特标志InfoBitFlags[graph_info[p].InvPerm[b*N1+d]]为0，将Local_R[d][0]置为1e+30。
示例程序如下：
Figure BDA0002568040640000071
Figure BDA0002568040640000081
索引为((p,b),0)，p＝0,1,...,P-1；b＝0,1,...,2n1-1的线程共有P*2n1个，这P*2n1线程可以并行执行。
步骤3：译码内核函数进行若干次循环迭代，最大循环次数由程序预先设定。每次循环包括L1阶段(步骤3.1)、L1-L2阶段间交换线程块共享内存(步骤3.2)、L2阶段(步骤3.3)、R1阶段(步骤3.4)、R1-R2阶段间交换线程块共享内存(步骤3.5)、R2阶段(步骤3.6)和判断是否有因子图满足早期终止条件或者已达到最大循环次数，并设置变量p_good(步骤3.7)。
步骤3.1：向左迭代的第一阶段，即L1阶段，级编号s＝n-1,...,n-n1.
每一级包含三个层次的并行：
(1)多因子图之间的并行：第p个因子图由索引为(p,0),...,(p,2n1-1)的线程块负责，p＝0,1,...,P-1。各因子图的线程块之间相互独立，可以并行进行；
(2)同一因子图多线程块的并行：每个因子图由2n1个线程块负责计算，第b(b＝0,1,...,2n1-1)个线程块使用的L和R的第一维索引集合为L1_Block(b)＝{ia*2n1+b,ia＝0,...,2n-n1-1}，因此各线程块之间没有数据依赖，可以并行运行。
(3)同一因子图线程块内的多线程并行：每个线程块在每一级的计算可以分解为N/N1/2＝2n-n1-1个子任务。编号为i(i＝0,1,...,2n-n1-1-1)的子任务在第s级使用的L和R的第一维索引集合为：
L1_Task(b,i,s)＝{floor(i/2s-n1)*2s+1+a*2s+(i％2s-n1)*2n1+b:a＝0,1}；
因此同一线程块内各子任务之间没有数据依赖关系，所以可以将这些子任务分成min(T,2n-n1-1)组，每组子任务由线程块内的一个线程负责，各线程可以并行执行。每个线程完成所负责的这组子任务后，进行线程块内的线程同步。
根据本发明的线程和子任务分工方案，线程块(p,b)在第s级的第i个子任务需要计算：
Lup,s＝f(Ldown,s+1+Rdown,s,Lup,s+1)
Ldown,s＝Ldown,s+1+f(Lup,s+1,Rup,s)
其中
up＝floor(i/2s-n1)*2s+1+(i％2s-n1)*2n1+b
down＝floor(i/2s-n1)*2s+1+2s+(i％2s-n1)*2n1+b
根据本发明的分布式存储方案，在L1阶段，up和down在共享内存中的索引分别为Local_up和Local_down,s和s+1在共享内存中的索引分别为s+1和s+2，其中Local_up＝floor(i’/2j)*2j+1+(i’％2j),Local_down＝floor(i’/2j)*2j+1+2j+(i’％2j),i’＝(i％2n -2n1)*2n-2n1+floor(i/2n-2n1),i’＝0,1,...,2n-n1-1-1,j＝s-n+n1，Local_up和Local_down是线程块(p,b)的共享内存地址。
因此以上运算用程序可以表示为：
Local_L[Local_up][s+1]＝f(Local_L[Local_down][s+2]+Local_R[Local_down][s+1],Local_L[up][s+2])；
Local_L[Local_down][s+1]＝Local_L[Local_down][s+2]+f(Local_L[Local_up][s+2],Local_R[up][s+1])；
函数f的定义为：f(x,y)＝2tanh-1(tanh(x/2)tanh(y/2))，在实际计算中，通常用f(x,y)≈0.9375sgn(x)sgn(y)min(|x|,|y|)来近似。
在程序中，可以用两重循环来实现本步骤，具体流程如下：
(1)第一重循环的循环索引变量为s，s＝n-1,...,n-n1。(注意在这个循环里，循环索引变量s每次递减1。)
(2)计算j＝s-n+n1。
(3)第二重循环的循环索引变量为i，i＝t,t+T,...,t+(floor(2n-n1-1-1-t)/T)*T。
(4)i的二进制表示包含n-n1-1位，将i的低j位，即i％(2j)赋给变量i_LSB；并将i的高(n-n1-1-j)位，即floor(i/(2j))赋给变量i_MSB。
(5)计算共享内存地址Local_up和Local_down，其中Local_up＝(i_MSB<<(j+1))+i_LSB，Local_down＝Local_up+(1<<j)。
(6)计算Local_L[Local_up][s+1]和Local_L[Local_down][s+1]，其中前者的值为f(Local_L[Local_down][s+2]+Local_R[Local_down][s+1],Local_L[up][s+2])，后者的值为Local_L[Local_down][s+1]＝Local_L[Local_down][s+2]+
f(Local_L[Local_up][s+2],Local_R[up][s+1])。
(7)第二重循环结束后，调用__syncthreads()进行线程块内的各线程同步。
示例程序如下：
Figure BDA0002568040640000101
步骤3.2：每个因子图的线程块之间通过全局内存交换共享内存中的Local_L[][n-n1+1]，即L1-L2阶段间交换线程块共享内存；具体如下：
步骤3.2.1每个线程块的0号线程，即线程((p,b),0)(p＝0,1,...,P-1；b＝0,1,...,2n1-1)将线程块(p,b)共享内存中的Local_L[d2*2n1+d1][n-n1+1]写入全局内存中的graph_info[p].swap[d1*2n-n1+d2*2n1+b],其中d1＝0,1,...,2n1-1；d2＝0,1,...,2n-2n1-1。编号为((p,b),0)，p＝0,1,...,P-1；b＝0,1,...,2n1-1的线程共有P*2n1个。因为各线程写入swap的地址之间没有重叠，所以这P*2n1线程可以并行执行。
步骤3.2.2每个因子图的2n1个线程块进行线程块之间的同步。不同因子图的线程块之间不需要同步。
步骤3.2.3线程块(p,d1)的0号线程((p,d1),0)将全局内存中的graph_info[p].swap[d1*2n-n1+d2*2n1+b]写入线程块(p,d1)共享内存中的Local_L[d2*2n1+b][n-n1],其中b＝0,1,...,2n1-1；d1＝0,1,...,2n1-1；d2＝0,1,...,2n-2n1-1。编号为((p,d1),0)，p＝0,1,...,P-1；d1＝0,1,...,2n1-1的线程共有P*2n1个，这P*2n1线程可以并行执行。
步骤3.3：向左迭代的第二阶段，即L2阶段，级编号s＝n-n1-1,...,0.
每一级也包含三个层次的并行：
(1)多因子图之间的并行：第p个因子图由编号为(p,0),...,(p,2n1-1)的线程块负责，p＝0,1,...,P-1。各因子图的线程块之间相互独立，可以并行进行；
(2)同一因子图多线程块的并行：每个因子图由2n1个线程块负责计算，第b(b＝0,1,...,2n1-1)个线程块使用的L和R的第一维索引集合为L2_Block(b)＝{b*2n-n1+ia,ia＝0,...,2n-n1-1}，因此各线程块之间没有数据依赖，可以并行运行。
(3)同一因子图线程块内的多线程并行：每个线程块在每一级的计算可以分解为N/N1/2＝2n-n1-1个子任务。编号为i(i＝0,1,...,2n-n1-1-1)的子任务在第s级使用的L和R的第一维索引集合为
L2_Task(b,i,s)＝{b*2n-n1+floor(i/2s)*2s+1+a*2s+(i％2s):a＝0,1}
因此同一线程块内各子任务之间没有数据依赖关系，所以可以将这些子任务分成min(T,2n-n1-1)组，每组子任务由线程块内的一个线程负责，各线程可以并行执行。每个线程完成所负责的这组子任务后，进行线程块内的线程同步。
根据本发明的线程和子任务分工方案，线程块(p,b)在第s级的第i个子任务需要计算：
Lup,s＝f(Ldown,s+1+Rdown,s,Lup,s+1)
Ldown,s＝Ldown,s+1+f(Lup,s+1,Rup,s)
其中，
up＝b*2n-n1+floor(i/2s)*2s+1+(i％2s)
down＝b*2n-n1+floor(i/2s)*2s+1+2s+(i％2s)
根据本发明的分布式存储方案，在L2阶段，up和down在共享内存中的索引分别为Local_up和Local_down,s和s+1在共享内存中的索引分别为s和s+1，其中Local_up＝floor(i/2s)*2s+1+(i％2s),Local_down＝floor(i/2s)*2s+1+2s+(i％2s)，Local_up和Local_down是线程块(p,b)的共享内存地址。
因此以上运算用程序可以表示为：
Local_L[Local_up][s]＝f(Local_L[Local_down][s+1]+Local_R[Local_down][s],Local_L[up][s+1])；
Local_L[Local_down][s]＝Local_L[Local_down][s+1]+f(Local_L[Local_up][s+1],Local_R[up][s])；
在程序中，可以用两重循环来实现本步骤，具体流程如下：
(1)第一重循环的循环索引变量为s，s＝n-n1,n-n1-1,...,0。(注意在这个循环里，循环索引变量s每次递减1。)
(2)第二重循环的循环索引变量为i，i＝t,t+T,...,t+(floor(2n-n1-1-1-t)/T)*T。
(3)i的二进制表示包含n-n1-1位，将i的低s位，即i％(2s)赋给变量i_LSB；并将i的高(n-n1-1-s)位，即floor(i/(2j))赋给变量i_MSB。
(4)计算共享内存地址Local_up和Local_down，其中Local_up＝(i_MSB<<(s+1))+i_LSB，Local_down＝Local_up+(1<<s)。
(5)计算Local_L[Local_up][s]和Local_L[Local_down][s]，其中前者的值为f(Local_L[Local_down][s+1]+Local_R[Local_down][s],Local_L[Local_up][s+1])，后者的值为Local_L[Local_down][s+1]＝L[Local_down][Local_s+1]+
f(L[Local_up][s+1],Local_R[Local_up][s])。
(6)第二重循环结束后，调用__syncthreads()进行线程块内的各线程同步。
示例程序如下：
Figure BDA0002568040640000121
步骤3.4：向右迭代的第一阶段，即R1阶段，级编号s＝0,...,n-n1-1.
每一级也包含三个层次的并行：
(1)多因子图之间的并行：第p个因子图由编号为(p,0),...,(p,2n1-1)的线程块负责，p＝0,1,...,P-1。各因子图的线程块之间相互独立，可以并行进行；
(2)同一因子图多线程块的并行：每个因子图由2n1个线程块负责计算，第b(b＝0,1,...,2n1-1)个线程块使用的L和R的第一维索引集合为R1_Block(b)＝{b*2n-n1+ia,ia＝0,...,2n-n1-1}，因此各线程块之间没有数据依赖，可以并行运行。
(3)同一因子图线程块内的多线程并行：每个线程块在每一级的计算可以分解为N/N1/2＝2n-n1-1个子任务。编号为i(i＝0,1,...,2n-n1-1-1)的子任务在第s级使用的L和R的第一维索引集合为R1_Task(b,i,s)＝{b*2n-n1+floor(i/2s)*2s+1+a*2s+(i％2s):a＝0,1}
因此同一线程块内各子任务之间没有数据依赖关系，所以可以将这些子任务分成min(T,2n-n1-1)组，每组子任务由线程块内的一个线程负责，各线程可以并行执行。每个线程完成所负责的这组子任务后，进行线程块内的线程同步。
根据本发明的线程和子任务分工方案，线程块(p,b)在第s级的第i个子任务需要计算：
Rup,s+1＝f(Ldown,s+1+Rdown,s,Rup,s)
Rdown,s+1＝Rdown,s+f(Lup,s+1,Rup,s)
其中，
up＝b*2n-n1+floor(i/2s)*2s+1+(i％2s)
down＝b*2n-n1+floor(i/2s)*2s+1+2s+(i％2s)
根据本发明的分布式存储方案，在R1阶段，up和down在共享内存中的索引分别为Local_up和Local_down,s和s+1在共享内存中的索引分别为s和s+1，其中Local_up＝floor(i/2s)*2s+1+(i％2s),Local_down＝floor(i/2s)*2s+1+2s+(i％2s)，Local_up和Local_down是线程块(p,b)的共享内存地址。
因此以上运算用程序可以表示为：
Local_R[Local_up][s+1]＝
f(Local_L[Local_down][s+1]+Local_R[Local_down][s],R[Local_up][s])；
Local_R[Local_down][s+1]＝Local_R[Local_down][s]+
f(L[Local_up][s+1],R[Local_up][s])；
在程序中，可以用两重循环来实现本步骤，具体流程如下：
(1)第一重循环的循环索引变量为s，s＝0,1,n-n1,...,n-n1-1。
(2)第二重循环的循环索引变量为i，i＝t,t+T,...,t+(floor(2n-n1-1-1-t)/T)*T。
(3)i的二进制表示包含n-n1-1位，将i的低s位，即i％(2s)赋给变量i_LSB；并将i的高(n-n1-1-s)位，即floor(i/(2s))赋给变量i_MSB。
(4)计算共享内存地址Local_up和Local_down，其中Local_up＝(i_MSB<<
(s+1))+i_LSB，Local_down＝Local_up+(1<<s)。
(5)计算Local_L[Local_up][s+1]和Local_L[Local_down][s+1]，其中前者的值为f(Local_L[Local_down][s+1]+Local_R[Local_down][s],Local_R[Local_up][s])，后者的值为Local_R[Local_down][s]+f(L[Local_up][s+1],R[Local_up][s])。
(6)第二重循环结束后，调用__syncthreads()进行线程块内的各线程同步。
示例程序如下：
Figure BDA0002568040640000141
步骤3.5：每个因子图的线程块之间通过全局内存交换共享内存中的Local_R[][n-n1]，即R1-R2阶段间交换线程块共享内存，具体如下：
步骤3.5.1每个线程块的0号线程，即线程((p,b),0)(p＝0,1,...,P-1；b＝0,1,...,2n1-1)将线程块(p,b)共享内存中的Local_R[d2*2n1+d1][n-n1]写入全局内存中的graph_info[p].swap[d1*2n-n1+d2*2n1+b],其中d1＝0,1,...,2n1-1；d2＝0,1,...,2n-2n1-1。编号为((p,b),0)，p＝0,1,...,P-1；b＝0,1,...,2n1-1的线程共有P*2n1个。因为各线程写入swap的地址之间没有重叠，所以这P*2n1线程可以并行执行。
步骤3.5.2每个因子图的2n1个线程块进行线程块之间的同步。不同因子图的线程块之间不需要同步。
步骤3.5.3线程块(p,d1)的0号线程((p,d1),0)将全局内存中的graph_info[p].swap[d1*2n-n1+d2*2n1+b]写入线程块共享内存中的Local_R[d2*2n1+b][n-n1+1],其中b＝0,1,...,2n1-1；d1＝0,1,...,2n1-1；d2＝0,1,...,2n-2n1-1。编号为((p,d1),0)，p＝0,1,...,P-1；d1＝0,1,...,2n1-1的线程共有P*2n1个，这P*2n1线程可以并行执行。
步骤3.6：向右迭代的第二阶段，即R2阶段，级编号s＝n-n1,...,n-1.
每一级也包含三个层次的并行：
(1)多因子图之间的并行：第p个因子图由编号为(p,0),...,(p,2n1-1)的线程块负责，p＝0,1,...,P-1。各因子图的线程块之间相互独立，可以并行进行；
(2)同一因子图多线程块的并行：每个因子图由2n1个线程块负责计算，第b(b＝0,1,...,2n1-1)个线程块使用的L和R的第一维索引集合为R2_Block(b)＝{ia*2n1+b,ia＝0,...,2n-n1-1}，因此各线程块之间没有数据依赖，可以并行运行。
(3)同一因子图线程块内的多线程并行：每个线程块在每一级的计算可以分解为N/N1/2＝2n-n1-1个子任务。编号为i(i＝0,1,...,2n-n1-1-1)的子任务在第s级使用的L和R的第一维索引集合为：
R2_Task(b,i,s)＝{floor(i/2s-n1)*2s+1+a*2s+(i％2s-n1)*2n1+b:a＝0,1}
因此同一线程块内各子任务之间没有数据依赖关系，所以可以将这些子任务分成min(T,2n-n1-1)组，每组子任务由线程块内的一个线程负责，各线程可以并行执行。每个线程完成所负责的这组子任务后，进行线程块内的线程同步。
根据本发明的线程和子任务分工方案，线程块(p,b)在第s级的第i个子任务需要计算：
Rup,s+1＝f(Ldown,s+1+Rdown,s,Rup,s)
Rdown,s+1＝Rdown,s+f(Lup,s+1,Rup,s)
其中，
up＝floor(i/2s-n1)*2s+1+(i％2s-n1)*2n1+b
down＝floor(i/2s-n1)*2s+1+2s+(i％2s-n1)*2n1+b
根据本发明的分布式存储方案，在R2阶段，up和down在共享内存中的索引分别为Local_up和Local_down,s和s+1在共享内存中的索引分别为s+1和s+2，其中Local_up＝floor(i’/2j)*2j+1+(i’％2j),Local_down＝floor(i’/2j)*2j+1+2j+(i’％2j),i’＝(i％2n -2n1)*2n-2n1+floor(i/2n-2n1),i’＝0,1,...,2n-n1-1-1,j＝s-n+n1，Local_up和Local_down是线程块(p,b)的共享内存地址。
因此以上运算用程序可以表示为：
Local_R[Local_up][s+2]＝f(Local_L[Local_down][s+2]+Local_R[Local_down][s+1],[Local_up][s+1])；
Local_R[Local_down][s+2]＝Local_R[Local_down][s+1]+
f(L[Local_up][s+2],R[Local_up][s+1])；
在程序中，可以用两重循环来实现本步骤，具体流程如下：
(1)第一重循环的循环索引变量为s，s＝n-n1,...,n-1。
(2)计算j＝s-n+n1。
(3)第二重循环的循环索引变量为i，i＝t,t+T,...,t+(floor(2n-n1-1-1-t)/T)*T。
(4)i的二进制表示包含n-n1-1位，将i的低j位，即i％(2j)赋给变量i_LSB；并将i的高(n-n1-1-j)位，即floor(i/(2j))赋给变量i_MSB。
(5)计算共享内存地址Local_up和Local_down，其中Local_up＝(i_MSB<<(j+1))+i_LSB，Local_down＝Local_up+(1<<j)。
(6)计算Local_L[Local_up][s+2]和Local_L[Local_down][s+2]，其中前者的值为f(Local_L[Local_down][s+2]+Local_R[Local_down][s+1],Local_R[Local_up][s+1])，后者的值为Local_R[Local_down][s+1]+f(L[Local_up][s+2],Local_R[Local_up][s+1])。
(7)第二重循环结束后，调用__syncthreads()进行线程块内的各线程同步。
示例程序如下：
Figure BDA0002568040640000161
Figure BDA0002568040640000171
步骤3.7：判断每个因子图的迭代结果是否符合早期终止条件。如果至少有一个因子图满足条件，将这个因子图的编号p记录在变量p_good中(如果有多个因子图满足条件，将任一满足条件的p记录在变量p_good中)，循环终止即跳转到步骤4。否则，即所有因子图都不满足条件，判断是否已经达到预设的最大循环次数，如果已达到最大循环次数，令p_good＝0(对应第一个因子图)，循环终止即跳转到步骤4；如果未达到预设的最大循环次数，继续循环，即跳转到步骤3.1。其中，所述的早期终止条件可以有多种条件，例如迭代没有改进，通过附加的CRC校验等。其中，最大循环次数可以选择和Polar码置信传播译码的串行化实现相同的最大循环次数，通常为50-200之间。
步骤4：因子图p_good的所有线程块的0号线程，即线程((p_good,b),0)，其中b＝0,1,...,N1-1，将其共享内存中的Local_L[][0]+Local_R[][0]经过逆置换后，作为译码结果，存放在uLLR中。有N1个索引为((p_good,b),0)的线程，这些线程可以并行执行。
在程序中，可以用一重循环来实现本步骤，具体流程如下：
(1)循环索引变量为d，d＝0,1,...,N/N1-1
(2)局部内存中的Local[L][0]和Local_R[d][0]对应的全局内存地址为b*N1+d，在原因子图(即置换前的因子图)的索引为graph_info[p_good].InvPerm[b*N1+d]。将Local_L[d][0]和Local_R[d][0]相加，并将结果存放到uLLR[graph_info[p_good].InvPerm[b*N1+d]]中。
示例程序如下：
for(d＝0；d<N/N1；d++)
uLLR[graph_info[p_good].InvPerm[b*N1+d]]＝
Local_L[d][0]+Local_R[d][0]；
步骤5：主机将译码结果即源比特的对数似然比uLLR从GPU传回到主机。

Claims (6)

1.一种基于GPU的Polar码高速并行译码方法，其特征在于：整个译码过程分为三个阶段：初始化阶段、译码阶段、结果回传阶段，其中初始化阶段包括以下步骤1和步骤2，译码阶段包括以下步骤3和步骤4，以下步骤5是结果回传阶段：
步骤1：主机初始化
依次包括：为信息比特标志、因子图置换和逆置换信息、接收机接收到的信号、译码的结果即源比特的对数似然比分配内存空间，信息和变量的初始化，存储接收到的信号并计算编码比特对数似然比；
步骤2：GPU初始化
依次包括：GPU全局内存分配，主机将数据发送给GPU，启动GPU的并行译码线程，GPU分配共享内存，初始化共享内存，根据全局内存给共享内存的数组赋值；
步骤3：译码内核函数进行若干次循环迭代，最大循环次数由程序预先设定
每次循环依次包括：L1阶段、L1-L2阶段间交换线程块共享内存、L2阶段、R1阶段、R1-R2阶段间交换线程块共享内存、R2阶段和循环终止条件判断：如果在循环过程中有因子图满足早期终止条件，或者已达到最大循环次数，则设置变量p_good，并终止循环，跳转到步骤4；
步骤4：对于变量p_good的所有线程块的0号线程，即线程((p_good,b),0)，其中b＝0,1,...,N1-1，
Figure FDA0003586617490000011
N为Polar码的码长，将其共享内存中的Local_L[][0]+Local_R[][0]经过逆置换后，作为译码结果；
步骤5：主机将译码结果从GPU传回到主机。
2.根据权利要求1所述的一种基于GPU的Polar码高速并行译码方法，其特征在于：所述步骤2在GPU初始化时，译码过程中使用的L和R数组分布式存储在各线程块共享内存，即在一次完整的循环过程中，只需要通过全局内存在线程块之间交换共享内存2次，其他所有操作都使用线程块内的共享内存。
3.根据权利要求1所述的一种基于GPU的Polar码高速并行译码方法，其特征在于：步骤2所述的全局内存的分配，具体为：同一个因子图使用的全局内存连续存放，并且用于线程块之间交换共享内存的全局内存空间按照线程块读取的顺序连续存放，即每个线程块从交换空间读取时，其读取的地址空间是连续的。
4.根据权利要求1所述的一种基于GPU的Polar码高速并行译码方法，其特征在于：所述步骤3循环迭代的每次循环包括以下步骤：
步骤3.1：向左迭代的第一阶段，即L1阶段，包括第n-1,...,n-n1级迭代，其中n1＝log2N1；
步骤3.2：每个因子图的线程块之间通过全局内存交换共享内存中的Local_L[][n-n1+1]；
步骤3.3：向左迭代的第二阶段，即L2阶段，包括第n-n1-1,...,0级迭代；
步骤3.4：向右迭代的第一阶段，即R1阶段，包括第0,...,n-n1-1级迭代；
步骤3.5：每个因子图的线程块之间通过全局内存交换共享内存中的Local_R[][n-n1]；
步骤3.6：向右迭代的第二阶段，即R2阶段，包括第n-n1,...,n-1级迭代；
步骤3.7：判断是否有因子图满足早期终止条件或者已达到最大循环次数，并设置变量p_good。
5.根据权利要求4所述的一种基于GPU的Polar码高速并行译码方法，其特征在于：步骤3所述的L1阶段、L2阶段、R1阶段和R2阶段，每个阶段都包括三个层次的并行：
第一个层次是多因子图之间的并行，每个因子图由N1个线程块负责，各因子图的线程块之间相互独立，不同因子图的线程块并行运行；
第二个层次是同一因子图多线程块的并行，每个因子图由N1个线程块负责计算，不同线程块之间没有数据依赖，并行运行；
第三个层次是同一线程块内的多线程并行，每个线程块在每一级的计算分为
Figure FDA0003586617490000021
个子任务，各子任务之间没有数据依赖关系，各子任务被分成
Figure FDA0003586617490000022
组,其中T是GPU上每个流多处理器包含的核心数，每组子任务由线程块内的一个线程负责，各线程并行执行；每个线程完成所负责的这组子任务后，进行线程块内的线程同步。
6.根据权利要求5所述的一种基于GPU的Polar码高速并行译码方法，其特征在于：第二个层次所述的同一因子图多线程块的并行和第三个层次所述的同一线程块内的多线程并行的分工细节如下：
(1)在L1阶段，不同线程块的数据之间没有依赖关系，各线程块并行运行，即L1阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即L1阶段同一线程块内多线程并行；
(2)在L2阶段，不同线程块的数据之间没有依赖关系，各线程块并行运行，即L2阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即L2阶段同一线程块内多线程并行；
(3)在R1阶段，不同线程块的数据之间没有依赖关系，各线程块并行运行，即R1阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即R1阶段同一线程块内多线程并行；
(4)在R2阶段，不同线程块的数据之间没有依赖关系，各线程块并行运行，即R2阶段同一因子图多线程块并行；同一线程块在第s级的各子任务之间没有数据依赖关系，将这些子任务分配到线程块内的多个线程并行执行，即R2阶段同一线程块内多线程并行。
CN202010629868.3A 2020-07-03 2020-07-03 一种基于GPU的Polar码高速并行译码方法 Active CN111966405B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202010629868.3A CN111966405B (zh) 2020-07-03 2020-07-03 一种基于GPU的Polar码高速并行译码方法

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202010629868.3A CN111966405B (zh) 2020-07-03 2020-07-03 一种基于GPU的Polar码高速并行译码方法

Publications (2)

Publication Number Publication Date
CN111966405A CN111966405A (zh) 2020-11-20
CN111966405B true CN111966405B (zh) 2022-07-26

Family

ID=73361314

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202010629868.3A Active CN111966405B (zh) 2020-07-03 2020-07-03 一种基于GPU的Polar码高速并行译码方法

Country Status (1)

Country Link
CN (1) CN111966405B (zh)

Families Citing this family (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN113014270B (zh) * 2021-02-22 2022-08-05 上海大学 码长可配置的部分折叠极化码译码器

Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN101814039A (zh) * 2010-02-02 2010-08-25 北京航空航天大学 一种基于GPU的Cache模拟器及其空间并行加速模拟方法
CN105843590A (zh) * 2016-04-08 2016-08-10 深圳航天科技创新研究院 一种并行指令集预译码方法及系统
CN108462495A (zh) * 2018-04-03 2018-08-28 北京航空航天大学 一种基于gpu的多元ldpc码高速并行译码器及其译码方法
CN111026444A (zh) * 2019-11-21 2020-04-17 中国航空工业集团公司西安航空计算技术研究所 一种gpu并行阵列simt指令处理模型

Family Cites Families (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US9362956B2 (en) * 2013-01-23 2016-06-07 Samsung Electronics Co., Ltd. Method and system for encoding and decoding data using concatenated polar codes
CN107545303B (zh) * 2016-01-20 2021-09-07 中科寒武纪科技股份有限公司 用于稀疏人工神经网络的计算装置和运算方法
GB2567507B (en) * 2018-04-05 2019-10-02 Imagination Tech Ltd Texture filtering with dynamic scheduling

Patent Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN101814039A (zh) * 2010-02-02 2010-08-25 北京航空航天大学 一种基于GPU的Cache模拟器及其空间并行加速模拟方法
CN105843590A (zh) * 2016-04-08 2016-08-10 深圳航天科技创新研究院 一种并行指令集预译码方法及系统
CN108462495A (zh) * 2018-04-03 2018-08-28 北京航空航天大学 一种基于gpu的多元ldpc码高速并行译码器及其译码方法
CN111026444A (zh) * 2019-11-21 2020-04-17 中国航空工业集团公司西安航空计算技术研究所 一种gpu并行阵列simt指令处理模型

Non-Patent Citations (1)

* Cited by examiner, † Cited by third party
Title
Polar码并行级联结构设计及性能分析;潘小飞;《通信技术》;20160210;全文 *

Also Published As

Publication number Publication date
CN111966405A (zh) 2020-11-20

Similar Documents

Publication Publication Date Title
US11048517B2 (en) Decoupled processor instruction window and operand buffer
US9477465B2 (en) Arithmetic processing apparatus, control method of arithmetic processing apparatus, and a computer-readable storage medium storing a control program for controlling an arithmetic processing apparatus
US9817919B2 (en) Agglomerative treelet restructuring for bounding volume hierarchies
US10007605B2 (en) Hardware-based array compression
CN100461094C (zh) 一种针对流处理器的指令控制方法
US20150205607A1 (en) Tree-based thread management
US8312227B2 (en) Method and apparatus for MPI program optimization
CN101717817B (zh) 对基于随机上下文无关文法的rna二级结构预测进行加速的方法
US9229717B2 (en) Register allocation for clustered multi-level register files
Ueno et al. Parallel distributed breadth first search on GPU
JP2018518776A (ja) プロセッサ命令ウィンドウへの命令ブロックのバルク割り当て
JP2018519597A (ja) ブロックサイズに基づくマッピング命令ブロック
WO2020220935A1 (zh) 运算装置
US10599638B2 (en) System and method for identifying maximal independent sets in parallel
Jin et al. A multi-level optimization method for stencil computation on the domain that is bigger than memory capacity of GPU
CN104375807A (zh) 基于众核协处理器的三级流水序列比对方法
CN111966405B (zh) 一种基于GPU的Polar码高速并行译码方法
US11429299B2 (en) System and method for managing conversion of low-locality data into high-locality data
CN118550697A (zh) 一种基于近内存处理架构的推荐模型推理加速系统
US9830161B2 (en) Tree-based thread management
Fuentes-Alventosa et al. Cuvle: Variable-length encoding on cuda
CN111966404B (zh) 一种基于gpu的规则稀疏码分多址scma高速并行译码方法
EP3304283B1 (en) System, apparatus, and method for temporary load instruction
US20210042111A1 (en) Efficient encoding of high fanout communications
KR20240064001A (ko) 인코딩된 인에이블 클록 게이터

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant