用CUDA加速多波束测深系统波束形成

摘 要

通常信号处理中的波束形成计算都是在DSP或者FPGA上实现的，普通的PC机上无法实现实时波束形成。本文在CUDA架构下用GPU加速多波束测深系统的波束形成计算过程，在移动工作站上实现了测深数据的实时波束形成处理，具有开发周期短、成本低、计算能力可扩充的优点。

【关键词】CUDA多波束 波束形成 实时处理

1 引言

多波束测深系统是一种高效的水底地形测量系统。设备通过向水底发射一次信号，接收信号的回波来进行处理，得到测量平台两侧一条带上几百个采样点的深度数据。波束形成是多波束测深系统信号处理中一个非常重要的任务，往往耗费大量的计算资源，通常对这类数据的处理采用DSP来实现。然而这种DSP实时处理的方式具有投入大，研发周期长，板卡多，升级困难的缺点。由英伟达公司提出的CUDA（Compute Unified Device Architecture）计算架构，能使用标准的C来进行GPU编程，对大规模数据并行计算，不仅运行速度快，并行化程度高，还具有开发周期短、无需硬件设计调试、便于升级的优点。

本文通过CUDA技术对多波束测深系统的波束形成进行并行优化处理，大大提高了信号处理速度，从而在移动工作站上实现了多波束测深数据的实时波束形成处理。

2 多波束测深系统波束形成

考虑多波束测深系统中基阵接收位于远场的情况，即信号波前到达基阵时可看作平面波。在远场情况下对基带解调后的信号，采用基带延迟加相移波束形成。对于波束形成需要的时延补偿，其整数点采样部分，用时延实现，剩余的小数采样部分，用相移实现。

在远场平面波的假设条件下，以任意阵型M元阵，入射单频脉冲信号为例。假设第i号阵元上的接收信号经过基带解调后为xi （t），其中i=1，2，…，M，它的希尔波特变换（Hilbert）（t）。设信号原始频率为f0，接收信号经过基带解调后频率为fc，采样频率为fs。若要形成θ方向的预成波束，根据常规波束形成的时延计算方法，可以计算出，第i号阵元需要补偿的时延，假定为τi （θ），其中i=1，2，…，M。采用基带信号的延迟加相移波束形成方法，则θ方向的预成波束输出为：

（1）

其中： （2）

（3）

ai为阵元的权值，ni（θ）为需要延时的整数采样点，τi（θ）为剩下的小数点采样对应的时延。

3 CUDA实现

3.1 CUDA编程模型

CUDA采用了“单指令多线程”（Single Instruction Multiple Thread，SIMT）的执行模型，只需一条指令，GPU的内部机制会保证其所有配置的线程都会执行同样的运算。CUDA的线程模型分为三层，从上到下依次为Grid，Block，Thread。Thread（线程）是最小的单元，每个Thread可以映射到一个流处理器上，通过合理安排threads的工作可以提高程序的并行化，提升程序性能。kernel（内核函数）是一个只能运行在GPU上的函数，其调用格式为__global__ kernel_function（param1，…），其中num_blocks和num_threads是dim3整形向量类型的变量，分别表示网格和块的维度，两者决定了内核函数执行时所需要的线程数。

3.2 CUDA程序波束形成设计

对于N通道的基带解调数据，数据以N*M的2维数组排列，X维表示时间序列，Y维表示通道号，对 方向进行波束形成：

定义dim3 threads_rect（x，y），表示一个块内的线程数，定义dim3 blocks_rect（x，y，z），表示线程块的数量。其中threads_rect.y* blocks_rect.y=N， blocks_rect.x=（M+threads_rect.x-1）/threads_rect.x，一个波束时z=1，对应的CUDA内部kernel分配图如图1所示。

其中X维的索引IDx=（blockIdx.x*blockDim.x）+threadIdx.x；

Y维的索引IDy=（blockIdx.y*blockDim.y）+threadIdx.y。

3.3 性能优化

3.3.1 线程数的选择

表1显示了在不同计算能力的设备上，每个线程块开启不同线程数时的设备利用率。

通过表1可以看出，当线程数为128～512之间时，在不同的设备上都有比较好的利用率。但是高利用率并不一定意味着高性能，CUDA是以线程束（Warp）的方式访问全局内存，线程束的大小通常为32，每次以Half-Warp的16个独立线程对全局内存中的一段连续数据进行访问，可以最大化IO的效率。因此我们线程块开启的线程数设置为16x16，对应threads_rect.x=16， threads_rect.y=16。

3.3.2 常量内存

在CUDA中，常量内存其实只是全局内存的一种虚拟地址形式，但是它具有高速缓存，访问速度要比没有缓存的全局存储器快，因此把系统中所有只读常量保存在常量内存中可以提高性能。

3.3.3 锁页内存

为了对某一数据集进行计算，需要把数据从主机传输到GPU设备上，计算完成后再把数据结果传输回主机，主机与GPU设备的数据传输是依靠PCI-E总线来实现的。PCI-E传输实际上只能使用基于DMA的传输方式，在不使用锁页内存时，驱动程序会在后台先分配一块锁页内存，实行一个从常规内存到锁页内存的主机端复制操作，这些时间额外消耗了宝贵的CPU周期。直接使用cudaHostAlloc或cudaMallocHost分配锁页内存，可以增加传输带宽，提高通信效率。

3.3.4 异步传输

在主机与GPU设备的数据传输过程中，GPU的会出现大量空闲，这样导致GPU执行效率不高。因此需要通过CUDA流和cudaMemcpyAsync来使CPU和GPU都尽量处于满负荷状态，提高GPU使用效率，掩盖传输带来的延迟。

4 性能测试

4.1 实验环境

实验在移动工作站上进行，CPU为Intel Core i7-4600M，主频2.9GHz，2核4线程，内存16GB，GPU为GeForce GT 745M，384 个流处理器，2GB显存；软件支持环境为：Win7-64位操作系统、MicrosoftVS 2010、NVIDIA GPU Computing SDK 7.5。

4.2 GPU性能测试

选取128通道一个量程扫描周期的数据进行256个波束形成速度比较。

不同通道点数的CPU和GPU加速度对比，左边为消耗的时间，右边为加速比。

4.3 实验分析

从表3和图2可以看出，相对于CPU，GPU对多波束测深系统的波束形成计算具有明显的加速作用，并且随着每通道点数的增加，CPU的计算速度大幅下降，而GPU的耗时却平稳增长，到每通道8192点数据时，GPU对CPU的加速比达到了接近20X。

5 总结

本文介绍了一种将CUDA技术应用到多波束测深系统波束形成处理的方法，在移动工作站上实现了信号实时波束形成的并行化加速，克服了DSP实时处理系统的缺点，这对将来构建低成本、开发方便、便携式的多波束测深系统具有非常重要的意义。

参考文献

[1]Shane Cook，CUDA并行程序设计[M].北京：机械工业出版社，2014.

[2]沈玉琳.通用 GPU 计算技术在高性能计算平台上的应用研究[D].兰州大学，2012.

[3]张舒等.GPU高性能运算之CUDA[M].北京：中国水利水电出版社，2009.

作者单位

杭州应用声学研究所 浙江省杭州市 310000