保存为double 类型

pull/3/head
陈增辉 2025-01-06 12:03:14 +08:00
parent 243af414f2
commit 822c839a40
7 changed files with 186 additions and 334 deletions

View File

@ -14,7 +14,9 @@
#define __CUDANVCC___ // 定义CUDA函数
//#define __PRFDEBUG__
#define __PRFDEBUG__
#define __PRFDEBUG_PRFINF__
//#define __ECHOTIMEDEBUG__
#define __TBPIMAGEDEBUG__

View File

@ -1945,7 +1945,7 @@ bool saveEigenMatrixXd2Bin(Eigen::MatrixXd data, QString dataStrPath)
Eigen::MatrixXd gt = Eigen::MatrixXd::Zero(2, 3);
gdalImage img=CreategdalImage(dataStrPath, data.rows(), data.cols(), 1, gt, "", false, false,true);
gdalImage img=CreategdalImage(dataStrPath, data.rows(), data.cols(), 1, gt, "", false, true,true);
img.saveImage(data, 0,0,1);
@ -2824,10 +2824,27 @@ void testOutAmpArr(QString filename, float* amp, long rowcount, long colcount)
std::cout << filename.toLocal8Bit().constData() << std::endl;
std::cout << "max:\t" << h_amp_img.maxCoeff() << std::endl;
std::cout << "min:\t" << h_amp_img.minCoeff() << std::endl;
}
void testOutAmpArr(QString filename, double* amp, long rowcount, long colcount)
{
Eigen::MatrixXd h_amp_img = Eigen::MatrixXd::Zero(rowcount, colcount);
for (long hii = 0; hii < rowcount; hii++) {
for (long hjj = 0; hjj < colcount; hjj++) {
h_amp_img(hii, hjj) = amp[hii * colcount + hjj];
}
}
QString ampPath = getDebugDataPath(filename);
saveEigenMatrixXd2Bin(h_amp_img, ampPath);
std::cout << filename.toLocal8Bit().constData() << std::endl;
std::cout << "max:\t" << h_amp_img.maxCoeff() << std::endl;
std::cout << "min:\t" << h_amp_img.minCoeff() << std::endl;
}
void testOutClsArr(QString filename, long* amp, long rowcount, long colcount) {
Eigen::MatrixXd h_amp_img = Eigen::MatrixXd::Zero(rowcount, colcount);

View File

@ -275,6 +275,8 @@ bool saveEigenMatrixXd2Bin(Eigen::MatrixXd data, QString dataStrPath);
// 测试
void testOutAntPatternTrans(QString antpatternfilename, float* antPatternArr, double starttheta, double deltetheta, double startphi, double deltaphi, long thetanum, long phinum);
void testOutAmpArr(QString filename, float* amp, long rowcount, long colcount);
void testOutAmpArr(QString filename, double* amp, long rowcount, long colcount);
void testOutClsArr(QString filename, long* amp, long rowcount, long colcount);

View File

@ -409,54 +409,8 @@ __global__ void CUDA_InterpSigma(
}
__global__ void CUDA_CalculationEchoAmp(float* sigma0, float* TransAnt, float* ReciveAnt, float* R,
float Pt,
float* ampArr, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
float r = R[idx];
float amptemp = Pt * TransAnt[idx] * ReciveAnt[idx] * sigma0[idx];
amptemp = amptemp / (powf(4 * LAMP_CUDA_PI, 2) * powf(r, 4)); // 反射强度
ampArr[idx] = amptemp;
}
}
__global__ void CUDA_CalculationEchoPhase(float* R, float lamda, float* phaseArr, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
float r = R[idx];
// 处理相位
float phi = (-4 * LAMP_CUDA_PI / lamda) * r;
phaseArr[idx] = phi;
}
}
__global__ void CUDA_CombinationEchoAmpAndPhase(float* R,
float* echoAmp, float* echoPhase,
float nearRange, float Fs, long plusepoints,
cuComplex* echo, long* FreqID, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
float r = R[idx];
float phase = echoPhase[idx];
float amp = echoAmp[idx];
cuComplex echophi = make_cuComplex(0, phase);
cuComplex echophiexp = cuCexpf(echophi);
float timeR = 2 * (r - nearRange) / LIGHTSPEED * Fs;
long timeID = floorf(timeR);
if (timeID < 0 || timeID >= plusepoints) {
timeID = 0;
amp = 0;
}
cuComplex echotemp = make_cuComplex(echophiexp.x * amp, echophiexp.y * amp);
echo[idx] = echotemp;
FreqID[idx] = timeID;
}
}
__global__ void CUDAKernel_RFPC_Caluation_R_Gain(
float antX, float antY, float antZ, // 天线的坐标
@ -472,18 +426,26 @@ __global__ void CUDAKernel_RFPC_Caluation_R_Gain(
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
float NearR, float FarR, // 距离范围
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
float* outR, // 输出距离
double* outR, // 输出距离
float* outAmp // 输出增益
) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
float RstX = antX - targetX[idx]; // 计算坐标矢量
float RstY = antY - targetY[idx];
float RstZ = antZ - targetZ[idx];
double tx = targetX[idx];
double ty = targetY[idx];
double tz = targetZ[idx];
double RstX = antX - tx; // 计算坐标矢量
double RstY = antY - ty;
double RstZ = antZ - tz;
float slopeX = demSlopeX[idx];
float slopeY = demSlopeY[idx];
float slopeZ = demSlopeZ[idx];
float RstR = sqrtf(RstX * RstX + RstY * RstY + RstZ * RstZ); // 矢量距离
double RstR2 = RstX * RstX + RstY * RstY + RstZ * RstZ;
double RstR = sqrt(RstR2); // 矢量距离
//printf("antX=%f;antY=%f;antZ=%f;targetX=%f;targetY=%f;targetZ=%f;RstR=%.6f;diffR=%.6f;\n",antX,antY,antZ,targetX,targetY,targetZ,RstR, RstR - 9.010858499003178e+05);
if (RstR<NearR || RstR>FarR) {
@ -524,6 +486,7 @@ __global__ void CUDAKernel_RFPC_Caluation_R_Gain(
float sigma0 = 0;
{
long clsid = demCls[idx];
//printf("clsid=%d\n", clsid);
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2) {
sigma0 = 0;
@ -561,96 +524,35 @@ __global__ void CUDAKernel_RFPC_Caluation_R_Gain(
}
}
__global__ void CUDARFPCKernel_Target_Freq_EchoData(
float* InR,
float* InlocalAngle,
float* InampGain,
long* Indemcls,
long len,
float Pt, float freq,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
cuComplex* OutechoArr) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
float sigma0 = 0;
{
long clsid = Indemcls[idx];
float localangle = InlocalAngle[idx];
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2) {
sigma0 = 0;
}
else {}
if (abs(tempsigma.p1) < PRECISIONTOLERANCE &&
abs(tempsigma.p2) < PRECISIONTOLERANCE &&
abs(tempsigma.p3) < PRECISIONTOLERANCE &&
abs(tempsigma.p4) < PRECISIONTOLERANCE &&
abs(tempsigma.p5) < PRECISIONTOLERANCE &&
abs(tempsigma.p6) < PRECISIONTOLERANCE
) {
sigma0 = 0;
}
else {
float sigma = GPU_getSigma0dB(tempsigma, localangle);
sigma0 = powf(10.0, sigma / 10.0);// 后向散射系数
}
}
float amp = Pt * InampGain[idx] * sigma0;
float phi = 4 * PI * freq / LIGHTSPEED * InR[idx];
// 欧拉公式 exp(ix)=cos(x)+isin(x)
// echo=Aexp(ix)=A*cos(x)+i*A*sin(x)
cuComplex echotemp = make_cuComplex(amp * cosf(phi), amp * sinf(phi));
OutechoArr[idx] = echotemp;
}
}
__global__ void CUDAKernel_SumPRF_Temp(cuComplex* d_dem_echo, long plusepoints, long grid_size, cuComplex* d_echo_PRF) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < plusepoints) {
cuComplex echo_PRF = make_cuComplex(0, 0);
for (long i = 0; i < grid_size; i++) {
echo_PRF = cuCaddf(echo_PRF, d_dem_echo[idx * grid_size + i]);
}
d_echo_PRF[idx] = echo_PRF;
}
}
__global__ void CUDAKernel_PRF_CalFreqEcho(
float* Rarr, float* ampArr, long pixelcount,
float* freqpoints, long freqnum,
double* Rarr, float* ampArr, long pixelcount,
double* freqpoints, long freqnum,
double dx, double nearR,
cuComplex* PRFEcho, long prfid) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < freqnum) {
float freq = freqpoints[idx];
float factoj = PI * 4 * freq / LIGHTSPEED;
float phi = 0;
float amptemp = 0;
cuComplex tempfreqEcho = PRFEcho[prfid * freqnum + idx];
double freq = freqpoints[idx];
double lamda = LIGHTSPEED / freq;// 波长
double fatorj = -4 * PI / lamda;
double phi = 0;
double amptemp = 0;
cuComplex tempfreqEcho = PRFEcho[prfid * freqnum + idx];
double R = 0;
for (long i = 0; i < pixelcount; i++) { // 区域积分
phi = factoj * Rarr[i]; // 相位
amptemp = ampArr[i];
// 欧拉公式 exp(ix)=cos(x)+isin(x)
// echo=Aexp(ix)=A*cos(x)+i*A*sin(x)
tempfreqEcho.x = tempfreqEcho.x + amptemp * cosf(phi); // 实部
tempfreqEcho.y = tempfreqEcho.y + amptemp * sinf(phi); // 虚部
R = Rarr[i];
//phi = (R = R - (floor(R / lamda) - 1) * lamda)* fatorj; // 相位
phi = R * fatorj; // 相位
amptemp = ampArr[i];
//printf("amp=%f\n", amptemp);
// Eular; exp(ix)=cos(x)+isin(x)
tempfreqEcho.x = tempfreqEcho.x + amptemp * cos(phi); // 实部
tempfreqEcho.y = tempfreqEcho.y + amptemp * sin(phi); // 虚部
//printf("freqid=%d;freq=%.10f;fatorj=%.12f;d_R=%.10f;phi=%.10f;echo=complex(%.5f,%.5f)\n", idx, freq, fatorj, Rarr[i], phi, tempfreqEcho.x, tempfreqEcho.y);
}
PRFEcho[prfid*freqnum+idx] = tempfreqEcho;
}
}
/** 对外封装接口 *******************************************************************************************************/
extern "C" void SatelliteAntDirectNormal(float* RstX, float* RstY, float* RstZ,
@ -658,8 +560,8 @@ extern "C" void SatelliteAntDirectNormal(float* RstX, float* RstY, float* RstZ,
float antYaxisX, float antYaxisY, float antYaxisZ,
float antZaxisX, float antZaxisY, float antZaxisZ,
float antDirectX, float antDirectY, float antDirectZ,
float* thetaAnt, float* phiAnt
, long len) {
float* thetaAnt, float* phiAnt,
long len) {
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
@ -728,67 +630,7 @@ extern "C" void calculationEcho(float* sigma0, float* TransAnt, float* ReciveAnt
cudaDeviceSynchronize();
}
extern "C" void CUDACalculationEchoAmp(float* sigma0, float* TransAnt, float* ReciveAnt, float* R, float Pt, float* ampArr, long len)
{
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_CalculationEchoAmp << <numBlocks, blockSize >> > (
sigma0, TransAnt, ReciveAnt, R, Pt, ampArr, len);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDA_RFPC_SiglePRF CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
extern "C" void CUDACalculationEchoPhase(float* R, float lamda, float* phaseArr, long len)
{
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_CalculationEchoPhase << <numBlocks, blockSize >> > (
R, lamda, phaseArr, len);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDA_RFPC_SiglePRF CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
extern "C" void CUDACombinationEchoAmpAndPhase(float* R,
float* echoAmp, float* echoPhase,
float nearRange, float Fs, long plusepoints, cuComplex* echo, long* FreqID, long len)
{
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_CombinationEchoAmpAndPhase << <numBlocks, blockSize >> > (
R,
echoAmp, echoPhase,
nearRange, Fs, plusepoints, echo, FreqID, len
);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDA_RFPC_SiglePRF CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
extern "C" void CUDAInterpSigma(
@ -825,7 +667,7 @@ extern "C" void CUDARFPC_Caluation_R_Gain(
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
float NearR, float FarR, // 距离范围
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
float* outR, // 输出距离
double* outR, // 输出距离
float* outAmp // 输出增益
)
{
@ -865,8 +707,9 @@ extern "C" void CUDARFPC_Caluation_R_Gain(
}
extern "C" void CUDA_PRF_CalFreqEcho(
float* Rarr, float* ampArr, long pixelcount,
float* freqpoints, long freqnum,
double* Rarr, float* ampArr, long pixelcount,
double* freqpoints, long freqnum,
double dx, double nearR,
cuComplex* PRFEcho, long prfid)
{
int blockSize = 256; // 每个块的线程数
@ -875,6 +718,7 @@ extern "C" void CUDA_PRF_CalFreqEcho(
CUDAKernel_PRF_CalFreqEcho << <numBlocks, blockSize >> > (
Rarr, ampArr, pixelcount,
freqpoints, freqnum,
dx,nearR,
PRFEcho, prfid
);
#ifdef __CUDADEBUG__

View File

@ -36,28 +36,6 @@ extern "C" void calculationEcho(float* sigma0, float* TransAnt, float* ReciveAnt
cuComplex* echoAmp, long* FreqID,
long len);
extern "C" void CUDACalculationEchoAmp(
float* sigma0,
float* TransAnt, float* ReciveAnt,
float* R,
float Pt,
float* ampArr,
long len
);
extern "C" void CUDACalculationEchoPhase(
float* R, float lamda,
float* phaseArr,
long len
);
extern "C" void CUDACombinationEchoAmpAndPhase(float* R, float* echoAmp, float* echoPhase,
float nearRange, float Fs,long plusepoints,
cuComplex* echo,long* FreqID,
long len
);
extern "C" void CUDAInterpSigma(
long* demcls, float* sigmaAmp, float* localanglearr, long len,
@ -78,15 +56,16 @@ extern "C" void CUDARFPC_Caluation_R_Gain(
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
float NearR,float FarR, // 距离范围
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
float* outR, // Êä³ö¾àÀë
double* outR, // 输出距离
float* outAmp // 输出增益
);
extern "C" void CUDA_PRF_CalFreqEcho(
float* Rarr,float* amp,long pixelcount,//
float* freqpoints,long freqnum,
cuComplex* PRFEcho,long prfid
double* Rarr, float* amp, long pixelcount,//
double* freqpoints, long freqnum,
double dx, double nearR,
cuComplex* PRFEcho, long prfid
);

View File

@ -74,11 +74,11 @@
<ClCompile>
<PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;_SILENCE_NONFLOATING_COMPLEX_DEPRECATION_WARNING;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<OpenMPSupport>true</OpenMPSupport>
<Optimization>Disabled</Optimization>
<WholeProgramOptimization>false</WholeProgramOptimization>
<Optimization>MaxSpeed</Optimization>
<WholeProgramOptimization>true</WholeProgramOptimization>
<EnableParallelCodeGeneration>true</EnableParallelCodeGeneration>
<DebugInformationFormat>EditAndContinue</DebugInformationFormat>
<FavorSizeOrSpeed>Neither</FavorSizeOrSpeed>
<DebugInformationFormat>ProgramDatabase</DebugInformationFormat>
<FavorSizeOrSpeed>Speed</FavorSizeOrSpeed>
</ClCompile>
<Link>
<LinkTimeCodeGeneration>Default</LinkTimeCodeGeneration>

View File

@ -142,7 +142,7 @@ ErrorCode RFPCProcessCls::Process(long num_thread)
//stateCode = this->RFPCMainProcess(num_thread);
// 初始化回波
this->EchoSimulationData->initEchoArr(std::complex<float>(0, 0));
stateCode = this->RFPCMainProcess_GPU( );
stateCode = this->RFPCMainProcess_GPU();
if (stateCode != ErrorCode::SUCCESS) {
return stateCode;
@ -200,7 +200,7 @@ ErrorCode RFPCProcessCls::DEMPreprocess()
this->demxyzPath = QDir(tmpfolderPath).filePath("demxyz.bin");
gdalImage demds(this->DEMTiffPath);
gdalImage demxyz = CreategdalImage(demxyzPath, demds.height, demds.width, 3, demds.gt, demds.projection, true, true,true);// X,Y,Z
gdalImage demxyz = CreategdalImage(demxyzPath, demds.height, demds.width, 3, demds.gt, demds.projection, true, true, true);// X,Y,Z
// 分块计算并转换为XYZ
@ -215,7 +215,7 @@ ErrorCode RFPCProcessCls::DEMPreprocess()
double rowidx = 0;
double colidx = 0;
for (int max_rows_ids = 0; max_rows_ids < demds.height; max_rows_ids = max_rows_ids + line_invert) {
Eigen::MatrixXd demdata = demds.getData(max_rows_ids, 0, line_invert, demds.width, 1);
Eigen::MatrixXd xyzdata_x = demdata.array() * 0;
@ -247,8 +247,8 @@ ErrorCode RFPCProcessCls::DEMPreprocess()
this->demsloperPath = QDir(tmpfolderPath).filePath("demsloper.bin");
this->demmaskPath = QDir(tmpfolderPath).filePath("demmask.bin");
gdalImage demsloperxyz = CreategdalImage(this->demsloperPath, demds.height, demds.width, 4, demds.gt, demds.projection, true, true,true);// X,Y,Z,cosangle
gdalImage demmask = CreategdalImage(this->demmaskPath, demxyz.height, demxyz.width, 1, demxyz.gt, demxyz.projection, true, true,true);// X,Y,Z
gdalImage demsloperxyz = CreategdalImage(this->demsloperPath, demds.height, demds.width, 4, demds.gt, demds.projection, true, true, true);// X,Y,Z,cosangle
gdalImage demmask = CreategdalImage(this->demmaskPath, demxyz.height, demxyz.width, 1, demxyz.gt, demxyz.projection, true, true, true);// X,Y,Z
line_invert = 1000;
long start_ids = 0;
@ -339,7 +339,7 @@ ErrorCode RFPCProcessCls::InitEchoMaskArray()
}
ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
{
double widthSpace = LIGHTSPEED / 2 / this->TaskSetting->getFs();
double prf_time = 0;
@ -359,17 +359,17 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
// 回波
long echoIdx = 0;
double NearRange = this->EchoSimulationData->getNearRange(); // 輪訇擒
double FarRange = this->EchoSimulationData->getFarRange();
float NearRange = this->EchoSimulationData->getNearRange(); // 近斜距
float FarRange = this->EchoSimulationData->getFarRange();
double TimgNearRange = 2 * NearRange / LIGHTSPEED;
double TimgFarRange = 2 * FarRange / LIGHTSPEED;
double Fs = this->TaskSetting->getFs(); // 擒燭砃粒欴薹
double Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 楷扞萇揤 1v
float TimgNearRange = 2 * NearRange / LIGHTSPEED;
float TimgFarRange = 2 * FarRange / LIGHTSPEED;
float dx = (FarRange - NearRange) / (PlusePoint - 1);
float Fs = this->TaskSetting->getFs(); // 距离向采样率
float Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 发射电压 1v
//double GainAntLen = -3;// -3dB 为天线半径
long pluseCount = this->PluseCount;
double lamda = this->TaskSetting->getCenterLamda(); // 疏酗
float lamda = this->TaskSetting->getCenterLamda(); // 波长
// 天线方向图
std::shared_ptr<AbstractRadiationPattern> TransformPattern = this->TaskSetting->getTransformRadiationPattern(); // 发射天线方向图
@ -423,7 +423,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
float* d_TantPattern = (float*)mallocCUDADevice(sizeof(float) * Tthetanum * Tphinum);
for (long i = 0; i < Tthetanum; i++) {
for (long j = Tphinum - 1; j >=0 ; j--) {
for (long j = Tphinum - 1; j >= 0; j--) {
//h_TantPattern[i * Tphinum + j] = TransformPattern->getGainLearThetaPhi(TstartTheta + i * Tdtheta, TstartPhi + j * Tdphi);
h_TantPattern[i * Tphinum + j] = TransformPattern->getGain(TstartTheta + i * Tdtheta, TstartPhi + j * Tdphi);
}
@ -432,10 +432,10 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
testOutAntPatternTrans("TransPattern.bin", h_TantPattern, TstartTheta, Tdtheta, TstartPhi, Tdphi, Tthetanum, Tphinum);
for (long i = 0; i < Tthetanum; i++) {
for (long j = 0; j < Tphinum; j++) {
h_TantPattern[i * Tphinum + j] = powf(10.0, h_TantPattern[i * Tphinum + j]/10);
h_TantPattern[i * Tphinum + j] = powf(10.0, h_TantPattern[i * Tphinum + j] / 10);
}
}
HostToDevice(h_TantPattern, d_TantPattern, sizeof(float)* Tthetanum* Tphinum);
HostToDevice(h_TantPattern, d_TantPattern, sizeof(float) * Tthetanum * Tphinum);
// 处理接收天线方向图
double Rminphi = ReceivePattern->getMinPhi();
@ -462,7 +462,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
}
}
testOutAntPatternTrans("ReceivePattern.bin", h_RantPattern, Rmintheta, Rdtheta, RstartPhi, Rdphi, Rthetanum, Rphinum);
for (long i = 0; i < Tthetanum; i++) {
for (long j = 0; j < Tphinum; j++) {
@ -516,22 +516,28 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
// 打印日志
std::cout << "sigma params:" << std::endl;
std::cout << "classid:\tp1\tp2\tp3\tp4\tp5\tp6"<<std::endl;
std::cout << "classid:\tp1\tp2\tp3\tp4\tp5\tp6" << std::endl;
for (long ii = 0; ii < clamapid; ii++) {
std::cout << ii << ":\t" << h_clsSigmaParam[ii].p1;
std::cout << "\t" << h_clsSigmaParam[ii].p2;
std::cout << "\t" << h_clsSigmaParam[ii].p3;
std::cout << "\t" << h_clsSigmaParam[ii].p4;
std::cout << "\t" << h_clsSigmaParam[ii].p5;
std::cout << "\t" << h_clsSigmaParam[ii].p6<<std::endl;
std::cout << "\t" << h_clsSigmaParam[ii].p6 << std::endl;
}
std::cout << "";
}
HostToDevice(h_clsSigmaParam, d_clsSigmaParam, sizeof(CUDASigmaParam) * clamapid);
long blockwidth = demxyz.width;
#ifdef __PRFDEBUG__
blokline = 1;
blockwidth = 1;
#endif
// 地面 XYZ
Eigen::MatrixXd dem_x = demxyz.getData(0, 0, blokline, demxyz.width, 1); // 華醱釴梓
Eigen::MatrixXd dem_x = demxyz.getData(0, 0, blokline, blockwidth, 1); // 地面坐标
long tempDemRows = dem_x.rows();
long tempDemCols = dem_x.cols();
@ -540,61 +546,44 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
Eigen::MatrixXd demsloper_x = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd demsloper_y = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd demsloper_z = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd sloperAngle = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
float* h_dem_x;
float* h_dem_y;
float* h_dem_z;
float* h_demsloper_x;
float* h_demsloper_y;
float* h_demsloper_z;
float* h_demsloper_angle;
float* h_dem_x = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* h_dem_y = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* h_dem_z = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* h_demsloper_x = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* h_demsloper_y = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* h_demsloper_z = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* d_dem_x;
float* d_dem_y;
float* d_dem_z;
float* d_demsloper_x;
float* d_demsloper_y;
float* d_demsloper_z;
float* d_demsloper_angle;
h_dem_x=(float* )mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
h_dem_y=(float* )mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
h_dem_z=(float* )mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
h_demsloper_x=(float* )mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
h_demsloper_y=(float* )mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
h_demsloper_z=(float* )mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
d_dem_x=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols); // 7
d_dem_y=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
d_dem_z=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
d_demsloper_x=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
d_demsloper_y=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
d_demsloper_z=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
float* d_dem_x = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols); // 7
float* d_dem_y = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
float* d_dem_z = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
float* d_demsloper_x = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
float* d_demsloper_y = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
float* d_demsloper_z = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
// 提前声明参数变量
float* h_R=(float* )mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* d_R= (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
double* h_R = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* d_R = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols);
float* h_amp = (float*)mallocCUDAHost( sizeof(float)* blokline* tempDemCols);
float* d_amp = (float*)mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
float* h_amp = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* d_amp = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
// 地面回波
cuComplex* h_echo=(cuComplex*)mallocCUDAHost(sizeof(cuComplex) * blokline * tempDemCols);
cuComplex* d_echo=(cuComplex*)mallocCUDADevice( sizeof(cuComplex) * blokline * tempDemCols); //19
cuComplex* h_echo = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * blokline * tempDemCols);
cuComplex* d_echo = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * blokline * tempDemCols); //19
long echoblockline = Memory1GB / 8 / 2 / PlusePoint*2;
long echoblockline = Memory1GB / 8 / 2 / PlusePoint * 3;
// 每一行的脉冲
cuComplex* h_PRFEcho = (cuComplex*)mallocCUDAHost(sizeof(cuComplex)* echoblockline * PlusePoint);
cuComplex* d_PRFEcho = (cuComplex*)mallocCUDADevice(sizeof(cuComplex)* echoblockline * PlusePoint);
cuComplex* h_PRFEcho = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * echoblockline * PlusePoint);
cuComplex* d_PRFEcho = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * echoblockline * PlusePoint);
float* h_freqpoints = (float*)mallocCUDAHost(sizeof(float) * freqlist.size());
float* d_freqpoints = (float*)mallocCUDADevice(sizeof(float) * freqlist.size());
double* h_freqpoints = (double*)mallocCUDAHost(sizeof(double) * freqlist.size());
double* d_freqpoints = (double*)mallocCUDADevice(sizeof(double) * freqlist.size());
for (long ii = 0; ii < freqlist.size(); ii++) {
h_freqpoints[ii] = freqlist[ii];
}
HostToDevice(h_freqpoints, d_freqpoints, sizeof(float) * freqlist.size());
HostToDevice(h_freqpoints, d_freqpoints, sizeof(double) * freqlist.size());
testOutAmpArr("freqpointslist.bin", h_freqpoints, freqlist.size(), 1);
// 地表覆盖类型
Eigen::MatrixXd landcover = Eigen::MatrixXd::Zero(blokline, tempDemCols);// 地面覆盖类型
long* h_demcls = (long*)mallocCUDAHost(sizeof(long) * blokline * tempDemCols);
@ -606,14 +595,14 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
newblokline = demRow - startline;
bloklineflag = true;
}
dem_x = demxyz.getData(startline, 0, newblokline, demxyz.width, 1); // 華醱釴梓
dem_y = demxyz.getData(startline, 0, newblokline, demxyz.width, 2);
dem_z = demxyz.getData(startline, 0, newblokline, demxyz.width, 3);
demsloper_x = demsloperxyz.getData(startline, 0, newblokline, demsloperxyz.width, 1);
demsloper_y = demsloperxyz.getData(startline, 0, newblokline, demsloperxyz.width, 2);
demsloper_z = demsloperxyz.getData(startline, 0, newblokline, demsloperxyz.width, 3);
dem_x = demxyz.getData(startline, 0, newblokline, blockwidth, 1); // 地面坐标
dem_y = demxyz.getData(startline, 0, newblokline, blockwidth, 2);
dem_z = demxyz.getData(startline, 0, newblokline, blockwidth, 3);
demsloper_x = demsloperxyz.getData(startline, 0, newblokline,blockwidth, 1);
demsloper_y = demsloperxyz.getData(startline, 0, newblokline,blockwidth, 2);
demsloper_z = demsloperxyz.getData(startline, 0, newblokline,blockwidth, 3);
landcover = demlandcls.getData(startline, 0, newblokline, demlandcls.width, 1);
landcover = demlandcls.getData(startline, 0, newblokline, blockwidth, 1);
if (bloklineflag) {
FreeCUDAHost(h_dem_x); FreeCUDADevice(d_dem_x);
@ -626,39 +615,49 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
FreeCUDAHost(h_amp); FreeCUDADevice(d_amp);
FreeCUDAHost(h_echo); FreeCUDADevice(d_echo);//19
FreeCUDAHost(h_demcls); FreeCUDADevice(d_demcls);
h_dem_x = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_dem_y = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_dem_z = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_demsloper_x = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_demsloper_y = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_demsloper_z = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_R = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_R = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_amp = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
h_echo = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * newblokline * tempDemCols);
h_demcls = (long*)mallocCUDAHost(sizeof(long) * newblokline * tempDemCols);
d_dem_x=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_dem_y=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_dem_z=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_demsloper_x=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_demsloper_y=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_demsloper_z=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);//6
d_amp =(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_echo=(cuComplex*)mallocCUDADevice(sizeof(cuComplex) * newblokline * tempDemCols);
d_dem_x = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_dem_y = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_dem_z = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_demsloper_x = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_demsloper_y = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_demsloper_z = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);//6
d_amp = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_echo = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * newblokline * tempDemCols);
d_demcls = (long*)mallocCUDADevice(sizeof(long) * newblokline * tempDemCols);
d_R = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
d_R = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);
}
//# pragma omp parallel for
//# pragma omp parallel for
for (long i = 0; i < newblokline; i++) {
for (long j = 0; j < demxyz.width; j++) {
h_dem_x[i * demxyz.width + j] = float(dem_x(i, j));
h_dem_y[i * demxyz.width + j] = float(dem_y(i, j));
h_dem_z[i * demxyz.width + j] = float(dem_z(i, j));
h_demsloper_x[i * demxyz.width + j] = float(demsloper_x(i, j));
h_demsloper_y[i * demxyz.width + j] = float(demsloper_y(i, j));
h_demsloper_z[i * demxyz.width + j] = float(demsloper_z(i, j));
h_demcls[i * demxyz.width + j] = clamap[long(landcover(i, j))];
for (long j = 0; j < blockwidth; j++) {
#ifdef __PRFDEBUG__
h_dem_x[i * blockwidth + j] = -2028380.6250000; float(dem_x(i, j));
h_dem_y[i * blockwidth + j] = 4139373.250000; float(dem_y(i, j));
h_dem_z[i * blockwidth + j] = 4393382.500000;float(dem_z(i, j));
h_demsloper_x[i * blockwidth + j] = 4393382.500000;float(demsloper_x(i, j));
h_demsloper_y[i * blockwidth + j] = 446.923950;float(demsloper_y(i, j));
h_demsloper_z[i * blockwidth + j] = -219.002213;float(demsloper_z(i, j));
h_demcls[i * blockwidth + j] = clamap[80] ;// clamap[long(landcover(i, j))];
#else
h_dem_x[i * blockwidth + j] = float(dem_x(i, j));
h_dem_y[i * blockwidth + j] = float(dem_y(i, j));
h_dem_z[i * blockwidth + j] = float(dem_z(i, j));
h_demsloper_x[i * blockwidth + j] = float(demsloper_x(i, j));
h_demsloper_y[i * blockwidth + j] = float(demsloper_y(i, j));
h_demsloper_z[i * blockwidth + j] = float(demsloper_z(i, j));
h_demcls[i * blockwidth + j] = clamap[long(landcover(i, j))];
#endif
}
}
@ -668,15 +667,15 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
HostToDevice((void*)h_demsloper_x, (void*)d_demsloper_x, sizeof(float) * newblokline * tempDemCols);
HostToDevice((void*)h_demsloper_y, (void*)d_demsloper_y, sizeof(float) * newblokline * tempDemCols);
HostToDevice((void*)h_demsloper_z, (void*)d_demsloper_z, sizeof(float) * newblokline * tempDemCols);
HostToDevice((void*)h_demcls, (void*)d_demcls, sizeof(long) * newblokline* tempDemCols);
HostToDevice((void*)h_demcls, (void*)d_demcls, sizeof(long) * newblokline * tempDemCols);
#ifdef __PRFDEBUG__
testOutAmpArr("h_dem_x.bin", h_dem_x, newblokline, tempDemCols);
testOutAmpArr("h_dem_y.bin", h_dem_y, newblokline, tempDemCols);
testOutAmpArr("h_dem_z.bin", h_dem_z, newblokline, tempDemCols);
testOutClsArr( "h_demcls.bin" , h_demcls, newblokline, tempDemCols);
#ifdef __PRFDEBUG__ && __PRFDEBUG_PRFINF__
printf("tatgetPs=[%f,%f,%f]\n", h_dem_x[0], h_dem_y[0], h_dem_z[0]);
std::shared_ptr<double> h_temp_R(new double[PluseCount],delArrPtr);
#endif // __PRFDEBUG__
long pixelcount = newblokline * tempDemCols;
@ -692,10 +691,11 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
}
HostToDevice(h_PRFEcho, d_PRFEcho, sizeof(cuComplex) * echoblockline * PlusePoint);
for (long tempprfid = 0; tempprfid < templine; tempprfid++) {
{// 计算
long prfid = tempprfid + startprfid;
std::cout << "[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] "<< prfid<<std::endl;
// 天线位置
float antpx = sateOirbtNodes[prfid].Px;
float antpy = sateOirbtNodes[prfid].Py;
@ -726,7 +726,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
antYaxisX, antYaxisY, antYaxisZ,// 天线坐标系的Y轴
antZaxisX, antZaxisY, antZaxisZ,// 天线坐标系的Z轴
antdirectx, antdirecty, antdirectz,// 天线的指向
Pt,
Pt,// 增益后发射能量
d_TantPattern, TstartTheta, TstartPhi, Tdtheta, Tdphi, Tthetanum, Tphinum, // 发射天线方向图
d_RantPattern, RstartTheta, RstartPhi, Rdtheta, Rdphi, Rthetanum, Rphinum,//接收天线方向图
NearRange, FarRange,
@ -738,11 +738,16 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
CUDA_PRF_CalFreqEcho(
d_R, d_amp, pixelcount,
d_freqpoints, PlusePoint,
dx, NearRange,
d_PRFEcho, tempprfid);
if (prfid % 1000 == 0) {
if (prfid % 100 == 0) {
std::cout << "[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] dem:\t" << startline << "\t-\t" << startline + newblokline << "\t:\t pluse :\t" << prfid << " / " << pluseCount << std::endl;
}
#ifdef __PRFDEBUG__ && __PRFDEBUG_PRFINF__
DeviceToHost(h_R, d_R, sizeof(double)* newblokline* tempDemCols);
h_temp_R.get()[prfid] = h_R[0];
#endif
}
}
@ -756,9 +761,12 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
}
this->EchoSimulationData->saveEchoArr(echotemp, startprfid, templine);
}
#ifdef __PRFDEBUG__ && __PRFDEBUG_PRFINF__
testOutAmpArr("test_out_D_R.bin", h_temp_R.get(), pluseCount, 1);
#endif
#ifdef __PRFDEBUG__
break;
#endif // __PRFDEBUG__
}
@ -842,7 +850,7 @@ std::shared_ptr<SatelliteOribtNode[]> RFPCProcessCls::getSatelliteOribtNodes(dou
antpos.reset();
qDebug() << "Ant position finished sucessfully !!!";
}
return sateOirbtNodes;
}