[Python知识库]C/CUDA雷达信号处理流程并行化设计

目的

将雷达信号处理流程脉冲压缩、MTI、加窗、MTD、CFAR过程进行并行化设计

代码

// An highlighted block
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <windows.h>
#include <ctime>
#include "Parameter.h"
#include "ReadSaveData.h"
#include "malloc.h"
#include <cufft.h>         
#include "ComputeEchoNoise.h"
#include "ComputeHFilterFFT.h"
#include "ComputeCoherentAccumulation.h"
#include <vector>
#include "TargetRecognition.h"
#include "AddWin.h"
#include <iostream>
#include "Parameter.h"
using namespace std;
using std::vector;

__global__ void MTI(cufftComplex *idata, cufftComplex *odata, const int size)
{
	int blockId = blockIdx.x + blockIdx.y * gridDim.x;
	int threadId = blockId * (blockDim.x * blockDim.y)
		+ (threadIdx.y * blockDim.x) + threadIdx.x;
	unsigned int index = threadId;

	if (N_jie == 1)//一次对消
	{
		if (index < size)
		{
			odata[index].x = idata[index + NWITH0].x - idata[index].x;
			odata[index].y = idata[index + NWITH0].y - idata[index].y;
		}
	}
	if (N_jie == 2)//二次对消
	{
		if (index < size)
		{
			odata[index].x = idata[index + 2 * NWITH0].x - 2 * idata[index + NWITH0].x + idata[index].x;
			odata[index].y = idata[index + 2 * NWITH0].y - 2 * idata[index + NWITH0].y + idata[index].y;
		}
	}

}





int main()
{


	/* 读取数据*/
	cufftComplex *data_Host = (cufftComplex*)malloc(NWITH0 * Nfft * sizeof(cufftComplex)); // 回波数据
	ReadData(data_Host, NWITH0 * Nfft); // 读取数据

	/* 开辟设备端的内存空间 */
	cufftComplex *H_dev; // 匹配滤波器，长度为 NWITH0, 1行
	cudaMalloc((void**)&H_dev, NWITH0 * sizeof(cufftComplex));
	cudaMemset(H_dev, 0, NWITH0 * sizeof(cufftComplex));

	cufftComplex *data_NWITH0_dev; // 每行长度为 NWITH0, Nfft 行
	cudaMalloc((void**)&data_NWITH0_dev, Nfft * NWITH0 * sizeof(cufftComplex));

	cufftComplex *data_MTI_dev; // MTI结果
	cudaMalloc((void**)&data_MTI_dev, Nfft * NWITH0 * sizeof(cufftComplex));
	cudaMemset(data_MTI_dev, 0, Nfft * NWITH0 * sizeof(cufftComplex));

	double *Win; // 窗
	cudaMalloc((void**)&Win, Nfft * sizeof(double));
	cudaMemset(Win, 0, Nfft * sizeof(double));

	cufftComplex *data_AddWin; // 加窗结果
	cudaMalloc((void**)&data_AddWin, Nfft * NWITH0 * sizeof(cufftComplex));
	cudaMemset(data_AddWin, 0, Nfft * NWITH0 * sizeof(cufftComplex));

	cufftComplex *data_N_dev; // MTD结果
	cudaMalloc((void**)&data_N_dev, Nfft * NWITH0 * sizeof(cufftComplex));
	cudaMemset(data_N_dev, 0, Nfft * NWITH0 * sizeof(cufftComplex));

	double *dataTAbs_N_dev; // MTD求abs结果
	cudaMalloc((void**)&dataTAbs_N_dev, Nfft * NWITH0 * sizeof(double));
	cudaMemset(dataTAbs_N_dev, 0, Nfft * NWITH0 * sizeof(double));

	double *targetYoN_dev; // 每行长度为 NWITH0, Nfft 行
	cudaMalloc((void**)&targetYoN_dev, Nfft * NWITH0 * sizeof(double));

	/* 开辟主机端的内存空间 */
	double *target_host = (double*)malloc(Nfft * NWITH0 *sizeof(double));
	memset(target_host, 0, Nfft * NWITH0  * sizeof(double));



	/* 申请 cufft 句柄*/
	cufftHandle plan_NWITH0_One, plan_NWITH0_Many, plan_Nfft_Many; // 创建cuFFT句柄
	cufftPlan1d(&plan_NWITH0_One, NWITH0, CUFFT_C2C, BATCH);
	int number_NWITH0[1] = { NWITH0 };
	int inembed[2];
	int onembed[2];
	inembed[0] = NWITH0;
	inembed[1] = Nfft;
	onembed[0] = NWITH0;
	onembed[1] = Nfft;


	/* 设置脉冲压缩FFT句柄*/
	cufftPlanMany(&plan_NWITH0_Many, 1, number_NWITH0, inembed, 1, NWITH0, onembed, 1, NWITH0, CUFFT_C2C, Nfft);

	int number_Nfft[1] = { Nfft };

	/* 设置MTD FFT句柄*/
	cufftPlanMany(&plan_Nfft_Many, 1, number_Nfft, inembed, NWITH0, 1, onembed, NWITH0, 1, CUFFT_C2C, NWITH0);


	/* 计算匹配滤波器并进行fft运算*/
	ComputeHFilterFFT(H_dev, plan_NWITH0_One); 

	cudaMemset(data_NWITH0_dev, 0, Nfft*NWITH0 * sizeof(cufftComplex));
	cudaMemcpy(data_NWITH0_dev, data_Host, Nfft*NWITH0 * sizeof(cufftComplex), cudaMemcpyHostToDevice);


	/* 整体脉冲压缩*/
	ComputeEchoNoise(H_dev, data_NWITH0_dev, plan_NWITH0_One, plan_NWITH0_Many, 0); 


	/* 加汉明窗*/
	AddWin(Win, data_NWITH0_dev, data_AddWin);

	/* MTI*/
	dim3 dimBlock2D(BLOCK_SIZE, BLOCK_SIZE);
	dim3 dimGrid2D_NWITH0_Nfft((NWITH0 + BLOCK_SIZE - 1) / dimBlock2D.x, (Nfft + BLOCK_SIZE - 1) / dimBlock2D.y);
	MTI << <dimBlock2D, dimGrid2D_NWITH0_Nfft >> >(data_AddWin, data_MTI_dev, (Nfft - 1)*NWITH0);


	/* MTD*/
	ComputeCoherentAccumulation(data_MTI_dev, data_MTI_dev, dataTAbs_N_dev, plan_Nfft_Many); 


	/* CFAR*/
	TargetRecognition(dataTAbs_N_dev, targetYoN_dev); 

	cudaMemcpy(target_host, targetYoN_dev, NWITH0 * Nfft * sizeof(double), cudaMemcpyDeviceToHost); // 数据从设备内存拷贝到主机内存


	cudaFree(H_dev);
	cudaFree(data_NWITH0_dev);
	cudaFree(data_MTI_dev);
	cudaFree(data_N_dev);
	cudaFree(dataTAbs_N_dev);
	cudaFree(targetYoN_dev);
	cudaFree(target_host); 
	cudaFree(data_Host);
	return 0;
}