From 8bd997f92425f9c3445cd7320060d55431ba19a6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E9=99=88=E5=A2=9E=E8=BE=89?= <3045316072@qq.com> Date: Sun, 23 Mar 2025 16:01:28 +0800 Subject: [PATCH] =?UTF-8?q?=E6=96=B0=E5=A2=9E=E5=A4=9A=E5=8D=A1=E6=94=AF?= =?UTF-8?q?=E6=8C=81?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../BaseTool/BaseConstVariable.h | 6 + BaseCommonLibrary/BaseTool/EchoDataFormat.cpp | 13 ++ BaseCommonLibrary/BaseTool/EchoDataFormat.h | 1 + GPUBaseLib/GPUBaseLib.vcxproj | 6 +- GPUBaseLib/GPUBaseLib.vcxproj.filters | 8 + GPUBaseLib/GPUTool/GPUBaseTool.cpp | 56 +++++ GPUBaseLib/GPUTool/GPUBaseTool.h | 25 +++ GPUBaseLib/GPUTool/GPUTool.cuh | 15 +- Toolbox/SimulationSARTool/Sigma0ClsReflect.cu | 132 +++++++++++ .../SimulationSAR/GPURFPC.cu | 122 +++++++++++ .../SimulationSAR/GPURFPC.cuh | 76 ++++++- .../SimulationSAR/RFPCProcessCls.cpp | 206 +++++++++++++++++- .../SimulationSAR/RFPCProcessCls.h | 4 + .../SimulationSARTool.vcxproj | 1 + .../SimulationSARTool.vcxproj.filters | 3 + 15 files changed, 660 insertions(+), 14 deletions(-) create mode 100644 GPUBaseLib/GPUTool/GPUBaseTool.cpp create mode 100644 GPUBaseLib/GPUTool/GPUBaseTool.h create mode 100644 Toolbox/SimulationSARTool/Sigma0ClsReflect.cu diff --git a/BaseCommonLibrary/BaseTool/BaseConstVariable.h b/BaseCommonLibrary/BaseTool/BaseConstVariable.h index 07dccd3..f001b98 100644 --- a/BaseCommonLibrary/BaseTool/BaseConstVariable.h +++ b/BaseCommonLibrary/BaseTool/BaseConstVariable.h @@ -339,6 +339,12 @@ inline long nextpow2(long n) { return pow(2,en); } +inline void releaseVoidArray(void* a) +{ + free(a); + a = NULL; +} + #endif \ No newline at end of file diff --git a/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp b/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp index b2824ba..250ef30 100644 --- a/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp +++ b/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp @@ -511,6 +511,19 @@ ErrorCode EchoL0Dataset::loadFromXml() { return ErrorCode::SUCCESS; } +std::shared_ptr EchoL0Dataset::getAntPosVelc() +{ + omp_lock_t lock; + omp_init_lock(&lock); + omp_set_lock(&lock); + long prfcount = this->PluseCount; + std::shared_ptr antposlist= SatelliteAntPosOperator::readAntPosFile(this->GPSPointFilePath, prfcount); + + omp_unset_lock(&lock); // + omp_destroy_lock(&lock); // + return antposlist; +} + std::shared_ptr EchoL0Dataset::getAntPos() { omp_lock_t lock; diff --git a/BaseCommonLibrary/BaseTool/EchoDataFormat.h b/BaseCommonLibrary/BaseTool/EchoDataFormat.h index 749843f..ec2f685 100644 --- a/BaseCommonLibrary/BaseTool/EchoDataFormat.h +++ b/BaseCommonLibrary/BaseTool/EchoDataFormat.h @@ -209,6 +209,7 @@ public: // public: // 读取文件 + std::shared_ptr< SatelliteAntPos> getAntPosVelc(); std::shared_ptr getAntPos(); std::shared_ptr> getEchoArr(long startPRF, long& PRFLen); std::shared_ptr> getEchoArr(); diff --git a/GPUBaseLib/GPUBaseLib.vcxproj b/GPUBaseLib/GPUBaseLib.vcxproj index 0e21f32..8ebd518 100644 --- a/GPUBaseLib/GPUBaseLib.vcxproj +++ b/GPUBaseLib/GPUBaseLib.vcxproj @@ -32,6 +32,7 @@ + @@ -40,6 +41,9 @@ {872ecd6f-30e3-4a1b-b17c-15e87d373ff6} + + + 17.0 Win32Proj @@ -180,7 +184,7 @@ true true true - NDEBUG;_CONSOLE;GPUBASELIB_API;%(PreprocessorDefinitions) + NDEBUG;_CONSOLE;GPUBASELIB_API;_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions) true stdcpp14 stdc11 diff --git a/GPUBaseLib/GPUBaseLib.vcxproj.filters b/GPUBaseLib/GPUBaseLib.vcxproj.filters index 1d8fae5..4fddae9 100644 --- a/GPUBaseLib/GPUBaseLib.vcxproj.filters +++ b/GPUBaseLib/GPUBaseLib.vcxproj.filters @@ -35,5 +35,13 @@ GPUTool + + GPUTool + + + + + GPUTool + \ No newline at end of file diff --git a/GPUBaseLib/GPUTool/GPUBaseTool.cpp b/GPUBaseLib/GPUTool/GPUBaseTool.cpp new file mode 100644 index 0000000..6832603 --- /dev/null +++ b/GPUBaseLib/GPUTool/GPUBaseTool.cpp @@ -0,0 +1,56 @@ +#include "GPUBaseTool.h" + + +// 获取文件大小 +extern "C" size_t getfsize(FILE* fp) +{ + fseek(fp, 0L, SEEK_END); + size_t size = ftell(fp); + fseek(fp, 0L, SEEK_SET); + return size; +} + +extern "C" unsigned char* loadBinFromPath(char* binPath, size_t* binpath_len) +{ + FILE* fd = fopen(binPath, "rb"); + if (NULL == fd) + { + perror("Failed to open file"); + return NULL; + } + + size_t f_len = getfsize(fd); + *binpath_len = f_len; + + unsigned char* buffer = (unsigned char*)malloc(f_len * sizeof(unsigned char)); + if (buffer == NULL) + { + perror("Failed to allocate memory"); + fclose(fd); + return NULL; + } + + size_t fread_count = fread(buffer, 1, f_len, fd); + fclose(fd); + + if (fread_count != f_len) + { + releaseVoidArray(buffer); + return NULL; + } + + return buffer; +} + +extern "C" void writeComplexDataBinFile(char* dataPath, size_t datalen, cuComplex* data) +{ + FILE* pd = fopen(dataPath, "w"); + double* tempdata = (double*)malloc(datalen * 2 * sizeof(double)); + for (long i = 0; i < datalen; i++) + { + tempdata[i * 2 + 0] = data[i].x; + tempdata[i * 2 + 1] = data[i].y; + } + fwrite(tempdata, sizeof(double), datalen * 2, pd); + fclose(pd); +} diff --git a/GPUBaseLib/GPUTool/GPUBaseTool.h b/GPUBaseLib/GPUTool/GPUBaseTool.h new file mode 100644 index 0000000..e6bf254 --- /dev/null +++ b/GPUBaseLib/GPUTool/GPUBaseTool.h @@ -0,0 +1,25 @@ +#ifndef __GPUBASETOOL_H__ +#define __GPUBASETOOL_H__ + +#include "GPUBaseLibAPI.h" +#include "BaseConstVariable.h" +#include +#include +#include +#include "GPUTool.cuh" + +extern "C" GPUBASELIBAPI size_t getfsize(FILE* fp); +extern "C" GPUBASELIBAPI unsigned char* loadBinFromPath(char* binPath, size_t* binpath_len); +extern "C" GPUBASELIBAPI void writeComplexDataBinFile(char* dataPath, size_t datalen, cuComplex* data); + +template +inline std::shared_ptr CPUToHost(std::shared_ptr CPUArr, size_t len) { + std::shared_ptr result = std::shared_ptr((T*)mallocCUDAHost(len*sizeof(T)), FreeCUDAHost); + for (size_t i = 0; i < len; i++) { + result.get()[i] = CPUArr.get()[i]; + } + return result; +} + + +#endif // !__GPUBASETOOL_H__ diff --git a/GPUBaseLib/GPUTool/GPUTool.cuh b/GPUBaseLib/GPUTool/GPUTool.cuh index 2a9a709..a1730d8 100644 --- a/GPUBaseLib/GPUTool/GPUTool.cuh +++ b/GPUBaseLib/GPUTool/GPUTool.cuh @@ -16,11 +16,13 @@ #define LAMP_CUDA_PI 3.141592653589793238462643383279 // SHAREMEMORY_FLOAT_HALF_STEP * BLOCK_SIZE = SHAREMEMORY_FLOAT_HALF +/** CUDA 调用参数 ************************************************************************************/ #define BLOCK_SIZE 256 -#define SHAREMEMORY_BYTE 49152 -#define SHAREMEMORY_FLOAT_HALF 6144 +#define SHAREMEMORY_BYTE 49152 +#define SHAREMEMORY_FLOAT_HALF 6144 #define SHAREMEMORY_FLOAT_HALF_STEP 24 - +#define SHAREMEMORY_DEM_STEP 768 +#define SHAREMEMORY_Reflect 612 @@ -110,6 +112,13 @@ extern "C" GPUBASELIBAPI void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long extern "C" GPUBASELIBAPI void FFTShift1D(cuComplex* d_data, int batch_size, int signal_length); extern "C" GPUBASELIBAPI void shared_complexPtrToHostCuComplex(std::complex* src, cuComplex* dst, size_t len); + + extern "C" GPUBASELIBAPI void HostCuComplexToshared_complexPtr(cuComplex* src, std::complex* dst, size_t len); + + + + + #endif #endif diff --git a/Toolbox/SimulationSARTool/Sigma0ClsReflect.cu b/Toolbox/SimulationSARTool/Sigma0ClsReflect.cu new file mode 100644 index 0000000..a33f484 --- /dev/null +++ b/Toolbox/SimulationSARTool/Sigma0ClsReflect.cu @@ -0,0 +1,132 @@ +#include "BaseConstVariable.h" +#include "GPURFPC.cuh" + +const double cls_sigma0[102][6] = { + {0, 0, 0, 0, 0, 0}, // 0 + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {28.15, -39.73, 0.0986, 2.863, 4.356, -6.279}, // 10 + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {36.13, -48.2, 0.1299, -1.838, 5.404, -4.015}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {183.5, -194.6, 0.0167, 2.952, -4.1, 6.026}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {50.97, -62.9, 0.0968, 1.604, 4.637, 6.108}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {-10.31, 15.96, 37.73, -4.584, 4.997, -3.569}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {-12.45, 0.1561, -3.187, -2.482, 8.244, 0.3632}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {-19.23, 0.3623, -2.209, 9.649, 0.1292, -0.264}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {0, 0, 0, 0, 0, 0}, + {56.61, -62.29, 0.01388, 2.767, -3.943, 5.995}, + {0, 0, 0, 0, 0, 0} }; + + + +/** 参数 ***************************************************************/ + + + +extern "C" double* hostSigmaData_toDevice(int devid) +{ + double* h_data = (double*)mallocCUDAHost(102 * 6 * sizeof(double)); + double* d_data = (double*)mallocCUDADevice(102 * 6 * sizeof(double), devid); + printf("copy to "); + for (long i = 0; i < 102; i++) + { + printf(" %d ,", i); + for (long j = 0; j < 6; j++) + { + h_data[i * 6 + j] = cls_sigma0[i][j]; + } + } + printf("host to device sigma data\n"); + HostToDevice(h_data, d_data, 102 * 6 * sizeof(double)); + FreeCUDAHost(h_data); + return d_data; +} + diff --git a/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cu b/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cu index cef1a04..cd86f67 100644 --- a/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cu +++ b/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cu @@ -463,6 +463,128 @@ void CUDA_RFPC_MainProcess( } + + + + + + + + + +/* 核函数 ****************************************************************************************************************************/ +__global__ void CUDA_Kernel_RFPC( + SateState* antlist, + long PRFCount, long Freqcount, // 整体的脉冲数， + GoalState* goallist, + long demLen, + double StartFreqGHz, double FreqStep, + double refPhaseRange, + double NearR, double FarR, + CUDASigmaParam clsSigma0, + cuComplex* echodata + ) +{ + + __shared__ GoalState Ts[SHAREMEMORY_DEM_STEP]; + long threadid = threadIdx.x; + + long idx = blockIdx.x * blockDim.x + threadIdx.x; // 获取当前的线程编码 + long prfid = floorf(idx / Freqcount); + long freqid = idx % Freqcount; + // printf("%d,%d ",prfid,freqid); + if (prfid < PRFCount && freqid < Freqcount) + { + SateState antPos = antlist[prfid]; + double factorjTemp = RFPCPIDIVLIGHT * (StartFreqGHz + freqid * FreqStep); + + double Tx = 0; + double Ty = 0; + double Tz = 0; + + double R = 0; + double incAngle = 0; + + double echo_real = 0; + double echo_imag = 0; + + cuComplex echo = make_cuComplex(0, 0); + + for (long startid = 0; startid < demLen; startid = startid + SHAREMEMORY_DEM_STEP) + { + __syncthreads(); // 确定所有待处理数据都已经进入程序中 + for (long i = 0; i < 3; i++) { + long ttid = startid + threadid + i * blockDim.x; + long stid = threadid + i * blockDim.x; + if ((stid < SHAREMEMORY_DEM_STEP) && (ttid < demLen)) { + Ts[stid] = goallist[ttid]; + } + } + __syncthreads(); // 确定所有待处理数据都已经进入程序中 + + for (long tid = 0; tid < SHAREMEMORY_DEM_STEP; tid++) + { + if ((tid + startid) < demLen) + { + GoalState p = Ts[tid]; + Tx = p.Tx; + Ty = p.Ty; + Tz = p.Tz; + + Tx = antPos.Px - Tx; // T->P + Ty = antPos.Py - Ty; + Tz = antPos.Pz - Tz; + + R = sqrt(Tx * Tx + Ty * Ty + Tz * Tz); + bool isNearFar = (R < NearR || R > FarR); + + + incAngle = sqrt(p.TsX * p.TsX + p.TsY * p.TsY + p.TsZ * p.TsZ); + incAngle = acos((Tx * p.TsX + Ty * p.TsY + Tz * p.TsZ) / (R * incAngle)); + incAngle = GPU_getSigma0dB_params(clsSigma0.p1, clsSigma0.p2, clsSigma0.p3, clsSigma0.p4, clsSigma0.p5, clsSigma0.p6, incAngle); // sigma + incAngle = pow(10.0, incAngle / 10.0); // amp + incAngle = incAngle / (powf(4 * LAMP_CUDA_PI, 2) * powf(R, 4)); // + + R = (R - refPhaseRange); + R = factorjTemp * R; + + echo_real = incAngle * cos(R)* isNearFar; + echo_imag = incAngle * sin(R)* isNearFar; + echo.x = echo.x + echo_real; + echo.y = echo.y + echo_imag; + } + } + } + + echodata[idx] = cuCaddf(echodata[idx], echo); + } +} + + +/** 分块处理 ****************************************************************************************************************/ + +extern "C" void ProcessRFPCTask(RFPCTask& task) +{ + long pixelcount = task.prfNum * task.freqNum; + long grid_size = (pixelcount + BLOCK_SIZE - 1) / BLOCK_SIZE; + printf("start %d ,%d,%d\n", task.targetnum, grid_size, BLOCK_SIZE); + CUDA_Kernel_RFPC << > > ( + task.antlist, + task.prfNum,task.freqNum, + task.goallist, + task.targetnum, + task.startFreq,task.stepFreq, + task.Rref,task.Rnear,task.Rfar, + task.sigma0_cls, + task.d_echoData + ); + PrintLasterError("ProcessRFPCTask"); + cudaDeviceSynchronize(); + printf("start %d \n", task.targetnum); +} + + + #endif diff --git a/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cuh b/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cuh index 6b5b4a2..b72623b 100644 --- a/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cuh +++ b/Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cuh @@ -8,8 +8,25 @@ #include #include +/** CUDA 调用参数 ************************************************************************************/ + + + #define RFPCPIDIVLIGHT -4*PI/(LIGHTSPEED/1e9) + +extern "C" struct SateState { + double Px, Py, Pz, Vx, Vy, Vz; +}; + + +extern "C" struct GoalState { + double Tx, Ty, Tz, TsX, TsY, TsZ; + size_t cls; +}; + + + extern "C" struct CUDASigmaParam { double p1; @@ -20,7 +37,60 @@ extern "C" struct CUDASigmaParam { double p6; }; - +extern "C" struct SloperDs +{ + double* SloperX; + double* SloperY; + double* SloperZ; + double* SloperAngle; + +}; + +extern "C" struct DEMDs +{ + double* demX; + double* demY; + double* demZ; +}; + +extern "C" struct LandDataDs +{ + unsigned char* landData; +}; + +extern "C" struct GPSPointsDs +{ + double* Pxs; + double* Pys; + double* Pzs; + double* Vxs; + double* Vys; + double* Vzs; +}; + + +extern "C" struct RFPCTask +{ + double startFreq; + double stepFreq; + long freqNum; + long prfNum; + + double Rnear; + double Rfar; + double Rref; + + + SateState* antlist = nullptr; + GoalState* goallist = nullptr; + cuComplex* d_echoData = nullptr; // 回波 + CUDASigmaParam sigma0_cls; + + + long targetnum; + +}; + extern __device__ double GPU_getSigma0dB_params( const double p1, const double p2, const double p3, const double p4, const double p5, const double p6, @@ -73,8 +143,8 @@ extern "C" void CUDA_RFPC_MainProcess( - - +extern "C" double* hostSigmaData_toDevice(int devid); +extern "C" void ProcessRFPCTask(RFPCTask& task); diff --git a/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.cpp b/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.cpp index 7412081..221d981 100644 --- a/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.cpp +++ b/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.cpp @@ -1,6 +1,7 @@ 锘 #include "stdafx.h" #include +#include "GPUBaseTool.h" #include "RFPCProcessCls.h" #include "BaseConstVariable.h" #include "SARSatelliteSimulationAbstractCls.h" @@ -306,8 +307,8 @@ ErrorCode RFPCProcessCls::Process(long num_thread) //return ErrorCode::SUCCESS; - stateCode = this->RFPCMainProcess_GPU(); - + //stateCode = this->RFPCMainProcess_GPU(); + stateCode = this->RFPCMainProcess_MultiGPU_NoAntPattern(); if (stateCode != ErrorCode::SUCCESS) { return stateCode; } @@ -542,17 +543,12 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() { long freqnum_temp = freqnum; float f0 = float(freqlist[0] / 1e9); float dfreq = float((freqlist[1] - freqlist[0]) / 1e9); - qDebug() << "freqnum: " << freqnum << " f0: " << f0 << " dfreq: " << dfreq; long PRFCount = this->EchoSimulationData->getPluseCount(); - double NearRange = this->EchoSimulationData->getNearRange(); // 杩戞枩璺 double FarRange = this->EchoSimulationData->getFarRange(); - double Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 鍙戝皠鐢靛帇 1v - double lamda = this->TaskSetting->getCenterLamda(); // 娉㈤暱 double refphaseRange = this->TaskSetting->getRefphaseRange(); // 鍙傝冪浉浣嶆枩璺 - double prf_time = 0; double dt = 1 / this->TaskSetting->getPRF();// 鑾峰彇姣忔鑴夊啿鐨勬椂闂撮棿闅 bool antflag = true; // 璁＄畻澶╃嚎鏂瑰悜鍥 @@ -935,6 +931,202 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() { return ErrorCode::SUCCESS; } +ErrorCode RFPCProcessCls::RFPCMainProcess_MultiGPU_NoAntPattern() +{ + int num_devices; + cudaGetDeviceCount(&num_devices); + PRINT("GPU Count : %d \n", num_devices); + + long prfcount = this->EchoSimulationData->getPluseCount(); + size_t prfblockcount = (prfcount + num_devices +2- 1) / num_devices; + + double prf_time = 0; + double dt = 1 / this->TaskSetting->getPRF();// 鑾峰彇姣忔鑴夊啿鐨勬椂闂撮棿闅 + bool antflag = true; // 璁＄畻澶╃嚎鏂瑰悜鍥 + long double imageStarttime = this->TaskSetting->getSARImageStartTime(); + std::shared_ptr sateOirbtNodes = this->getSatelliteOribtNodes(prf_time, dt, antflag, imageStarttime); + + + + +#pragma omp parallel for + for (int devid = 0; devid < num_devices; devid++) { + cudaSetDevice(devid); // 纭繚褰撳墠绾跨▼鎿嶄綔鎸囧畾鐨凣PU璁惧 + this->RFPCMainProcess_GPU_NoAntPattern(0, 0, devid); + size_t startTid = devid * prfblockcount; + size_t prf_devLen = prfblockcount; + prf_devLen = (startTid + prf_devLen) < prfcount ? prf_devLen : (prfcount - startTid); + this->RFPCMainProcess_GPU_NoAntPattern(startTid, prf_devLen, devid); + } + + + return ErrorCode::SUCCESS; +} + +ErrorCode RFPCProcessCls::RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, size_t prfcount, int devId) +{ + /// 鏄惧瓨涓嶉檺鍒 + cudaSetDevice(devId); // 纭繚褰撳墠绾跨▼鎿嶄綔鎸囧畾鐨凣PU璁惧 + POLARTYPEENUM polartype = this->TaskSetting->getPolarType(); + std::map clssigmaParamsDict = this->SigmaDatabasePtr->getsigmaParams(polartype);; + std::map clsCUDASigmaParamsDict; + for (const auto& pair : clssigmaParamsDict) { + clsCUDASigmaParamsDict.insert(std::pair(pair.first, + CUDASigmaParam{ + pair.second.p1, + pair.second.p2, + pair.second.p3, + pair.second.p4, + pair.second.p5, + pair.second.p6 + })); + } + + + // 璇诲彇绫诲埆 + gdalImage demxyz(this->demxyzPath);// 鍦伴潰鐐瑰潗鏍 + gdalImage demlandcls(this->LandCoverPath);// 鍦拌〃瑕嗙洊绫诲瀷 + gdalImage slpxyz(this->demsloperPath);// 鍦伴潰鍧″悜 + // 澶勭悊鍦伴潰鍧愭爣 + long demRow = demxyz.height; + long demCol = demxyz.width; + size_t demCount = size_t(demRow) * size_t(demCol); + + std::shared_ptr demX = readDataArr(demxyz, 0, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr demY = readDataArr(demxyz, 0, 0, demRow, demCol, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr demZ = readDataArr(demxyz, 0, 0, demRow, demCol, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr slpX = readDataArr(slpxyz, 0, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr slpY = readDataArr(slpxyz, 0, 0, demRow, demCol, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr slpZ = readDataArr(slpxyz, 0, 0, demRow, demCol, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr clsArr = readDataArr(demlandcls, 0, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + + + // 妫绱㈢被鍒暟閲 + std::map clsCountDict; + for (const auto& pair : clssigmaParamsDict) { + clsCountDict.insert(std::pair(pair.first, 0)); + } + + for (size_t i = 0; i < demCount; i++) { + long clsid = clsArr.get()[i]; + if (clsCountDict.find(clsid) != clsCountDict.end()) { + clsCountDict[clsid] = clsCountDict[clsid] + 1; + } + } + + std::map> clsGoalStateDict; + for (const auto& pair : clsCountDict) { + clsGoalStateDict.insert( + std::pair>( + pair.first, + std::shared_ptr((GoalState*)mallocCUDAHost(sizeof(GoalState) * pair.second), FreeCUDAHost))); + } + + // 鍒嗗潡澶勭悊澶у皬 + size_t blocksize = 1000; + + std::map clsCountDictTemp; + for (const auto& pair : clsCountDict) { + clsCountDictTemp.insert(std::pair(pair.first, pair.second)); + } + + for (long i = 0; i < demCount; i++) { + long clsid = clsArr.get()[i]; + size_t Currentclscount = clsCountDictTemp[clsid]; + size_t allclscount = clsCountDict[clsid]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].Tx = demX.get()[i]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].Ty = demY.get()[i]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].Tz = demZ.get()[i]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].TsX = slpX.get()[i]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].TsY = slpY.get()[i]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].TsZ = slpZ.get()[i]; + clsGoalStateDict[clsid].get()[Currentclscount-allclscount].cls = clsArr.get()[i]; + Currentclscount = Currentclscount - 1; + } + + + + RFPCTask task; + // 鍙傛暟澹版槑 + task.freqNum = this->EchoSimulationData->getPlusePoints(); + task.prfNum = prfcount; + task.Rref = this->EchoSimulationData->getRefPhaseRange(); + task.Rnear = this->EchoSimulationData->getNearRange(); + task.Rfar = this->EchoSimulationData->getFarRange(); + task.startFreq = this->EchoSimulationData->getCenterFreq() - this->EchoSimulationData->getBandwidth() / 2; + task.stepFreq = this->EchoSimulationData->getBandwidth() / (task.freqNum - 1); + task.d_echoData = (cuComplex*)mallocCUDADevice(prfcount * task.freqNum * sizeof(cuComplex), devId); + + // 澶╃嚎浣嶇疆 + { + std::shared_ptr antplise = this->EchoSimulationData->getAntPosVelc(); + std::shared_ptr h_antlist((SateState*)mallocCUDAHost(prfcount * sizeof(SateState)), FreeCUDAHost); + + for (long i = 0; i < prfcount; i++) { + h_antlist.get()[i].Px = antplise.get()[i + startprfid].Px; + h_antlist.get()[i].Py = antplise.get()[i + startprfid].Py; + h_antlist.get()[i].Pz = antplise.get()[i + startprfid].Pz; + h_antlist.get()[i].Vx = antplise.get()[i + startprfid].Vx; + h_antlist.get()[i].Vy = antplise.get()[i + startprfid].Vy; + h_antlist.get()[i].Vz = antplise.get()[i + startprfid].Vz; + } + + task.antlist = (SateState*)mallocCUDADevice(prfcount * sizeof(SateState), devId); + HostToDevice(h_antlist.get(), task.antlist, sizeof(double) * prfcount); + + } + + // 鍒嗗潡璁＄畻 + for (const auto& pair : clsCUDASigmaParamsDict) { + long clsid = pair.first; + size_t clscount = clsCountDict[clsid]; + PRINT("Process Class ID : %d , Count: %d\n", clsid, clscount); + task.targetnum = clscount; + task.goallist = (GoalState*)mallocCUDADevice(clscount * sizeof(GoalState), devId); + HostToDevice(clsGoalStateDict[clsid].get(), task.goallist, sizeof(GoalState) * clscount); + task.sigma0_cls = pair.second; + ProcessRFPCTask(task); + FreeCUDADevice(task.goallist); + } + + + + // 鏂囦欢璇诲啓 + omp_lock_t lock; + omp_init_lock(&lock); + omp_set_lock(&lock); + + + cuComplex* h_echoData = (cuComplex*)mallocCUDAHost(prfcount * task.freqNum * sizeof(cuComplex)); + DeviceToHost(h_echoData, task.d_echoData, prfcount* task.freqNum * sizeof(cuComplex)); + + long prfcount_read = prfcount; + std::shared_ptr> fileEchoArr = this->EchoSimulationData->getEchoArr(startprfid, prfcount_read); + + for (size_t i = 0; i < prfcount; i++) { + for (size_t j = 0; j < task.freqNum; j++) { + std::complex temp = fileEchoArr.get()[i * task.freqNum + j]; + fileEchoArr.get()[i * task.freqNum + j] = std::complex( + temp.real() + h_echoData[i * task.freqNum + j].x, + temp.imag() + h_echoData[i * task.freqNum + j].y + ); + } + } + + + this->EchoSimulationData->saveEchoArr(fileEchoArr, startprfid, prfcount_read); + + + omp_unset_lock(&lock); // 閿熼叺鏀句紮鎷锋枻鎷 + omp_destroy_lock(&lock); // 鍔紮鎷锋枻鎷 + + FreeCUDADevice(task.d_echoData); + FreeCUDADevice(task.antlist); + //FreeCUDADevice(task.goallist); + + return ErrorCode::SUCCESS; +} + diff --git a/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.h b/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.h index 75c9667..0dafc01 100644 --- a/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.h +++ b/Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.h @@ -79,6 +79,10 @@ private: // ErrorCode InitEchoMaskArray(); //ErrorCode RFPCMainProcess(long num_thread); ErrorCode RFPCMainProcess_GPU(); + ErrorCode RFPCMainProcess_MultiGPU_NoAntPattern(); // 多GPU处理,不考虑天线方向图 + + ErrorCode RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, size_t prfcount,int devId=0); + std::shared_ptr getSatelliteOribtNodes(double prf_time, double dt, bool antflag, long double imageStarttime); diff --git a/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj b/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj index e18556f..1072389 100644 --- a/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj +++ b/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj @@ -226,6 +226,7 @@ + diff --git a/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj.filters b/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj.filters index c565244..6ca9c53 100644 --- a/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj.filters +++ b/Toolbox/SimulationSARTool/SimulationSARTool.vcxproj.filters @@ -217,5 +217,8 @@ SimulationSAR + + SimulationSAR + \ No newline at end of file