一种高吞吐量大带宽的通用信道化GPU算法

摘要

本发明公开了一种高吞吐量大带宽的通用信道化GPU算法，基于CUDA计算平台，包括以下步骤：S1：利用计算初始化模块预先申请信道化流程处理中所需的输入、输出及中间变量的存储空间；S2：利用参数配置模块设置目标窄带信号相关参数并根据这些相关参数将其与计算初始化模块中分配的存储资源进行配置处理；S3：利用核心功能模块执行实际的信道化计算，根据步骤S2给各路窄带信号分配对应的计算资源并且开启信道化流程内部各子功能模块的执行处理。本发明能保持较高的计算性能，满足相关应用场景的实时性处理指标。本发明在支持CUDA的多代GPU硬件平台上均能部署配置，并且算法性能可以自适应匹配GPU硬件本身的计算能力，能支撑信道化场景的灵活配置。

权利要求

1.一种高吞吐量大带宽的通用信道化GPU算法，其特征在于，基于CUDA计算平台，包括以下步骤：S1：利用计算初始化模块预先申请信道化流程处理中所需的输入、输出及中间变量的存储空间，所述存储空间包括主机内存区域和设备内存区域；

所述存储空间的组织形式为：所有窄带信号的输入、输出以及中间结果分配连续的存储空间，其中每路信号依次占用dataStep大小的间隔段落，并确保信道化业务流程中各功能模块处理时对应的输入信号长度countIn、输出信号长度countOut不会超过dataStep大小；

S2：利用参数配置模块设置目标窄带信号相关参数并根据这些相关参数将其与计算初始化模块中分配的存储资源进行配置处理，所述相关参数包括中心频率、带宽、输出采样率、目标增益，所述配置处理包括计算前批量配置模式和计算中动态配置模式；

所述计算中动态配置，包括计算中动态删除窄带信号和动态添加窄带信号，动态添加的信号有全局唯一的编号，其与存储空间区域标识将建立一个动态映射的关系，在动态删除时，根据所删除信号的编号查找与其映射的存储区域编号，并将其置为空闲状态，当删除的信号较多时，模块会对这些空闲存储区域按标识进行升序排列管理，后期动态添加目标窄带信号时，则优先选取队列头部的空闲存储区域标识与该信号建立映射关系，同时取消其空闲状态标记；

S3：利用核心功能模块执行实际的信道化计算，根据步骤S2给各路窄带信号分配对应的计算资源并且开启信道化流程内部各子功能模块的执行处理，所述计算资源包括CUDA中的线程网格Grid以及线程块Block。

2.根据权利要求1所述的一种高吞吐量大带宽的通用信道化GPU算法，其特征在于，所述步骤S3子功能模块包括混频子模块、半带滤波子模块、重采样处理子模块和低通滤波子模块，具体操作步骤：S301：将宽带输入信号数据从主机内存缓冲区传输到对应的设备内存缓冲区，主机内存缓冲区使用页锁定内存；

S302：将步骤301中设备内存缓冲区的信号数据依据各路信号的混频因子进行外差混频操作，各路信号对应混频因子在参数配置模块设置每路信号信息时提前计算；

S303：将步骤S302中各路信号对应外差混频计算中间结果作为输入进行半带滤波处理；

S304：将步骤S303中各路信号对应半带滤波处理中间结果作为输入进行重采样处理；

S305：将步骤S304中各路信号对应重采样处理中间结果作为输入进行低通滤波处理；

S306：将步骤S305中各路信号对应低通滤波处理结果由对应设备内存缓冲区传输到主机内存缓冲区并作为最终的信道化处理结果。

3.根据权利要求1所述的一种高吞吐量大带宽的通用信道化GPU算法，其特征在于，所述计算初始化模块只需根据设定参数执行一次，其申请的相关存储资源在后续核心功能模块调用中进行复用。

4.根据权利要求1所述的一种高吞吐量大带宽的通用信道化GPU算法，其特征在于，所述参数配置模块在计算中动态添加某个窄带信号参数的场景中，采用锁同步的机制，只有在参数配置模块设置完毕并解锁后，核心功能模块才能安全地访问动态更新后的信号参数信息。

5.根据权利要求1所述的一种高吞吐量大带宽的通用信道化GPU算法，其特征在于，所述核心功能模块在下次调用处理前将本次输出结果进行计算处理或者拷贝存储。

说明书

一种高吞吐量大带宽的通用信道化GPU算法

技术领域

[0001] 本发明涉及通信技术领域，具体涉及一种高吞吐量大带宽的通用信道化GPU算法。

背景技术

[0002] 信道化技术是数字宽带接收机中的关键技术之一，其用于提取包含在数字宽带采样信号中单个或多个相互独立的目标窄带信号，其中涉及数字下变频（DDC）、滤波（Filter）、采样率变换（SRC）等处理流程。

[0003] 随着现代通信技术的不断进步，在相关信道化处理业务中的数字宽带采样信号的输入带宽越来越大、处理的目标窄带信号越来越多，使得早期基于专用硬件模块实现的信道化平台方案所耗费的资源越来越难接受。而且，现代软件无线电中对功能设计灵活可变的需求促使信道化处理趋向于利用程序算法实现，不仅减小了系统设计复杂性，也使得后期功能调整更加灵活、方便。

[0004] 大规模科学研究和工程应用对计算性能的需求不断增加，过去十多年计算机技术历经“多核”体系发展到如今的“众核”体系，尤其以CPU + GPU为代表的异构并行加速平台得到了越来越广泛的应用。

[0005] GPU拥有成百上千的处理核心，访存带宽也远高于CPU，其理论峰值浮点计算性能相比CPU有着数量级的提升，随着各种异构编程语言的出现，在GPU上进行通用并行计算方面的编程技术逐渐普及到开发者大众中，其已经从专用的图像加速处理器逐渐演变成为如今的通用计算加速器，其独有的体系架构针对类似SIMD、SPMD风格的数据大规模并行计算场景进行了专门优化，尤其是NVIDIA公司推出的CUDA计算平台和编程模型，给计算机业界带来一场疾风骤雨般的“算力革命”，其推动了HPC、AI、自动驾驶、金融、医疗等众多行业的加速发展和生态繁荣。

[0006] 现代无线电通信、数字信号处理越来越朝着高、精、尖的方向前行，数字宽带接收机前端所面临的“算力瓶颈”越来越明显，因此开发信道化GPU算法有着极其重要的现实意义。

[0007] 现有技术中存在以下缺陷：

[0008] 在传统CPU中（单核、多核）实现的信道化解决方案虽然具有灵活易用、开发成本低的优势，但是由于CPU本身的算力限制，耗时严重、实时性很差，无法满足重度计算负载情形下的信道化业务；

[0009] 基于相关专用硬件（如ASIC、FPGA）实现的信道化平台解决方案虽然能满足实时性的性能要求，但是开发成本高、周期长、资源开销大，往往不能满足实际应用场景中苛刻的条件限制；

[0010] 随着现在通信应用处理系统集成功能越来越丰富，业务灵活多变，需要频繁对相关的功能模块做调整，因而基于专用硬件机制的方案不能满足这种需求。

发明内容

[0011] 针对上述问题，本发明提供一种高吞吐量大带宽的通用信道化GPU算法，基于CUDA计算平台、编程模型实现，通过在系统平台中搭载NVIDIA公司的GPU产品，为相关信道化任务负载处理提供辅助加速处理。

[0012] 本发明采用下述的技术方案：

[0013] 一种高吞吐量大带宽的通用信道化GPU算法，包括以下步骤：

[0014] S1：利用计算初始化模块设定支持的并行最大信道数量、单次处理的信号长度等参数值，然后依据这些参数值预先申请信道化流程处理中所需的输入、输出及中间变量的存储空间，所述存储空间包括主机内存区域和设备内存区域；

[0015] S2：利用参数配置模块设置目标窄带信号相关参数并根据这些相关参数将其与计算初始化模块中分配的存储资源进行配置处理，所述相关参数包括中心频率、带宽、输出采样率、目标增益，所述配置处理包括计算前批量配置模式和计算中动态配置模式；

[0016] S3：利用核心功能模块执行实际的信道化计算，根据步骤S2给各路窄带信号分配对应的计算资源并且开启信道化流程内部各子功能模块的执行处理，所述计算资源包括CUDA中的线程网格Grid以及线程块Block，所述步骤S3子功能模块包括混频子模块、半带滤波子模块、重采样处理子模块和低通滤波子模块，具体操作步骤：

[0017] S301：将宽带输入信号数据从主机内存缓冲区传输到对应的设备内存缓冲区，为了加快数据传输效率，主机内存缓冲区使用页锁定内存；

[0018] S302：将步骤301中设备内存缓冲区的信号数据依据各路信号的混频因子进行外差混频操作，为了提高性能，各路信号对应混频因子在参数配置模块设置每路信号信息时提前计算；

[0019] S303：将步骤S302中各路信号对应外差混频计算中间结果作为输入进行半带滤波处理；

[0020] S304：将步骤S303中各路信号对应半带滤波处理中间结果作为输入进行重采样处理；

[0021] S305：将步骤S304中各路信号对应重采样处理中间结果作为输入进行低通滤波处理；

[0022] S306：将步骤S305中各路信号对应低通滤波处理结果由对应设备内存缓冲区传输到主机内存缓冲区并作为最终的信道化处理结果。

[0023] 优选的，所述步骤S1存储空间的组织形式为：所有窄带信号的输入、输出以及中间结果分配一大段连续的存储空间，其中每路信号依次占用dataStep大小的间隔段落，并确保信道化业务流程中各功能模块处理时对应的输入信号长度countIn、输出信号长度countOut不会超过dataStep大小，从而保证算法内部在数据并行处理时，各路窄带信号之间不会产生数据冲突、污染。

[0024] 优选的，所述计算初始化模块只需根据设定参数执行一次，其申请的相关存储资源在后续核心功能模块调用中进行复用。

[0025] 优选的，所述参数配置模块在计算中动态添加某个窄带信号参数的场景中，采用锁同步的机制，只有在参数配置模块设置完毕并解锁后，核心功能模块才能安全地访问动态更新后的信号参数信息。

[0026] 优选的，所述核心功能模块每次调用内存缓冲区的数据都会被重写覆盖，因而需要在下次调用核心功能模块处理前将本次输出结果进行其他后续计算处理或者拷贝存储。

[0027] 本发明的有益效果是：通过利用GPU强大的并行数据处理能力，使得在高吞吐量、大带宽、目标“窄带”信号总数较多的信道化处理情形中仍能保持较高的计算性能，满足相关应用场景的实时性处理指标。另外，由于CUDA编程模型的兼容性和可扩展性，本发明在支持CUDA的多代GPU硬件平台上均能部署配置，而且算法性能可以自适应匹配GPU硬件本身的计算能力，这样就能支撑信道化场景的灵活配置，即简单场景配置低成本的GPU硬件、复杂场景配置更高端的GPU硬件。

附图说明

[0028] 为了更清楚地说明本发明实施例的技术方案，下面将对实施例的附图作简单地介绍，显而易见地，下面描述中的附图仅仅涉及本发明的一些实施例，而非对本发明的限制。

[0029] 图1为本发明系统结构框图；

[0030] 图2为本发明中计算初始化模块内部分配的设备存储空间组织形式；

[0031] 图3为本发明中参数配置模块内部信号与对应存储资源配置形式；

[0032] 图4为本发明中核心功能模块内部线程计算资源组织形式；

[0033] 图5为本发明中核心功能模块内部各路信号处理流程框图。

具体实施方式

[0034] 为使本发明实施例的目的、技术方案和优点更加清楚，下面将结合本发明实施例的附图，对本发明实施例的技术方案进行清楚、完整地描述。显然，所描述的实施例是本发明的一部分实施例，而不是全部的实施例。基于所描述的本发明的实施例，本领域普通技术人员在无需创造性劳动的前提下所获得的所有其他实施例，都属于本发明保护的范围。

[0035] 除非另外定义，本公开使用的技术术语或者科学术语应当为本公开所属领域内具有一般技能的人士所理解的通常意义。本公开中使用的“包括”或者“包含”等类似的词语意指出现该词前面的元件或者物件涵盖出现在该词后面列举的元件或者物件及其等同，而不排除其他元件或者物件。“上”、“下”、“左”、“右”等仅用于表示相对位置关系，当被描述对象的绝对位置改变后，则该相对位置关系也可能相应地改变。

[0036] 下面结合附图和实施例对本发明进一步说明。

[0037] 如图1所示，一种高吞吐量大带宽的通用信道化GPU算法，基于CUDA计算平台、编程模型实现，通过在系统平台中搭载NVIDIA公司的GPU产品，为相关信道化任务负载处理提供辅助加速处理。本实施例通过GTX1660、RTX3060Ti两个执行平台进行以下步骤：

[0038] S1：利用计算初始化模块设定支持的并行最大信道数量、单次处理的信号长度等参数值，然后依据这些参数值预先申请信道化流程处理中所需的输入、输出及中间变量的存储空间，所述存储空间包括主机内存区域和设备内存区域，存储空间的组织形式如图2中所示：所有窄带信号的输入、输出以及中间结果分配一大段连续的存储空间，其中每路信号依次占用dataStep大小的间隔段落，并确保信道化业务流程中各功能模块处理时对应的输入信号长度countIn、输出信号长度countOut不会超过dataStep大小，从而保证算法内部在数据并行处理时，各路窄带信号之间不会产生数据冲突、污染。

[0039] S2：利用参数配置模块设置目标窄带信号相关参数并根据这些相关参数将其与计算初始化模块中分配的存储资源进行配置处理，所述相关参数包括中心频率、带宽、输出采样率、目标增益，所述配置处理包括计算前批量配置模式和计算中动态配置模式如图3所示：

[0040] 所述计算前批量配置是指宽带输入信号数据进行信道化流程处理前已知所有目标窄带信号的参数信息并提前设置，在开始调用核心功能模块进行DDC流程业务处理前，提前配置好所有目标窄带信号的参数，根据批量配置时添加的先后顺序，各路信号与对应的存储空间一一对应；

[0041] 所述计算中动态配置是指宽带输入信号数据流正在执行信道化流程处理，但是需要新增或者删除窄带信号并即时生效的一种参数配置模式，包括计算中动态删除、计算中动态添加，从图中可见在该模式下，添加的信号有全局唯一的编号，其与存储空间区域标识将建立一个动态映射的关系，在动态删除时，根据所删除信号的编号查找与其映射的存储区域编号，并将其置为空闲状态，在之后调用核心功能模块时该路信号将不会被分配相应的线程计算资源。当删除的信号较多时，模块会对这些空闲存储区域按标识进行升序排列管理，后期动态添加目标窄带信号时，则优先选取队列头部的空闲存储区域标识与该信号建立映射关系，同时取消其空闲状态标记，这样之后调用核心功能模块时，该路信号将会被分配相应的线程计算资源，开启对应的DDC业务流程处理。

[0042] S3：利用核心功能模块执行实际的信道化计算，根据步骤S2给各路窄带信号分配对应的计算资源并且开启信道化流程内部各子功能模块的执行处理，所述计算资源包括CUDA中的线程网格Grid以及线程块Block，如图4所示，将多路窄带信号的信道化并行处理组织成为二维线程网格形式，另外考虑处理的是一维信号，因而线程块也选择为一维形式。二维线程网格中y方向索引(blockIdx.y)相同的所有线程块负责相应存储空间区域标识（非“空闲”状态）所映射信号的信道化处理任务，x方向不同索引（blockIdx.x）的众多线程块则负责实现对不同信号时段的数据点进行处理，通过这种多层级的数据并行，利用CUDA中海量线程机制能够充分利用GPU硬件中成百上千的CUDA cores核心的强大并行计算能力。

[0043] 所述步骤S3子功能模块如图5所示，包括混频子模块、半带滤波子模块、重采样处理子模块和低通滤波子模块，具体操作步骤：

[0044] S301：将宽带输入信号数据从主机内存缓冲区传输到对应的设备内存缓冲区，为了加快数据传输效率，主机内存缓冲区使用页锁定内存。页锁定内存是由CUDA函数cudaHostAlloc在主机内存上分配的，页锁定内存的重要属性是主机的操作系统将不会对这块内存进行分页和交换操作，确保该内存始终驻留在物理内存中；

[0045] S302：将步骤301中设备内存缓冲区的信号数据依据各路信号的混频因子进行外差混频操作，为了提高性能，各路信号对应混频因子在参数配置模块设置每路信号信息时提前计算，经过混频处理，各路目标窄带信号对应频谱结构被搬移到基带，即信号中心频率变为“0”；

[0046] S303：将步骤S302中各路信号对应外差混频计算中间结果作为输入进行半带滤波处理，每路信号根据自身的输出采样率与原始输入采样率的关系将配置若干次半带滤波循环，由于半带滤波系数相对固定，采用CUDA编程模型中“常量内存”进行存储，减少kernel核函数计算时过多的全局内存访问操作(CUDA模型中的长延时、高开销操作)，另外，为了提高性能，滤波处理中对应的线性卷积操作将按照线程块的x方向索引组织形式进行分段处理，对应各分段的数据可采用CUDA编程模型中“共享内存”进行存储，同样避免过多的全局内存访问操作，经过半带滤波处理，已经滤除了大部分“干扰”信号成分，而且结果信号的实际采样率也非常接近目标输出采样率参数；

[0047] S304：将步骤S303中各路信号对应半带滤波处理中间结果作为输入进行重采样处理，重采样处理子模块内部根据预先计算好的因子对其进行内插/抽样处理，经过重采样处理后，各路对应输出的结果信号满足输出采样率；

[0048] S305：将步骤S304中各路信号对应重采样处理中间结果作为输入进行低通滤波处理，低通滤波子模块根据预先设置参数设计的FIR低通滤波器系数对其进行卷积操作，采用“共享内存”机制进行性能优化处理，经过低通滤波处理后，各路输出信号结果中已经将对应目标窄带信号之外的其他所有信号频谱成分滤除，只包含目标窄带信号自身完整的频谱结构；

[0049] S306：将步骤S305中各路信号对应低通滤波处理结果由对应设备内存缓冲区传输到主机内存缓冲区并作为最终的信道化处理结果，由于各路目标窄带输出采样率参数不同，因而最终各路输出结果长度大小不一，如果直接按照图2中所示将整片内存区间从设备传输到主机，必然有大量的无效数据被传输，因此为了提高数据处理结果传输效率，核心模块内部会做额外处理，将各路信道化处理结果重新组织成首尾衔接的连续存储形式，然后再传回主机内存缓冲区，主调方根据的结果长度、数据首地址偏移等回传参数到对应主机内存缓冲区索引、访问各路信号对应的最终信道化处理结果。

[0050] 所述计算初始化模块只需根据设定参数执行一次，其申请的相关存储资源在后续核心功能模块调用中进行复用。

[0051] 所述参数配置模块在计算中动态添加某个窄带信号参数的场景中，采用锁同步的机制，只有在参数配置模块设置完毕并解锁后，核心功能模块才能安全地访问动态更新后的信号参数信息。

[0052] 所述核心功能模块每次调用内存缓冲区的数据都会被重写覆盖，因而需要在下次调用核心功能模块处理前将本次输出结果进行其他后续计算处理或者拷贝存储。

[0053] GTX1660、RTX3060Ti两个执行平台执行结果如下表所示：

[0054]

[0055] 通过上表可知，在不同处理路数场景下，本发明执行信道化任务负载所达性能指标均较为理想，而且如前所述算法的性能会随着GPU硬件平台计算能力的提升而增强。

[0056] 以上所述，仅是本发明的较佳实施例而已，并非对本发明作任何形式上的限制，虽然本发明已以较佳实施例揭露如上，然而并非用以限定本发明，任何熟悉本专业的技术人员，在不脱离本发明技术方案范围内，当可利用上述揭示的技术内容做出些许更动或修饰为等同变化的等效实施例，但凡是未脱离本发明技术方案的内容，依据本发明的技术实质对以上实施例所做的任何简单修改、等同变化与修饰，均仍属于本发明技术方案的范围内。