KJ135
/
RasterProcessTool
mirror of http://172.16.0.12:5000/LAMPSARToolSoftware/RasterProcessTool.git
4
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'RFPC-dev' of http://172.16.0.12:5000/LAMPSARToolSoftware/RasterProcessTool into RFPC-dev

Release-dev
chenzenghui 2025-04-08 00:14:32 +08:00
parent c45fedc513 947146320e
commit a0d3e68035
5 changed files with 235 additions and 293 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
BaseCommonLibrary/ImageOperatorFuntion.cpp
View File

@ -1153,7 +1153,7 @@ int ResampleGDAL(const char* pszSrcFile, const char* pszOutFile, double* gt, int
GDALWarpOptions* psWo = GDALCreateWarpOptions();
CPLSetConfigOption("GDAL_NUM_THREADS", "ALL_CPUS"); // 使用所有可用的CPU核心
CPLSetConfigOption("GDAL_CACHEMAX", "16000"); // 设置缓存大小为500MB
CPLSetConfigOption("GDAL_CACHEMAX", "4000"); // 设置缓存大小为500MB
// psWo->papszWarpOptions = CSLDuplicate(NULL);
psWo->eWorkingDataType = dataType;
psWo->eResampleAlg = eResample;

4
Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cu
View File

@ -110,7 +110,7 @@ __global__ void processPulseKernel(
im_final[idx].x += phCorr.x;
im_final[idx].y += phCorr.y;
//printf("r_start=%e;dr=%e;nR=%d\n", r_start, dr, nR);
if (abs(phCorr.x) > 1e-100 || abs(phCorr.y > 1e-100)) {
//if (abs(phCorr.x) > 1e-100 || abs(phCorr.y > 1e-100)) {
//printf(
// "[DEBUG] prfid=%-4ld | idx=%-8lld\n"
// " Ant: X=%-18.10e Y=%-18.10e Z=%-18.10e\n"
@ -133,7 +133,7 @@ __global__ void processPulseKernel(
// phCorr.x, phCorr.y,
// im_final[idx].x, im_final[idx].y
//);
}
//}
}
void bpBasic0CUDA(GPUDATA& data, int flag,double* h_R) {

105
Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cu
View File

@ -519,12 +519,12 @@ __global__ void Kernel_Computer_R_amp_NoAntPattern(
RstY = RstY / RstR;
RstZ = RstZ / RstR;
double slopeX = gp.TsX;
double slopeY = gp.TsY;
double slopeZ = gp.TsZ;
float slopeX = gp.TsX;
float slopeY = gp.TsY;
float slopeZ = gp.TsZ;
double slopR = sqrtf(slopeX * slopeX + slopeY * slopeY + slopeZ * slopeZ); //
if (abs(slopR - 0) > 1e-3) {
float slopR = sqrtf(slopeX * slopeX + slopeY * slopeY + slopeZ * slopeZ); //
if (slopR > 1e-3) {
float localangle = acosf((RstX * slopeX + RstY * slopeY + RstZ * slopeZ) / ( slopR));
@ -553,10 +553,11 @@ __global__ void Kernel_Computer_R_amp_NoAntPattern(
ampGain=2 * maxGain * (1 - (powf(diectAngle,2) / 6)
+ (powf(diectAngle, 4) / 120)
- (powf(diectAngle, 6) / 5040)); //dB
ampGain = powf(10.0, ampGain / 10.0);
ampGain = ampGain / (PI4POW2 * powf(RstR, 4)); // 反射强度
double sigma = GPU_getSigma0dB(sigma0Params, localangle);
float sigma = GPU_getSigma0dB(sigma0Params, localangle);
sigma = powf(10.0, sigma / 10.0);
double temp_amp = double(ampGain * Pt * sigma);
@ -573,76 +574,6 @@ __global__ void Kernel_Computer_R_amp_NoAntPattern(
}
}
__global__ void CUDA_Kernel_Computer_echo_NoAntPattern(
double* d_temp_R, double* d_temp_amps, long posNum,
double f0, double dfreq,
long FreqPoints, // 当前频率的分块
long maxfreqnum, // 最大脉冲值
cuComplex* echodata,
long temp_PRF_Count
) {
__shared__ float s_R[SHAREMEMORY_FLOAT_HALF]; // 注意一个完整的block_size 共享相同内存
__shared__ float s_amp[SHAREMEMORY_FLOAT_HALF];
long long tid = threadIdx.x;
long long bid = blockIdx.x;
long long idx = bid * blockDim.x + tid;
long long prfId = idx / FreqPoints; // 脉冲ID
long long fId = idx % FreqPoints;//频率ID
long long psid = 0;
long long pixelId = 0;
for (long ii = 0; ii < SHAREMEMORY_FLOAT_HALF_STEP; ii++) { // SHAREMEMORY_FLOAT_HALF_STEP * BLOCK_SIZE=SHAREMEMORY_FLOAT_HALF
psid = tid * SHAREMEMORY_FLOAT_HALF_STEP + ii;
pixelId = prfId * posNum + psid; //
if (psid < posNum) {
s_R[psid] = d_temp_R[pixelId];
s_amp[psid] = d_temp_amps[pixelId];
}
else {
s_R[psid] = 0;
s_amp[psid] = 0;
}
}
__syncthreads(); // 确定所有待处理数据都已经进入程序中
if (fId < maxfreqnum && prfId < temp_PRF_Count) {
long echo_ID = prfId * maxfreqnum + fId; // 计算对应的回波位置
float factorjTemp = RFPCPIDIVLIGHT * (f0 + fId * dfreq);
cuComplex echo = make_cuComplex(0, 0);
float temp_phi = 0;
float temp_amp = 0;
for (long dataid = 0; dataid < SHAREMEMORY_FLOAT_HALF; dataid++) {
temp_phi = s_R[dataid] * factorjTemp;
temp_amp = s_amp[dataid];
echo.x += (temp_amp * cosf(temp_phi));
echo.y += (temp_amp * sinf(temp_phi));
//if (dataid > 5000) {
// printf("echo_ID=%d; dataid=%d;ehodata=(%f,%f);R=%f;amp=%f;\n", echo_ID, dataid, temp_real, temp_imag, s_R[0], s_amp[0]);
//}
//if (isnan(temp_phi) || isnan(temp_amp) || isnan(echo.x) || isnan(echo.y)
// || isinf(temp_phi) || isinf(temp_amp) || isinf(echo.x) || isinf(echo.y)
// ) {
// printf("[amp,phi,real,imag]=[%f,%f,%f,%f];\n", temp_amp, temp_phi, echo.x, echo.y);
//}
}
echodata[echo_ID] = cuCaddf(echodata[echo_ID], echo);
}
}
__global__ void CUDA_Kernel_Computer_echo_NoAntPattern_Optimized(
double* d_temp_R, double* d_temp_amps, long posNum,
double f0, double dfreq,
@ -652,15 +583,15 @@ __global__ void CUDA_Kernel_Computer_echo_NoAntPattern_Optimized(
long temp_PRF_Count
) {
// 使用动态共享内存，根据线程块大小调整
extern __shared__ float s_data[];
float* s_R = s_data;
float* s_amp = s_data + blockDim.x;
extern __shared__ double s_data[];
double* s_R = s_data;
double* s_amp = s_data + blockDim.x;
const int tid = threadIdx.x;
const int prfId = blockIdx.x;
const int fId = tid; // 每个线程处理一个频率点
float factorjTemp = RFPCPIDIVLIGHT * (f0 + fId * dfreq);
double factorjTemp = RFPCPIDIVLIGHT * (f0 + fId * dfreq);
cuComplex echo = make_cuComplex(0.0f, 0.0f);
// 分块加载数据并计算
@ -670,8 +601,8 @@ __global__ void CUDA_Kernel_Computer_echo_NoAntPattern_Optimized(
// 加载当前块到共享内存
if (psid < posNum) {
s_R[tid] = static_cast<float>(d_temp_R[pixelId]);
s_amp[tid] = static_cast<float>(d_temp_amps[pixelId]);
s_R[tid] = static_cast<double>(d_temp_R[pixelId]);
s_amp[tid] = static_cast<double>(d_temp_amps[pixelId]);
}
else {
s_R[tid] = 0.0f;
@ -681,7 +612,7 @@ __global__ void CUDA_Kernel_Computer_echo_NoAntPattern_Optimized(
// 计算当前块的贡献
for (int dataid = 0; dataid < blockDim.x; ++dataid) {
float temp_phi = s_R[dataid] * factorjTemp;
float temp_phi =fmod( s_R[dataid] * factorjTemp,2*PI);
float temp_amp = s_amp[dataid];
float sin_phi, cos_phi;
sincosf(temp_phi, &sin_phi, &cos_phi);
@ -743,7 +674,7 @@ extern "C" void ProcessRFPCTask(RFPCTask& task, long devid)
dim3 blocks(task.prfNum);
dim3 threads(BLOCK_SIZE);
size_t shared_mem_size = 2 * BLOCK_SIZE * sizeof(float);
size_t shared_mem_size = 2 * BLOCK_SIZE * sizeof(double);
CUDA_Kernel_Computer_echo_NoAntPattern_Optimized << <blocks, threads, shared_mem_size >> > (
d_R, d_amps, SHAREMEMORY_FLOAT_HALF,
@ -761,11 +692,11 @@ extern "C" void ProcessRFPCTask(RFPCTask& task, long devid)
// task.d_echoData,
// task.prfNum
// );
//PrintLasterError("CUDA_Kernel_Computer_echo");
PrintLasterError("CUDA_Kernel_Computer_echo");
cudaDeviceSynchronize();
if ((sTi * 100.0 / task.targetnum) - process >= 1) {
if ((sTi * 100.0 / task.targetnum) - process >= 10) {
process = sTi * 100.0 / task.targetnum;
PRINT("TargetID [%f]: %d / %d finished %d\n", sTi * 100.0 / task.targetnum, sTi, task.targetnum,devid);
PRINT("device ID : %d , TargetID [%f]: %d / %d finished %d\n",devid, sTi * 100.0 / task.targetnum, sTi, task.targetnum,devid);
}
}

2
Toolbox/SimulationSARTool/SimulationSAR/GPURFPC.cuh
View File

@ -91,7 +91,7 @@ extern "C" struct RFPCTask
cuComplex* d_echoData = nullptr; // »Ø²¨
CUDASigmaParam sigma0_cls;
double maxGain=48;
double GainWeight=20; // 2śČˇśÎ§
double GainWeight=10; // 2śČˇśÎ§
size_t targetnum;

37
Toolbox/SimulationSARTool/SimulationSAR/RFPCProcessCls.cpp
View File

@ -993,19 +993,25 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, si
gdalImage demxyz(this->demxyzPath);// 地面点坐标
gdalImage demlandcls(this->LandCoverPath);// 地表覆盖类型
gdalImage slpxyz(this->demsloperPath);// 地面坡向
// 处理地面坐标
long demRow = demxyz.height;
long allDemRow = Memory1MB/demxyz.width/8/3*6000;
//allDemRow = allDemRow < demxyz.height ? allDemRow : demxyz.height;
for(long demId=0;demId< demxyz.height;demId=demId+ allDemRow){
PRINT("dem cover processbar: [%f precent]\n", demId * 100.0 / demxyz.height);
long demRow = allDemRow;
demRow = demRow + demId < demxyz.height ? demRow : demxyz.height - demId;
long demCol = demxyz.width;
size_t demCount = size_t(demRow) * size_t(demCol);
std::shared_ptr<double> demX = readDataArr<double>(demxyz, 0, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> demY = readDataArr<double>(demxyz, 0, 0, demRow, demCol, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> demZ = readDataArr<double>(demxyz, 0, 0, demRow, demCol, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> slpX = readDataArr<double>(slpxyz, 0, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> slpY = readDataArr<double>(slpxyz, 0, 0, demRow, demCol, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> slpZ = readDataArr<double>(slpxyz, 0, 0, demRow, demCol, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<long> clsArr = readDataArr<long>(demlandcls, 0, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
long long demCount = long long(demRow) * long long(demCol);
std::shared_ptr<double> demX = readDataArr<double>(demxyz, demId, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> demY = readDataArr<double>(demxyz, demId, 0, demRow, demCol, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> demZ = readDataArr<double>(demxyz, demId, 0, demRow, demCol, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> slpX = readDataArr<double>(slpxyz, demId, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> slpY = readDataArr<double>(slpxyz, demId, 0, demRow, demCol, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> slpZ = readDataArr<double>(slpxyz, demId, 0, demRow, demCol, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<long> clsArr = readDataArr<long>(demlandcls, demId, 0, demRow, demCol, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
PRINT("demRow: %d , demCol:%d \n", demRow, demCol);
// 检索类别数量
std::map<long, size_t> clsCountDict;
@ -1013,7 +1019,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, si
clsCountDict.insert(std::pair<long, size_t>(pair.first, 0));
}
for (size_t i = 0; i < demCount; i++) {
for (long long i = 0; i < demCount; i++) {
long clsid = clsArr.get()[i];
if (clsCountDict.find(clsid) != clsCountDict.end()) {
clsCountDict[clsid] = clsCountDict[clsid] + 1;
@ -1081,6 +1087,8 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, si
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);
CUDA_MemsetBlock(task.d_echoData, make_cuComplex(0, 0), prfcount * task.freqNum);
PRINT("Dev:%d ,freqnum：%d , prfnum:%d ,Rref: %e ,Rnear : %e ,Rfar: %e , StartFreq: %e ,DeletFreq: %e \n",
devId, task.freqNum, task.prfNum, task.Rref, task.Rnear, task.Rfar, task.startFreq, task.stepFreq);
@ -1113,7 +1121,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, si
for (const auto& pair : clsGoalStateDict) {
long clsid = pair.first;
size_t clscount = clsCountDict[clsid];
PRINT("Process Class ID : %d , Count: %d\n", clsid, clscount);
PRINT("Process Class ID : %d , Count: %d Device: %d\n", clsid, clscount,devId);
task.targetnum = clscount;
task.goallist = (GoalState*)mallocCUDADevice(clscount * sizeof(GoalState), devId);
HostToDevice(clsGoalStateDict[clsid].get(), task.goallist, sizeof(GoalState) * clscount);
@ -1130,6 +1138,9 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU_NoAntPattern(size_t startprfid, si
FreeCUDADevice(task.antlist);
//FreeCUDADevice(task.goallist);
}
PRINT("dem cover processbar: [100 precent]\n");
return ErrorCode::SUCCESS;
}