具体实施方式

下面将结合本发明实施例中的附图，对本发明实施例中的技术方案进行清楚、完整地描述，显然，所描述的实施例仅仅是本发明一部分实施例，而不是全部的实施例。基于本发明中的实施例，本领域普通技术人员在没有做出创造性劳动前提下所获得的所有其他实施例，都属于本发明保护的范围。

参照图2，为基于多服务器多GPU的星载SAR成像实时处理方法的整体连接图所示，分为三个模块：多CPU主机端间的免密通信模块、星载SAR数据预处理模块以及GPU设备端BP成像模块，具体包括以下功能：

1)多CPU主机端间的免密通信模块：多CPU主机端间的免密通信模块通过MPI实现CPU主机端的免密通信。

2)星载SAR数据预处理模块：星载SAR数据预处理模块主要包括CPU主机端读取雷达回波信号数据参数与GPU设备端信息、GPU设备端从CPU主机端读取雷达回波信号数据、GPU设备端生成复数回波以及GPU设备端的距离向脉冲压缩功能。

3)GPU设备端BP成像模块：GPU设备端BP成像模块完全在GPU设备端进行，主要包括：(1)构建针对雷达以及目标的坐标系，即雷达目标位置的计算，以便于针对目标和雷达实现投影；(2)选取BP成像的有效成像区域，以便于后续确定回波相干积累的区域；(3)划分地面网格，即对所有的网格点执行同样的投影和积累操作，实现图像聚焦；(4)划分BP成像有效区域，包括起始、雷达位置有效区域的选取，以及有效孔径的选取；(5)利用sinc插值进行插值计算，以确定回波脉冲对各个网格点的贡献；(6)进行距离向相位补偿，实现BP成像。

在本发明中，一个CPU主机端可以控制多个GPU设备端，一个CPU主机端可以读取多个GPU设备端信息，一个CPU主机端具体控制GPU设备端的数目取决于CPU主机端主机上接口数目。

每个GPU设备端根据对应的CPU主机端为其分配的内存按照方位向分块从对应的CPU主机端读取相应数据块大小的雷达回波信号数据(距离向优先存储)以及雷达后向投影成像BP算法参数信息，将仿真参数从对应的CPU主机端读取至GPU设备端的常量存储区中，将自适应分块雷达回波信号数据从对应的CPU主机端读取至GPU设备端的全局内存中；通过多个CPU主机端控制的多个GPU设备端，在每个GPU设备端上完成距离向脉冲压缩、进行BP成像所需的参数计算、利用插值实现回波脉冲的升采样，CPU主机端将雷达后向投影成像BP算法中的参数计算结果以及升采样后的回波脉冲分给不同的GPU设备端，最终通过GPU设备端实现BP成像。

为了满足较快的成像处理要求，考虑到星载SAR处理的数据量大，因此使用基于多服务器多GPU的星载SAR成像实时处理方法，参照图3，为基于多CPU和多GPU的异构平台BP成像的流程示意图，其包括以下具体步骤：

步骤1：在任意一个服务器CPU主机端上读取星载SAR雷达回波数据，利用多个服务器CPU主机端之间的免密通信，将雷达回波信号数据划分到多个服务器CPU主机端上。

具体步骤为，通过MPI(Massage Passing Interface)实现任意一个服务器CPU主机端与其他服务器CPU主机端的免密通信。(1)通过ssh(Secure Shell)网络安全协议登陆各个CPU主机端；(2)在各个服务器CPU主机端生成相同的雷达回波信号数据存储路径；(3)将相同的执行配置文件放在服务器CPU主机端相同存储路径的同一个目录下；(4)在不同的服务器CPU主机端执行(3)中相同的执行配置文件时同时启动各个服务器CPU主机端，获取更多的服务器CPU主机端线程；(5)利用OpenMP将雷达后向投影成像BP算法中的任务分配给不同的服务器CPU主机端线程，实现任务级并行。

其中，OpenMP是一种用于共享内存并行系统的多线程程序设计方案，OpenMP提供了对并行算法的高层抽象描述，特别适合在多核CPU机器上的并行程序设计。编译器根据程序中添加的pragma指令，自动将程序并行处理，使用OpenMP降低了并行编程的难度和复杂度。当编译器不支持OpenMP时，程序会退化成普通(串行)程序。程序中已有的OpenMP指令不会影响程序的正常编译运行。

具体的OpenMP工作模型如图4，为OpenMP工作模型示意图所示，是一种用于共享内存并行系统的多线程程序设计方案，支持的编程语言包括C、C++和Fortran。开始的时候只存在一个主线程，当需要进行并行计算的时候，派生出若干个分支线程来执行并行任务。当并行代码执行完成之后，分支线程会合，并把控制流程交给单独的主线程。OpenMP提供了对并行算法的高层抽象描述，特别适合在多核CPU机器上的并行程序设计。编译器根据程序中添加的pragma指令，自动将程序并行处理，使用OpenMP降低了并行编程的难度和复杂度。当编译器不支持OpenMP时，程序会退化成普通(串行)程序。程序中已有的OpenMP指令不会影响程序的正常编译运行。OpenMP采用fork-join的执行模式。

通过上述MPI技术即可完成各个服务器CPU主机端之间的免密通信配置，实现各个服务器CPU主机端之间的免密通信，使得服务器CPU主机端可以调用更多的线程控制GPU设备端，提高服务器CPU主机端与GPU设备端的数据传输效率。

步骤2：在每个服务器CPU主机端读取相对应的雷达回波信号数据参数、相对应多个GPU设备端信息，并设置该服务器CPU主机端线程数。

每个服务器CPU主机端读取的相对应的每个GPU设备端信息包括：显卡型号、设备计算能力、全局内存总量、设备的网格块线程划分的上下限、以及需要处理的星载SAR雷达回波数据的距离向点数(假设需要处理的星载SAR回波点数为Nr*Na，分别代表回波数据的距离向Nr和方位向Na，其中*代表乘运算)。

服务器CPU主机端根据读取的GPU设备端信息设置与GPU设备端数目相同的主机端线程，即一个服务器CPU主机端线程控制一个GPU设备端，通过OpenMP将雷达后向投影成像BP算法中的雷达回波信号数据处理任务与参数计算任务分配给不同的CPU主机端线程，实现任务级并行。

具体的，根据雷达后向投影成像算法的性能要求，将雷达后向投影成像BP算法相关参数封装为类，根据具体的成像需求更改这部分成像相关参数。

步骤3：每个GPU设备端从其对应的服务器CPU主机端读取为其针对方位向划分的雷达回波信号数据，以及该GPU设备端的分配内存。

每个GPU设备端读取其对应服务器CPU主机端的雷达方位向分块的雷达回波信号数据以及雷达后向投影成像算法需要的参数信息。

根据步骤2，服务器CPU在主机端基于获取的相对应的各个GPU计算能力、全局内存总量，以及需要处理的星载SAR回波数据的距离向点数(假设需要处理的星载SAR回波点数为Nr*Na，分别代表回波数据的距离向Nr和方位向Na)，进行针对方位向的星载SAR雷达回波数据划分，以及设备端内存分配。

其中，在进行星载SAR雷达回波数据数据划分时应注意，在后续步骤实现距离向的脉冲压缩操作前需要对雷达回波信号数据做距离向傅里叶变化、距离向傅里叶变化会额外开辟GPU设备端内存，会导致GPU设备端真正处理进行傅里叶变化的数据空间临时扩大为原来仅为雷达回波信号数据分配空间的三倍，因此，在CPU主机端进行雷达回波信号数据的方位向数据分块时需要在GPU设备端预留一定大小的GPU内存空间。

具体操作为：服务器CPU主机端读取对应各个GPU设备信息，根据GPU设备端信息全局内存总量，用GPU设备端信息全局内存总量乘以0.3设定真正可用的GPU内存总量。具体的方位向分块如图5所示，图5为雷达回波信号数据根据数据量大小的方位向分块示意图。由于回波数据是二维的，分为距离向和方位向，在图中从上至下为距离向，从左至右为方位向，在该发明中，传入的雷达回波数据是距离向优先存储的。为避免内存的跳跃读取从而增加耗时，因此，进行针对方位向的雷达回波数据分块假设将回波数据按照方位向分为n块，表示为blk_num，具体分块数取决于雷达回波数据大小、后续步骤进行的操作与GPU显存大小。

CPU主机端根据GPU设备端数目gpu_num开启主机端线程数cpu_trd_num，如图5中GPU设备端数目等于4，每个主机端线程负责n/4块雷达回波数据。根据步骤2读取到的GPU设备端信息，CPU主机端自适应的根据GPU设备端的数目，完成针对雷达回波信号数据方位向的自适应分块，并将不同的分块分配给CPU主机端的不同线程，由CPU主机端的各个线程控制传入不同GPU设备端中。

步骤4：每个GPU设备端接入相对应服务器CPU主机端传入雷达回波信号的复数回波。

每个服务器CPU主机端根据接收到的雷达回波信号数据的不同类型进行数据处理，雷达回波信号数据传入服务器CPU主机端有两种情况，一种是实部和虚部分离传入，一种是完整的复数信号传入。

(1)若雷达回波信号是分为实部和虚部分离传入，服务器CPU主机端根据传入星载SAR雷达回波信号数据的大小自适应的在GPU设备端开辟内存，并将雷达回波信号的实部和虚部分别传输到GPU设备端，然后在GPU设备端利用核函数GnrEcho对雷达回波信号在GPU设备端完成实部和虚部拼接；(2)若雷达回波信号数据是完整的复数回波信号，则该步骤略过。

步骤5：每个GPU设备端实现复数回波的距离向的脉冲压缩。

GPU设备端调用英伟达CUDA(Compute Unified Device Architecture)库中快速傅里叶FFT变换相关函数，对GPU设备端已经读取到复数回波数据做距离向快速傅里叶FFT变换，将雷达回波复数数据变化到方位时域对应的距离频域，在距离频域的雷达回波信号数据乘以匹配滤波器，以实现雷达回波信号脉冲压缩。

具体步骤是：基于步骤3所做的雷达回波信号数据划分，在服务器CPU主机端利用openMP实现CPU的多线程控制，在服务器CPU主机端线程内进一步控制多个GPU设备端进行相关并行计算：(1)在GPU设备端调用cuda自带的库函数cuFFT，对从服务器CPU主机端读取的已经完成方位向分块的雷达回波信号数据做距离向FFT，将回波信号变换到距离频域，为后续脉冲压缩准备；(2)根据雷达仿真参数，在方位时域、距离频域对回波数据做匹配滤波，即在每一个距离向nrn(为距离向Nr的索引)生成一个匹配滤波器，公式为exp(j*pi*f*f/Bp*Tp)，其中f＝nrn/Nr*fr；其中f为距离向采样频率、Bp为信号带宽、Tp为脉冲重复间隔、fr为雷达脉冲频率；(3)最后把经过步骤(1)操作的雷达回波数据与匹配滤波器进行复数相乘完成距离向脉压。在GPU设备端利用核函数Pul_press实现雷达回波信号数据的距离向脉压。

步骤6：在每个服务器CPU主机端所控制的GPU设备端上实现BP成像。

在每个服务器CPU主机端上所控制的GPU设备端上，进行BP成像所需的参数计算，其中如图6，为关于BP成像的具体操作流程示意图所示。

BP成像过程均在GPU设备端完成：(1)首先构建针对雷达以及目标的坐标系，即雷达目标位置的计算，以便于针对目标和雷达实现投影；(2)选取BP成像的有效成像区域，以便于后续确定回波相干积累的区域；(3)划分地面网格，对所有的网格点执行同样的投影和积累操作，实现图像聚焦；(4)选取雷达位置有效区域；(5)选取有效孔径；(6)通过slnc插值计算回波脉冲对网格点的贡献在BP成像前利用插值实现回波脉冲的升采样；(7)将服务器CPU主机端的参数计算结果以及升采样后的回波脉冲分给不同的GPU设备端，通过GPU设备端的反复执行投影与积累，实现BP成像。

具体的，步骤6.1，BP成像参数及设置。

雷达的三维位置在笛卡尔坐标系中的关系如图7，为条带式SAR雷达获取数据几何模型示意图中所示，在图6，为关于BP成像的具体操作流程示意图所示的Kernel3核内计算得到雷达的三维位置信息。

在为雷达位置建立三维坐标时，有两种数据组织方式，不同的数据组织方式会对核性能有不同的影响，如图9(a)，为AoS和SoA两种数据结构体组织方式，对应距离向和方位向优先存储方式示意图所示：第一种方式为数据结构体(AoS)，对应着雷达回波数据的方位向优先存储；第二种方式为结构体数组(SoA)，对应着雷达回波数据的距离向优先存储。AoS数据组织方式会导致在访问雷达坐标位置x数组、y数组、z数组过程中出现非对齐非连续访存，严重影响内存的读写速度，降低核性能。而使用SoA数据组织方式有效避免AoS的弊端，在存储雷达位置坐标的同时可以进行对齐连续的内存访问，使得核性能更加高效。

如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示，GPU核函数内部存在三种索引，第一种是索引线程的横向分布，表示目标场景距离向的目标点，记作idx，计算方式为nrn＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应图9(b)中各个线程块，blockDim.x对应图9(b)中各个线程块所对应的维度，threadIdx.x则对应图9(b)中各个维度线程块中的线程束索引，线程束数量为16的倍数即可；第二种是索引线程的纵向分布，表示目标场景方位向分块的目标点，记作idy，计算方式为nan＝blockIdx.y*blockDim.y+threadIdx.y；其中blockIdx.y对应图9(c)，为核函数Kernel9的线程分配示意图中各个线程块，blockDim.y对应图9(c)中各个线程块所对应的维度，threadIdx.y则对应图9(c)中各个维度线程块中的线程束索引，线程束数量为16的倍数即可；第三种是通过索引方位向分块之后每块的起始点坐标确定分块后各个目标在整体雷达回波数据中的偏移，记作Echo_offset，计算方式为Echo_offset＝Na_offset*len[i]，其中len[i]表示雷达信号回波目标区域的第几个分块,Na_offset表示为每一个分块的方位向偏移，通过与len[i]相乘得到各个不同分块的方位向偏移。

由线程的索引确定雷达目标点的三维坐标，将计算结果按行存储在GPU显存中开辟好的数组里。调用核函数计算时按行调用，形成对齐连续的访存，减少了跨内存访问，保证了处理距离向优先存储雷达回波数据的核函数有良好的核性能。根据线程索引确定每个雷达回波数据方位向分块的起始点坐标。

在程序中将具体的BP成像的相关参数封装为类，根据具体的成像需求更改这部分成像相关参数。

步骤6.2，完成BP成像。

分析星载SAR雷达回波数据在距离向的脉冲压缩、BP成像所需参数，利用插值实现回波脉冲的升采样，实现BP成像处理前后的依赖性，对有数据依赖性且for循环体量较大的操作流程选择通过CPU主机端使用openMP进行多线程加速；反之对没有数据依赖性且循环体量较大的操作流程选择在GPU设备端使用CUDA进行加速。

其中，常见的插值方式有上采样后线性插值、spline插值以及sinc插值等，可综合考虑成像质量与实时性要求，选取合适的插值方式。本发明中选择sinc插值。

在BP成像前利用sinc插值实现回波脉冲的升采样。CPU主机端将各个参数计算结果以及升采样后的回波脉冲分给不同的GPU设备端，通过GPU设备端反复执行投影与积累，实现BP成像。

另外，针对BP成像以及BP成像前的参数计算过程中没有数据依赖性的操作，在GPU设备端通过以下几个核函数实现；

Kernel1：GnrEcho：如果星载SAR雷达回波数据为实部虚部分离的形式，通过该核函数实现星载SAR雷达回波数据的实部虚部结合；

Kernel2：Pul_Press：针对做完距离向傅里叶变化的雷达回波数据，复乘一个匹配滤波器，实现雷达回波数据的距离向脉冲压缩；

Kernel3：InitSensorPos：计算雷达位置信息，构建雷达与目标的坐标系；

Kernel4，Init_Sig_stav：选取起始位置、雷达位置有效区域；

Kernel5、Kernel6：Init_Grdx以及Init_Grd_dlt_y：划分地面网格；

Kernel7、Kernel8：Init_Senpos以及Init_AzSE：选取有效孔径；

Kernel9：ProcessPoint：完成对回波数据的sinc插值，实现升采样；完成雷达天线中心相位与目标瞬时斜距的计算；同时完成距离向相位补偿；并最终完成BP成像。

如图3，为基于多CPU和多GPU的异构平台BP成像的流程示意图以及图6，为关于BP成像的具体操作流程示意图所示，在BP成像阶段对BP成像进行任务划分，并分析数据依赖性，若Kernel1、Kernel2、Kernel3、Kernel4所处理的数据具有数据依赖性，可以先将其分配到的GPU设备端上通过逐个核函数依次运行；

Kernel5、Kernel6的输入数据与Kernel1、Kernel2、Kernel3、Kernel4四个核函数处理的数据有关，在Kernel1、Kernel2、Kernel3、Kernel4四个核函数运算过后将运算结果置于GPU设备端上：Kernel5与Kernel6需要等待上述核函数Kernel1-Kernel4运算完成后再进行运算，可考虑在Kernel5核函数执行前加上阻塞；

Kernel7与Kernel8需要等待Kernel3完成后再进行运算，可以将Kernel3与Kernel7、Kernel8置于同一工作流下，尽可能减少GPU设备端与CPU主机端数据传输的耗时；

完成上述核函数操作后、利用Kernel9完成BP成像。考虑到对内存的消耗和各个GPU设备端之间的可重叠计算，可以计算好基于方位向优先存储的数据偏移，包括行列的偏移，实现多GPU的并行处理，进一步提高算法效率。

步骤7：获取所有服务CPU主机端相对应的所有GPU设备端上的BP成像，并写入雷达回波数据处理文件，利用MATLAB读取处理后的雷达回波数据并根据所处理雷达回波数据绘图，分析成像结果。

下面结合仿真实验对本发明的效果做进一步说明。

(1)未采用本发明基于多服务器多GPU的星载SAR成像实时处理方法的实测数据仿真。

仿真实验参数：脉冲回波点数12288、脉冲回波周期2.5e-6s、脉冲回波带宽2.997e+9、脉冲回波频率3.5975e+5、雷达速度7400m/s。

仿真步骤为：

1)主机端回波数据读取；

2)主机端距离向脉冲压缩；

3)主机端BP成像的参数计算：包括构建针对雷达以及目标的坐标系、BP成像的有效成像区域的选取、地面网格的划分、其中BP成像有效区域的划分包括有效孔径的选取、以及起始、雷达位置有效区域的选取；

4)主机端BP成像。

如图1，为由CPU主机端串行完成SAR回波数据的流程示意图所示，在CPU主机端上串行处理雷达回波信号数据时，1)首先读取仿真参数，将SAR回波形成需要的仿真参数封装为类，程序中包含头文件直接调用即可；2)CPU主机端在星载SAR回波数据文件中读取回波数据，在回波数据中选取需要的成像区域作为成像有效区域，根据具体成像需求，确定目标点数；3)对方位向和距离向的各目标点完成计算后，实现距离向脉压；4)将回波数据导入距离频域，方位时域，进一步继续仿真参数；5)然后根据仿真参数计算在雷达运行时间内各个目标点的实时距离，并根据实时距离判断目标点是否恰好在雷达天线发射波束的照射范围之内，计算在波束照射范围内目标点的回波，在距离向进行相位补偿，方位向进行补偿相位后的目标点回波数据累加。

仿真数据结果如图10，为未采用本发明基于多服务器多GPU的星载SAR成像实时处理方法的实测数据仿真图所示。

(2)采用本发明基于多服务器多GPU的星载SAR成像实时处理方法的实测数据仿真。

本仿真中的参数设置与仿真1中的设置相同。仿真步骤为：

1)CPU主机端读取回波数据；

2)GPU设备端进行距离向脉冲压缩；

3)GPU设备端进行BP成像的参数计算：包括构建针对雷达以及目标的坐标系、BP成像的有效成像区域的选取、地面网格的划分、其中BP成像有效区域的划分包括有效孔径的选取、以及起始、雷达位置有效区域的选取；

4)GPU设备端完成BP成像。具体步骤为：(1)在主机端端读取设备端信息，主机端完成数据的自适应分块，调用核函数GnrEcho，循环遍历所有的块，循环的次数由数据大小决定；(2)完成这数据的读取后，调用核函数pul_press，完成距离向脉冲压缩，设置异步阻塞，保证下一步处理前所有的回波数据完成距离向脉压；(3)基于完成距离向脉压后的数据分块依次调用核函数InitSensorPos构建雷达与目标的坐标系；(4)核函数Init_Sig_stav完成对起始位置、雷达位置有效区域的选取；(5)通过核函数Init_Grdx以及核函数Init_Grd_Slt_y完成对地面网格的划分；(6)通过核函数Init_Senpos以及核函数Init_AzSE完成成像坐标系的构建以及对有效孔径的选取；(7)根据设备端GPU的情况，数据大小以及参数计算选择最优的并行化方式，减少数据传输的耗时，实现BP成像设备端的最优化处理结果。

下面就具体的，针对整个流程出现核函数展开介绍，重点介绍GPU设备端针对各个核函数中网格和块的划分：TrdNum和BlkNum；明确各个核函数的意义以及核函数各个接口的封装，且明确分配的GPU设备端内存大小。

GPU设备端数量为4，GPU的显存大小为32G，该配置小于需要处理的雷达回波数据量大小，考虑到在快速傅里叶FFT变换、相位补偿和插值计算时会临时计算辅助数据，辅助数据的计算占用额外的GPU内存，此时需要考虑显存的最大利用率，利用CPU主机端读取的GPU设备端信息，合理分配GPU设备端内存。

为每个核函数配置线程，实现线程级并行，在遵循一定原则的基础上保证核函数拥有最优的性能；明确核内索引内容和数据组织方式，优化核函数内部计算，使核函数性能达到最优。为每个核函数配置网格块线程，进行配置应遵循以下原则，保证核函数有较好的性能：

(1)为核函数配置的网格和块在x维度，y维度，z维度的长度不超过GPU设备端的极限，并不是配置越接近GPU设备极限性能就越好。

(2)为每个核函数配置的线程数应为32的整倍数，且最好大于4*32个线程。因为GPU在硬件流多处理器SM上处理线程是以一个线程束(warp)为单位来处理的。一个线程束由32个线程组成，将四个线程束置于硬件流多处理器SM上来并行处理，可以有效隐藏访存延迟，提升核性能。放置在硬件流多处理器SM上的线程束越多，隐藏访存延迟的效果越好。

(3)每个核函数的性能会受到配置的影响，想得到最佳的配置，需要经过多次的测试。使用不同的配置运行核函数，利用工具多次测试核函数的运行时间，可得到耗时最短时的配置。

下面，GPU设备的内存分布与网格、块、线程之间的关系如图8，为GPU内存与线程关系示意图所示：通过事先声明核函数，设置线程网格数目、线程块数和线程数。当一个线程块的网格被启动后，网格中的线程块分布在SM(Streaming Multiprocessor)中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。一个核函数中的线程块可能会被调度到不同的SM中。线程束是SM中基本的执行单元。一个线程束由多个个连续的线程组成，在一个线程束中，所有的线程同时执行。在GPU内存层次结构中，最主要的两种内存是全局内存和共享内存。全局类似于CPU的系统内存，而共享内存类似于CPU的缓存。GPU的共享内存可以由CUDAC的内核直接控制，即核函数内部声明；而GPU的全局内存通过CUDAC的数据操作指令将CPU主机端的数据搬移到GPU的全局内存中，再利用核函数各个线程访问GPU全局内存进行并行计算。在GPU上所使用的核函数，核函数功能、核函数变量和核函数分配空间如下。Kernel1核函数名称为GenEcho，功能为将读取得到的雷达回波信号数据的实部和虚部进行组合，生成完整的复数信号；块内配置线程为GenEchoThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx、idy，分别对应距离向和方位向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定，核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应GenEchoblocks.x中各个线程块索引，blockDim.x对应GenEchoblocks.x中各个线程块所对应的维度，threadIdx.x则对应GenEchoblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Nr]为距离向索引；idy索引由块的大小数目决定，核函数idy索引计算方式为idy＝blockIdx.y，其中blockIdx.y对应GenEchoblocks.y中各个线程块索引，blockDim.y对应GenEchoblocks.y中各个线程块所对应的维度，threadIdx.y则对应GenEchoblocks.y中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Na]为方位向索引,代表方位向。使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。为该Kernel1核函数分配的数组变量分别为d_Recho,d_Iecho,d_echo,Nr，其中Recho代表数据实部，Iecho代表数据虚部，Echo代表复数数据，Nr代表数据存储的距离向长度；分配的GPU显存空间为sizeof(float)*Na*Nr，sizeof(float)*Na*Nr，sizeof(c_float)*Na*Nr；其中c_float为cuda_complex.cuh定义的实部虚部为float的复数信号。

Kernel2核函数名称为pul_press，块内配置线程为pul_pressThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示所示。在核内建立线程索引，设置线程索引为idx、idy，分别对应距离向和方位向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定，核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应pul_pressblocks.x中各个线程块，blockDim.x对应pul_pressblocks.x中各个线程块所对应的维度，threadIdx.x则对应pul_pressblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Nr]；核函数idy索引计算方式idy＝blockIdx.y，其中blockIdx.y对应pul_pressblocks.y中各个线程块索引，blockDim.y对应pul_pressblocks.y中各个线程块所对应的维度，threadIdx.y则对应pul_pressblocks.y中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Na]。使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能为完成雷达回波数据在距离向脉压，在一个核函数Pul_press内生成一个temp_MatchFun，根据匹配滤波器的计算公式为exp(i*pi*f*f/Bp*Tp)，其中f＝nrn/Nr*fr，其为各个距离向的匹配滤波器，对各个距离向的匹配滤波器与各个距离向的频域信号作复数相乘，完成距离向脉压。其中，为该函数分配的数组变量是d_echo，Nr，Fr，Br，Tr，其中，echo为回波数据，Nr为方位向长度，Fr为雷达脉冲频率、Br为雷达脉冲带宽；分配的空间为sizcof(c_float)*Na*Nr，sizeof(c_float)*Nr，c_float表示复数的float信号，Nr与Na分别表示距离向和方位向大小，利用匹配滤波器的计算公式对雷达回波数据做方位向匹配滤波，完成距离向脉冲压缩。

Kernel3核函数名为InitSensorPos，块内配置线程为InitSensorPosThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx、idy，分别对应距离向和方位向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定，核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应InitSensorPosblocks.x中各个线程块索引，blockDim.x对应InitSensorPosblocks.x中各个线程块所对应的维度，threadIdx.x则对应InitSensorPosblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Nr]为距离向索引；核函数idy索引计算方式为idy＝blockIdx.y，其中blockIdx.y对应InitSensorPosblocks.y中各个线程块索引，blockDim.y对应InitSensorPosblocks.y中各个线程块所对应的维度，threadIdx.y则对应InitSensorPosblocks.y中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Na]为方位向索引,代表方位向。使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能为完成如图6，为关于BP成像的具体操作流程示意图所示的BP成像的坐标系构建。其中雷达的三维位置在笛卡尔坐标系中的关系如图7，为条带式SAR雷达获取数据几何模型示意图所示，在核内计算得到雷达的三维位置信息。各个变量的含义为：sensor_pos为坐标系保存，Vr为雷达的方位向速度，H为雷达的飞行高度，PRF为脉冲重复采样频率,为sensor_pos分配的空间大小为sizeof(double)*3，生成的sensor_pos按照顺序存储的方向分别存储瞬时斜距、各个方位向的成像大小、以及雷达飞行高度H，通过这个核函数完成基于目标和雷达的坐标系的构建。

Kernel4核函数名称为Init_Sig_staV，块内配置线程为Init_Sig_staVThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx、idy，分别对应距离向和方位向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定，核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应Init_Sig_staVPosblocks.x中各个线程块索引，blockDim.x对应Init_Sig_staVPosblocks.x中各个线程块所对应的维度，threadIdx.x则对应Init_Sig_staVPosblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Nr]为距离向索引；核函数idy索引计算方式为idy＝blockIdx.y，其中blockIdx.y对应Init_Sig_staVPosblocks.y中各个线程块索引，blockDim.y对应Init_Sig_staVPosblocks.y中各个线程块所对应的维度，threadIdx.y则对应Init_Sig_staVPosblocks.y中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，Na]为方位向索引,代表方位向。使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能为完成图6，为关于BP成像的具体操作流程示意图中所示的BP成像有效区域的选取模块。为该核函数分配的数组变量分别为Sig_staV、LenNa、Sig_start、Tr；

其中，Sig_staV表示方位向成像区域大小，分配的空间大小为sizeof(double)*LenNa，LenNa为选取的方位向有效长度，Sig_start为方位向成像开始的索引，Tr为雷达脉冲周期，通过该核函数可以完成对有效的方位向成像区域的选取。

Kernel5核函数名称为Init_Grdx，块内配置线程为Init_GrdxThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx，对应距离向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定。核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应Init_GrdxPosblocks.x中各个线程块索引，blockDim.x对应Init_GrdxPosblocks.x中各个线程块所对应的维度，threadIdx.x则对应Init_GrdxPosblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，2*Hf_az]；使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能是完成图6，为关于BP成像的具体操作流程示意图中所示的对地面网格的X方向的方位划分模块。为这个核函数分配的数组变量分别为Grd_x、Cen_pos2、Hf_az、Pix_az；Grd_x含义为创建的方位向网格，分配的空间为sizeof(double)*2*Hf_az，其中Hf_az为方位向网格点个数，Cen_pos2表示网格中心点距雷达的方位距离，Pix_az表示方位向分辨率。

Kernel6核函数名称为Init_Grd_Slt_y，块内配置线程为Init_Grd_Slt_yThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx，对应距离向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定。核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应Init_Grd_Slt_yblocks.x中各个线程块索引，blockDim.x对应Init_Grd_Slt_yblocks.x中各个线程块所对应的维度，threadIdx.x则对应Init_Grd_Slt_yblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，2*Hf_rg]；使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能是完成图6，为关于BP成像的具体操作流程示意图中所示的地面网格的Y方向的方位划分模块。为这个核函数Init_Grd_Slt_y分配的数组变量分别为Grd_Slt、Grd_y、Sqrt_temp，Hf_rg、Pix_rg、H，各参数具体含义如下：Grd_Slt表示斜距平面网格，Grd_y表示距离向网格，Grd_Slt与Grd_y分配的数组空间为sizeof(double)*2*Hf_rg；Hf_rg表示距离向网格点个数，Pix_rg表示距离向分辨率，H表示地面据雷达的高度，通过核函数Init_Grdx与核函数Init_Grd_Slt_y，完成了地面网格的划分。

Kernel7核函数名称为Init_SenPos，块内配置线程为Init_SenPosThread(512,1)，网格内配置块为含义是创建一个大小为的线程网格，其中线程网格的单位是线程块，块内配置512*1个线程，可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx，对应距离向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定。核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应Init_SenPos.x中各个线程块索引，blockDim.x对应Init_SenPos.x中各个线程块所对应的维度，threadIdx.x则对应Init_SenPos.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，2*len]；使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能是完成图6，为关于BP成像的具体操作流程示意图中所示的起始雷达位置有效区域的选取模块。为这个核函数分配的数组变量分别为SensPos，sen_sor_pos、StaEchoAz、len、Na；sensPos与sensor_pos分配的空间均为sizeof(double)*3*len，其中，SensPos与sen_sor_pos用来表示雷达扫描各个时刻的位置信息，StaEchoAz为数据的方位向有效起始，len为选取的方位向的有效长度，Na为回波数据的方位向长度，通过该核函数完成有效成像坐标系的生成。

Kernel8核函数名称为Init_AzSE，块内配置线程为Init_AzSEThread(512,1)，网格内配置块为可根据线程的配置而自动对块的配置进行调整，示意图如图9(b)，为核函数Kernel1-Kernel8的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx，对应距离向索引，根据核的配置，idx索引由线程X方向偏移、块X方向偏移以及块的大小共同决定。核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x，其中blockIdx.x对应Init_AzSEblocks.x中各个线程块索引，blockDim.x对应Init_AzSEblocks.x中各个线程块所对应的维度，threadIdx.x则对应Init_AzSEblocks.x中各个维度线程块中的线程束索引，线程束数量为16的倍数即可，索引范围是[0，2*len]；使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能是完成图6，为关于BP成像的具体操作流程示意图中所示的有效孔径的选取模块。为该核函数Init_AzSE分配的数组变量分别为Apt_StaAz、Apt_EndAz、SensPos，Rc、Ant_Squint、La_hf、len；len为选取的方位向有效长度，Apt_StaAz与Apt_EndAz分别表示起始孔径与截至孔径；SensPos表示雷达扫描时刻的位置信息，SensPos分配的空间大小为sizeof(doule)*3*len，Ant_Squint表示雷达斜视角，La_hf表示合成孔径天线长度的一半，Rc表示场景中心距离。

通过核函数Init_SenPos与核函数Init_AzSE，完成了有效孔径的选取。

最后的也是关键步骤完成图6，为关于BP成像的具体操作流程示意图中所示的GPU设备端BP成像部分工作。Kernel9核函数名称为ProcessPoint，块内配置线程为ProcessPointThread(512,1,1)，网格内配置块为，含义是创建一个大小为 的线程网格，其中线程网格的单位是线程块，块内配置512*1*1个线程；可根据线程的配置而自动对块的配置进行调整，示意图如图9(c)，为核函数Kerne9的线程分配示意图所示。在核内建立线程索引，设置线程索引为idx、idy、idz，根据核的配置，各个线程索引由线程方向偏移、块方向偏移以及块的大小共同决定；核函数idx索引计算方式为idx＝blockIdx.x*blockDim.x+threadIdx.x、idy＝blockIdx.y、idz＝blockIdx.z+offset_Na、索引范围是[0，2*len]，len为方位向选取的有效长度，其中blockIdx.x、blockIdx.y、blockIdx.z对应ProcessPointblocks.x、ProcessPointblocks.y、ProcessPointblocks.z中不同维度各个线程块索引，blockDim.x、blockDim.y对应ProcessPointblocks.x、ProcessPointblocks.y中各个线程块所对应的维度，threadIdx.x、threadIdx.y则对应ProcessPointblocks.x、ProcessPointblocks.y中各个维度线程块中的线程束索引。使用上述索引遍历GPU显存中的变量地址空间，为GPU显存所开辟数组的数组元素赋值和数据计算。

功能是实现雷达回波数据的sinc插值(升采样)、完成雷达与目标瞬时斜距的计算、同时完成距离向相位补偿、并最终完成BP成像。

分析SAR图像BP成像的流程，明确各个任务之间的关系，依据任务的依赖性和独立性对任务进行划分。依赖性是指两个任务之间存在因果关系，部分或全部的数据互用。独立性是指两个任务交换处理顺序，对结果没有任何影响。

将经过上述BP成像处理的雷达回波数据由GPU设备端读取到CPU主机端，在CPU端将所得到的数据写入到.data类型文件中，在matlab软件中读取展示，如图11(a)，为本发明中星载SAR获取的雷达回波信号所示。对回波信号进行处理得到的图像如图11(b)，为本发明中对雷达回波信号进行处理得到的图像所示。

相对于CPU主机端串行处理雷达回波数据，在GPU设备端并行处理雷达回波数据时，GPU设备端与CPU主机端不共享内存，且GPU不适合直接大量地读取CPU内存数据，所以需要通过CPU主机端将雷达回波数据传输至GPU设备端内存中，在多个GPU设备端存在的情况下，需要将仿真参数根据需要从CPU主机端传输至GPU设备端中的全局内存。在雷达回波数据计算过程中，结合GPU特点，调用GPU核函数并给核函数分配大量设备端线程，每个线程对应于每个雷达目标点在某个雷达扫描时刻的雷达回波数据，这些线程可在GPU核函数内部并行计算，完成大量数据的并行处理。在GPU设备端处理SAR回波数据过程中，为了防止GPU设备端内存冲突，在CPU主机端额外配置核函数调用多个线程在GPU设备端完成相关计算。

仿真效果对比：图10，为未采用本发明基于多服务器多GPU的星载SAR成像实时处理方法的实测数据仿真图的成像结果相较于图11(b)，为本发明中对雷达回波信号进行处理得到的图像无明显差别。针对相同点数的SAR雷达回波数据仿真，串行BP成像耗时5434.18s，多服务器多GPU耗时4s，加速比1358.5倍，有效减少仿真耗时。

虽然，本说明书中已经用一般性说明及具体实施方案对本发明作了详尽的描述，但在本发明基础上，可以对之作一些修改或改进，这对本领域技术人员而言是显而易见的。因此，在不偏离本发明精神的基础上所做的这些修改或改进，均属于本发明要求保护的范围。