251 lines
6.9 KiB
Plaintext
251 lines
6.9 KiB
Plaintext
#include <cstdio>
|
|
#include <cufft.h>
|
|
#include <cmath>
|
|
#include <cuda_runtime.h>
|
|
#include <iostream>
|
|
#include <memory>
|
|
#include <vector>
|
|
#include <cmath>
|
|
#include <complex>
|
|
#include <device_launch_parameters.h>
|
|
#include <cuda_runtime.h>
|
|
#include <cublas_v2.h>
|
|
#include <cuComplex.h>
|
|
#include <cufft.h>
|
|
#include <cufftw.h>
|
|
#include <cufftXt.h>
|
|
#include <cublas_v2.h>
|
|
#include <cuComplex.h>
|
|
#include "BaseConstVariable.h"
|
|
#include "GPUTool.cuh"
|
|
#include "GPUBPTool.cuh"
|
|
#include "BPBasic0_CUDA.cuh"
|
|
#include "GPUBpSimulation.cuh"
|
|
#include "GPURFPC_single.cuh"
|
|
#include "GPURFPC.cuh"
|
|
|
|
|
|
double* getFreqPoints_mallocHost(double startFreq, double endFreq, long freqpoints)
|
|
{
|
|
long double dfreq = (endFreq - startFreq) / (freqpoints - 1);
|
|
double* freqlist = (double*)mallocCUDAHost(sizeof(double) * freqpoints);
|
|
for (long i = 0; i < freqpoints; i++) {
|
|
freqlist[i] = startFreq + dfreq * i;
|
|
}
|
|
return freqlist;
|
|
}
|
|
|
|
cuComplex* createEchoPhase_mallocHost(long Np, long Nf)
|
|
{
|
|
cuComplex* phdata = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * Np * Nf);
|
|
for (long i = 0; i < Np; i++) {
|
|
for (long j = 0; j < Nf; j++) {
|
|
phdata[i * Nf + j] = make_cuComplex(0, 0);
|
|
}
|
|
}
|
|
return phdata;
|
|
}
|
|
|
|
|
|
__global__ void kernel_RFPCProcess(
|
|
double* Sx,double* Sy,double* Sz,
|
|
double Tx, double Ty,double Tz,
|
|
double Tslx,double Tsly,double Tslz, // 目标的坡面向量
|
|
double p1,double p2, double p3, double p4, double p5, double p6,
|
|
long Np,long Nf,
|
|
double minF,double dFreq,double RefRange,
|
|
cuComplex* phdata
|
|
) {
|
|
long prfid = blockIdx.x * blockDim.x + threadIdx.x; // 获取当前的线程编码
|
|
if (prfid >= Np || prfid < 0) { return; }
|
|
else {}
|
|
|
|
// 距离
|
|
Vector3 S{ Sx[prfid],Sy[prfid],Sz[prfid] };
|
|
Vector3 T{ Tx,Ty,Tz };
|
|
Vector3 slp{ Tslx,Tsly,Tslz };
|
|
//printf("S=(%e,%e,%e)\n", S.x, S.y, S.z);
|
|
// 入射角
|
|
Vector3 TS = vec_sub(S, T);//T-->S
|
|
double localIncAngle = angleBetweenVectors(TS, slp, false);// 入射角
|
|
double sigma0 = GPU_getSigma0dB_params(p1, p2, p3, p4, p5, p6, localIncAngle);// 后向散射系数
|
|
sigma0 = powf(10.0, sigma0 / 10.0);
|
|
|
|
// 距离
|
|
double R = sqrt(vec_dot(TS, TS));
|
|
|
|
// 计算增益
|
|
double amp_echo = sigma0 / (powf(4 * LAMP_CUDA_PI, 2) * powf(R, 4)); // 反射强度
|
|
double phi = 0;
|
|
for (long fid = 0; fid < Nf; fid++) {
|
|
phi = -4.0 * PI / LIGHTSPEED * (minF + dFreq * fid) * (R - RefRange);
|
|
phdata[prfid * Nf + fid].x += amp_echo * cos(phi);
|
|
phdata[prfid * Nf + fid].y += amp_echo * sin(phi);
|
|
//printf("phdata(%d,%d)=complex(%e,%e)\n", prfid, fid, phdata[prfid * Nf + fid].x, phdata[prfid * Nf + fid].y);
|
|
}
|
|
//printf("Nf:%d\n", Nf);
|
|
//printf("amp_echo:%f\n", amp_echo);
|
|
//printf("sigma0:%f\n", sigma0);
|
|
//printf("R:%e\n", R);
|
|
}
|
|
|
|
|
|
void RFPCProcess(double Tx, double Ty, double Tz,
|
|
double Tslx, double Tsly, double Tslz, // 目标的坡面向量
|
|
double p1, double p2, double p3, double p4, double p5, double p6,
|
|
GPUDATA& d_data)
|
|
{
|
|
double* AntX = (double*)mallocCUDADevice(sizeof(double) * d_data.Np);
|
|
double* AntY = (double*)mallocCUDADevice(sizeof(double) * d_data.Np);
|
|
double* AntZ = (double*)mallocCUDADevice(sizeof(double) * d_data.Np);
|
|
|
|
HostToDevice(d_data.AntX, AntX, sizeof(double) * d_data.Np); //printf("antX host to device finished!!\n");
|
|
HostToDevice(d_data.AntY, AntY, sizeof(double) * d_data.Np); //printf("antY host to device finished!!\n");
|
|
HostToDevice(d_data.AntZ, AntZ, sizeof(double) * d_data.Np); //printf("antZ host to device finished!!\n");
|
|
double minF = d_data.minF[0];
|
|
long grid_size = (d_data.Np + BLOCK_SIZE - 1) / BLOCK_SIZE;
|
|
double dfreq = d_data.deltaF;
|
|
double R0 = d_data.R0;
|
|
kernel_RFPCProcess<<<grid_size , BLOCK_SIZE >>>(
|
|
AntX, AntY, AntZ,
|
|
Tx, Ty, Tz,
|
|
Tslx, Tsly, Tslz, // 目标的坡面向量
|
|
p1, p2, p3, p4, p5, p6,
|
|
d_data.Np, d_data.Nfft,
|
|
minF, dfreq,R0,
|
|
d_data.phdata
|
|
);
|
|
|
|
PrintLasterError("RFPCProcess");
|
|
FreeCUDADevice(AntX);
|
|
FreeCUDADevice(AntY);
|
|
FreeCUDADevice(AntZ);
|
|
|
|
return ;
|
|
}
|
|
|
|
|
|
|
|
/** 单精度成像测试 *****************************************************************************************************************************************/
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
__global__ void kernel_RFPCProcess_single(
|
|
float* Sx, float* Sy, float* Sz,
|
|
float Tx, float Ty, float Tz,
|
|
float Tslx, float Tsly, float Tslz, // 目标的坡面向量
|
|
float p1, float p2, float p3, float p4, float p5, float p6,
|
|
long Np, long Nf,
|
|
float minF, float dFreq, float RefRange,
|
|
cuComplex* phdata
|
|
) {
|
|
long prfid = blockIdx.x * blockDim.x + threadIdx.x; // 获取当前的线程编码
|
|
if (prfid >= Np || prfid < 0) { return; }
|
|
else {}
|
|
|
|
// 距离
|
|
Vector3_single S{ Sx[prfid],Sy[prfid],Sz[prfid] };
|
|
Vector3_single T{ Tx,Ty,Tz };
|
|
Vector3_single slp{ Tslx,Tsly,Tslz };
|
|
//printf("S=(%e,%e,%e)\n", S.x, S.y, S.z);
|
|
// 入射角
|
|
Vector3_single TS = vec_sub_single(S, T);//T-->S
|
|
float localIncAngle = angleBetweenVectors_single(TS, slp, false);// 入射角
|
|
float sigma0 = GPU_getSigma0dB_single(p1, p2, p3, p4, p5, p6, localIncAngle);// 后向散射系数
|
|
sigma0 = powf(10.0, sigma0 / 10.0);
|
|
|
|
// 距离
|
|
float R = sqrtf(vec_dot_single(TS, TS));
|
|
|
|
// 计算增益
|
|
float amp_echo = sigma0 / (powf(4 * LAMP_CUDA_PI, 2) * powf(R, 4)); // 反射强度
|
|
float phi = 0;
|
|
for (long fid = 0; fid < Nf; fid++) {
|
|
phi = -4.0 * PI / LIGHTSPEED * (minF + dFreq * fid) * (R - RefRange);
|
|
phdata[prfid * Nf + fid].x += amp_echo * cosf(phi);
|
|
phdata[prfid * Nf + fid].y += amp_echo * sinf(phi);
|
|
//printf("phdata(%d,%d)=complex(%e,%e)\n", prfid, fid, phdata[prfid * Nf + fid].x, phdata[prfid * Nf + fid].y);
|
|
}
|
|
//printf("Nf:%d\n", Nf);
|
|
//printf("amp_echo:%f\n", amp_echo);
|
|
//printf("sigma0:%f\n", sigma0);
|
|
//printf("R:%e\n", R);
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
float* getFreqPoints_mallocHost_single(float startFreq, float endFreq, long freqpoints)
|
|
{
|
|
float dfreq = (endFreq - startFreq) / (freqpoints - 1);
|
|
float* freqlist = (float*)mallocCUDAHost(sizeof(float) * freqpoints);
|
|
for (long i = 0; i < freqpoints; i++) {
|
|
freqlist[i] = startFreq + dfreq * i;
|
|
}
|
|
return freqlist;
|
|
}
|
|
|
|
void RFPCProcess_single(
|
|
float Tx, float Ty, float Tz,
|
|
float Tslx, float Tsly, float Tslz,
|
|
float p1, float p2, float p3, float p4, float p5, float p6,
|
|
GPUDATA_single& d_data)
|
|
{
|
|
float* AntX = (float*)mallocCUDADevice(sizeof(float) * d_data.Np);
|
|
float* AntY = (float*)mallocCUDADevice(sizeof(float) * d_data.Np);
|
|
float* AntZ = (float*)mallocCUDADevice(sizeof(float) * d_data.Np);
|
|
|
|
HostToDevice(d_data.AntX, AntX, sizeof(float) * d_data.Np); printf("antX host to device finished!!\n");
|
|
HostToDevice(d_data.AntY, AntY, sizeof(float) * d_data.Np); printf("antY host to device finished!!\n");
|
|
HostToDevice(d_data.AntZ, AntZ, sizeof(float) * d_data.Np); printf("antZ host to device finished!!\n");
|
|
float minF = d_data.minF[0];
|
|
long grid_size = (d_data.Np + BLOCK_SIZE - 1) / BLOCK_SIZE;
|
|
float dfreq = d_data.deltaF;
|
|
float R0 = d_data.R0;
|
|
kernel_RFPCProcess_single << <grid_size, BLOCK_SIZE >> > (
|
|
AntX, AntY, AntZ,
|
|
Tx, Ty, Tz,
|
|
Tslx, Tsly, Tslz, // 目标的坡面向量
|
|
p1, p2, p3, p4, p5, p6,
|
|
d_data.Np, d_data.Nfft,
|
|
minF, dfreq, R0,
|
|
d_data.phdata
|
|
);
|
|
|
|
PrintLasterError("RFPCProcess");
|
|
FreeCUDADevice(AntX);
|
|
FreeCUDADevice(AntY);
|
|
FreeCUDADevice(AntZ);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|