diff --git a/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp b/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp index dab330b..e78b9d8 100644 --- a/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp +++ b/BaseCommonLibrary/BaseTool/EchoDataFormat.cpp @@ -251,6 +251,16 @@ QString EchoL0Dataset::getEchoDataFilename() return GPSPointFilePath; } +QString EchoL0Dataset::getGPSPointFilePath() +{ + return this->GPSPointFilePath; +} + +QString EchoL0Dataset::getEchoDataFilePath() +{ + return this->echoDataFilePath; +} + void EchoL0Dataset::initEchoArr(std::complex init0) { long blockline = Memory1MB / 8 / 2 / this->PlusePoints * 8000; @@ -676,26 +686,26 @@ std::shared_ptr SatelliteAntPosOperator::readAntPosFile(QString long colcount = 19; std::shared_ptr antlist = readDataArr(antimg, 0, 0, rowcount, colcount, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); std::shared_ptr antpos(new SatelliteAntPos[rowcount], delArrPtr); - for (long i = 0; i < colcount; i++) { - antpos.get()[i].time = antlist.get()[i * 19 + 1]; - antpos.get()[i].Px = antlist.get()[i * 19 + 2]; - antpos.get()[i].Py = antlist.get()[i * 19 + 3]; - antpos.get()[i].Pz = antlist.get()[i * 19 + 4]; - antpos.get()[i].Vx = antlist.get()[i * 19 + 5]; - antpos.get()[i].Vy = antlist.get()[i * 19 + 6]; - antpos.get()[i].Vz = antlist.get()[i * 19 + 7]; //7 - antpos.get()[i].AntDirectX = antlist.get()[i * 19 + 8]; - antpos.get()[i].AntDirectY = antlist.get()[i * 19 + 9]; - antpos.get()[i].AntDirectZ = antlist.get()[i * 19 + 10]; - antpos.get()[i].AVx = antlist.get()[i * 19 + 11]; - antpos.get()[i].AVy = antlist.get()[i * 19 + 12]; - antpos.get()[i].AVz = antlist.get()[i * 19 + 13]; - antpos.get()[i].ZeroAntDiectX = antlist.get()[i * 19 + 14]; - antpos.get()[i].ZeroAntDiectY = antlist.get()[i * 19 + 15]; - antpos.get()[i].ZeroAntDiectZ = antlist.get()[i * 19 + 16]; - antpos.get()[i].lon = antlist.get()[i * 19 + 17]; - antpos.get()[i].lat = antlist.get()[i * 19 + 18]; - antpos.get()[i].ati = antlist.get()[i * 19 + 19]; // 19 + for (long i = 0; i < rowcount; i++) { + antpos.get()[i].time = antlist.get()[i * 19 + 0]; + antpos.get()[i].Px = antlist.get()[i * 19 + 1]; + antpos.get()[i].Py = antlist.get()[i * 19 + 2]; + antpos.get()[i].Pz = antlist.get()[i * 19 + 3]; + antpos.get()[i].Vx = antlist.get()[i * 19 + 4]; + antpos.get()[i].Vy = antlist.get()[i * 19 + 5]; + antpos.get()[i].Vz = antlist.get()[i * 19 + 6]; //7 + antpos.get()[i].AntDirectX = antlist.get()[i * 19 + 7]; + antpos.get()[i].AntDirectY = antlist.get()[i * 19 + 8]; + antpos.get()[i].AntDirectZ = antlist.get()[i * 19 + 9]; + antpos.get()[i].AVx = antlist.get()[i * 19 + 10]; + antpos.get()[i].AVy = antlist.get()[i * 19 + 11]; + antpos.get()[i].AVz = antlist.get()[i * 19 + 12]; + antpos.get()[i].ZeroAntDiectX = antlist.get()[i * 19 + 13]; + antpos.get()[i].ZeroAntDiectY = antlist.get()[i * 19 + 14]; + antpos.get()[i].ZeroAntDiectZ = antlist.get()[i * 19 + 15]; + antpos.get()[i].lon = antlist.get()[i * 19 + 16]; + antpos.get()[i].lat = antlist.get()[i * 19 + 17]; + antpos.get()[i].ati = antlist.get()[i * 19 + 18]; // 19 } return antpos; } diff --git a/BaseCommonLibrary/BaseTool/EchoDataFormat.h b/BaseCommonLibrary/BaseTool/EchoDataFormat.h index 49fea34..ca0e5ce 100644 --- a/BaseCommonLibrary/BaseTool/EchoDataFormat.h +++ b/BaseCommonLibrary/BaseTool/EchoDataFormat.h @@ -132,7 +132,8 @@ public: QString getxmlName(); QString getGPSPointFilename(); QString getEchoDataFilename(); - + QString getGPSPointFilePath(); + QString getEchoDataFilePath(); void initEchoArr(std::complex init0); diff --git a/GPUBaseLib/GPUTool/GPUTool.cu b/GPUBaseLib/GPUTool/GPUTool.cu index be15ded..300acf7 100644 --- a/GPUBaseLib/GPUTool/GPUTool.cu +++ b/GPUBaseLib/GPUTool/GPUTool.cu @@ -252,7 +252,7 @@ extern "C" void FFTShift1D(cuComplex* d_data, int batch_size, int signal_length) cudaDeviceSynchronize(); } -extern "C" void shared_complexPtrToHostCuComplex(std::complex* src, cuComplex* dst, long len) +extern "C" void shared_complexPtrToHostCuComplex(std::complex* src, cuComplex* dst, size_t len) { for (long i = 0; i < len; i++) { dst[i] = make_cuComplex(src[i].real(), src[i].imag()); @@ -260,7 +260,7 @@ extern "C" void shared_complexPtrToHostCuComplex(std::complex* src, cu return ; } -extern "C" void HostCuComplexToshared_complexPtr( cuComplex* src, std::complex* dst, long len) +extern "C" void HostCuComplexToshared_complexPtr( cuComplex* src, std::complex* dst, size_t len) { double maxvalue = src[0].x; for (long i = 0; i < len; i++) { @@ -318,14 +318,14 @@ extern "C" void checkCudaError(cudaError_t err, const char* msg) { } // 主机参数内存声明 -extern "C" void* mallocCUDAHost(long memsize) { +extern "C" void* mallocCUDAHost(size_t memsize) { void* ptr; cudaMallocHost(&ptr, memsize); #ifdef __CUDADEBUG__ cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { - printf("mallocCUDAHost CUDA Error: %s\n", cudaGetErrorString(err)); + printf("mallocCUDAHost CUDA Error: %s, malloc memory : %d byte\n", cudaGetErrorString(err),memsize); exit(2); } #endif // __CUDADEBUG__ @@ -344,7 +344,7 @@ extern "C" void FreeCUDAHost(void* ptr) { #ifdef __CUDADEBUG__ cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { - printf("FreeCUDAHost CUDA Error: %s\n", cudaGetErrorString(err)); + printf("FreeCUDAHost CUDA Error: %s,\n", cudaGetErrorString(err)); exit(2); } #endif // __CUDADEBUG__ @@ -354,13 +354,13 @@ extern "C" void FreeCUDAHost(void* ptr) { } // GPU参数内存声明 -extern "C" void* mallocCUDADevice(long memsize) { +extern "C" void* mallocCUDADevice(size_t memsize) { void* ptr; cudaMalloc(&ptr, memsize); #ifdef __CUDADEBUG__ cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { - printf("mallocCUDADevice CUDA Error: %s\n", cudaGetErrorString(err)); + printf("mallocCUDADevice CUDA Error: %s, malloc memory : %d byte\n", cudaGetErrorString(err), memsize); exit(2); } #endif // __CUDADEBUG__ @@ -386,7 +386,7 @@ extern "C" void FreeCUDADevice(void* ptr) { } // GPU 内存数据转移 -extern "C" void HostToDevice(void* hostptr, void* deviceptr, long memsize) { +extern "C" void HostToDevice(void* hostptr, void* deviceptr, size_t memsize) { cudaMemcpy(deviceptr, hostptr, memsize, cudaMemcpyHostToDevice); #ifdef __CUDADEBUG__ @@ -400,7 +400,7 @@ extern "C" void HostToDevice(void* hostptr, void* deviceptr, long memsize) { cudaDeviceSynchronize(); } -extern "C" void DeviceToHost(void* hostptr, void* deviceptr, long memsize) { +extern "C" void DeviceToHost(void* hostptr, void* deviceptr, size_t memsize) { cudaMemcpy(hostptr, deviceptr, memsize, cudaMemcpyDeviceToHost); #ifdef __CUDADEBUG__ cudaError_t err = cudaGetLastError(); @@ -412,7 +412,7 @@ extern "C" void DeviceToHost(void* hostptr, void* deviceptr, long memsize) { cudaDeviceSynchronize(); } -void DeviceToDevice(void* s_deviceptr, void* t_deviceptr, long memsize) +void DeviceToDevice(void* s_deviceptr, void* t_deviceptr, size_t memsize) { cudaMemcpy(t_deviceptr, s_deviceptr, memsize, cudaMemcpyDeviceToDevice); #ifdef __CUDADEBUG__ @@ -585,6 +585,7 @@ void PrintLasterError(const char* s) extern "C" void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long InRowCount, long InColCount, long outColCount) { + cufftHandle plan; cufftResult result; @@ -612,7 +613,7 @@ extern "C" void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long InRowCount, cuComplex* in_ptr = inArr; cuComplex* out_ptr = outArr; result = cufftExecC2C(plan, (cufftComplex*)in_ptr, (cufftComplex*)out_ptr, CUFFT_INVERSE); - + if (result != CUFFT_SUCCESS) { cufftDestroy(plan); return; @@ -622,7 +623,8 @@ extern "C" void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long InRowCount, cudaDeviceSynchronize(); cufftDestroy(plan); -} + +} diff --git a/GPUBaseLib/GPUTool/GPUTool.cuh b/GPUBaseLib/GPUTool/GPUTool.cuh index cdaf903..d16af3c 100644 --- a/GPUBaseLib/GPUTool/GPUTool.cuh +++ b/GPUBaseLib/GPUTool/GPUTool.cuh @@ -75,13 +75,13 @@ extern "C" GPUBASELIBAPI void printDeviceInfo(int deviceId); extern "C" GPUBASELIBAPI void checkCudaError(cudaError_t err, const char* msg); // GPU 内存函数 -extern "C" GPUBASELIBAPI void* mallocCUDAHost(long memsize); // 主机内存声明 +extern "C" GPUBASELIBAPI void* mallocCUDAHost(size_t memsize); // 主机内存声明 extern "C" GPUBASELIBAPI void FreeCUDAHost(void* ptr); -extern "C" GPUBASELIBAPI void* mallocCUDADevice(long memsize); // GPU内存声明 +extern "C" GPUBASELIBAPI void* mallocCUDADevice(size_t memsize); // GPU内存声明 extern "C" GPUBASELIBAPI void FreeCUDADevice(void* ptr); -extern "C" GPUBASELIBAPI void HostToDevice(void* hostptr, void* deviceptr, long memsize);//GPU 内存数据转移 设备 -> GPU -extern "C" GPUBASELIBAPI void DeviceToHost(void* hostptr, void* deviceptr, long memsize);//GPU 内存数据转移 GPU -> 设备 -extern "C" GPUBASELIBAPI void DeviceToDevice(void* s_deviceptr, void* t_deviceptr, long memsize);//GPU 内存数据转移 GPU -> 设备 +extern "C" GPUBASELIBAPI void HostToDevice(void* hostptr, void* deviceptr, size_t memsize);//GPU 内存数据转移 设备 -> GPU +extern "C" GPUBASELIBAPI void DeviceToHost(void* hostptr, void* deviceptr, size_t memsize);//GPU 内存数据转移 GPU -> 设备 +extern "C" GPUBASELIBAPI void DeviceToDevice(void* s_deviceptr, void* t_deviceptr, size_t memsize);//GPU 内存数据转移 GPU -> 设备 extern "C" GPUBASELIBAPI void CUDA_MemsetBlock(cuComplex* data, cuComplex init0, long len); // 矢量基础运算函数 @@ -109,7 +109,7 @@ extern "C" GPUBASELIBAPI void CUDAIFFTScale(cuComplex* inArr, cuComplex* outArr, extern "C" GPUBASELIBAPI void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long InRowCount, long InColCount, long outColCount); 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, long len); -extern "C" GPUBASELIBAPI void HostCuComplexToshared_complexPtr(cuComplex* src, std::complex* dst, long len); +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/ImageshowTool/Imageshow/ImageShowDialogClass.cpp b/ImageshowTool/Imageshow/ImageShowDialogClass.cpp index 40baeb7..f2fa26a 100644 --- a/ImageshowTool/Imageshow/ImageShowDialogClass.cpp +++ b/ImageshowTool/Imageshow/ImageShowDialogClass.cpp @@ -36,12 +36,16 @@ ImageShowDialogClass ::ImageShowDialogClass(QWidget *parent) ImageShowDialogClass::~ImageShowDialogClass() -{} +{ + if (nullptr != this->desCursor) { + this->desCursor->close(); + } +} void ImageShowDialogClass::on_action_cursor_enable_trigged() { - this->desCursor = new ImageShowCursorDesClass(this); + this->desCursor = new ImageShowCursorDesClass; this->desCursorflag = true; for (size_t i = 0; i < this->getGraphCount(); i++) { @@ -313,11 +317,14 @@ void ImageShowDialogClass::updateCursor(QMouseEvent *event) QPoint pos = event->pos(); double x=this->m_plot->xAxis->pixelToCoord(pos.x()); // 灏嗛紶鏍囦綅缃槧灏勫埌鍥捐〃鍧愭爣绯讳腑 double y = this->m_plot->yAxis->pixelToCoord(pos.y()); - this->statusbar->showMessage(u8"X: "+QString::number(x,'f', 6)+" y: "+QString::number(y, 'f', 6)); + QCPColorMap* colorMap = dynamic_cast(this->ui->m_plot->plottable(0)); + double dataValue = colorMap->data()->data(x, y); + this->statusbar->showMessage(u8"X: "+QString::number(x,'f', 6)+" y: "+QString::number(y, 'f', 6)+ + " value: "+QString::number(dataValue,'e',6)); if (this->desCursorflag) { if (this->graphclass == LAMPDATASHOWCLASS::LAMPColorMap) { - QCPColorMap* colorMap = dynamic_cast(this->ui->m_plot->plottable(0)); - double dataValue = colorMap->data()->data(x, y); + + this->desCursor->updateCursorContent(u8"X: " + QString::number(x, 'f', 6) + " y: " + QString::number(y, 'f', 6) +u8"\n" +u8"DataValue: "+QString::number(dataValue)); diff --git a/Toolbox/SimulationSARTool/GPUBpSimulation.cu b/Toolbox/SimulationSARTool/GPUBpSimulation.cu index d92436c..a39ca8f 100644 --- a/Toolbox/SimulationSARTool/GPUBpSimulation.cu +++ b/Toolbox/SimulationSARTool/GPUBpSimulation.cu @@ -56,32 +56,37 @@ __global__ void kernel_RFPCProcess( double minF,double dFreq,double RefRange, cuComplex* phdata ) { - //long prfid = blockIdx.x * blockDim.x + threadIdx.x; // 获取当前的线程编码 - //if (prfid >= Np || prfid < 0) { return; } - //else {} + 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 }; + // 距离 + 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(p1, p2, p3, p4, p5, p6, localIncAngle);// 后向散射系数 + sigma0 = powf(10.0, sigma0 / 10.0); - //// 入射角 - //Vector3 TS = vec_sub(S, T);//T-->S - //double localIncAngle = angleBetweenVectors(TS, slp, false);// 入射角 - //double sigma0 = GPU_getSigma0dB(p1, p2, p3, p4, p5, p6, localIncAngle);// 后向散射系数 - //sigma0 = powf(10.0, sigma0 / 10.0); + // 距离 + double R = sqrt(vec_dot(TS, TS)); - //// 距离 - //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); - //} + // 计算增益 + 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); } @@ -96,9 +101,9 @@ void RFPCProcess(double Tx, double Ty, double Tz, 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); - HostToDevice(d_data.AntY, AntY, sizeof(double) * d_data.Np); - HostToDevice(d_data.AntZ, AntZ, 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; diff --git a/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cu b/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cu index 4c0d849..0c542b3 100644 --- a/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cu +++ b/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cu @@ -31,8 +31,8 @@ __global__ void phaseCompensationKernel(cufftComplex* phdata, const double* Freq int idx = pulseIdx * K + freqIdx; double phase = 4 * PI * Freq[freqIdx] * r / c; - double cos_phase = cosf(phase); - double sin_phase = sinf(phase); + double cos_phase = cos(phase); + double sin_phase = sin(phase); cufftComplex ph = phdata[idx]; double new_real = ph.x * cos_phase - ph.y * sin_phase; @@ -79,8 +79,8 @@ __global__ void processPulseKernel( double dz = AntZ - z_mat[idx]; //printf("processPulseKernel xmat !!\n"); - - double dR = sqrtf(dx * dx + dy * dy + dz * dz) - R0; + double R = sqrt(dx * dx + dy * dy + dz * dz); + double dR = R - R0; if (dR < r_start || dR >= (r_start + dr * (nR - 1))) return; // Linear interpolation @@ -109,29 +109,29 @@ __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-14 || abs(phCorr.y > 1e-14)) { - printf( - "[DEBUG] prfid=%-4ld | idx=%-8lld\n" - " Ant: X=%-18.10e Y=%-18.10e Z=%-18.10e\n" - " Pix: X=%-18.10e Y=%-18.10e Z=%-18.10e\n" - " dR=%-18.10e | pos=%-8.4f[%-6d+%-8.6f]\n" - " RC: low=(%-18.10e,%-18.10e) high=(%-18.10e,%-18.10e)\n" - " => interp=(%-18.10e,%-18.10e)\n" - " Phase: val=%-18.10e | corr=(%-18.10e,%-18.10e)\n" - " Final: im=(%-18.10e,%-18.10e)\n" - "----------------------------------------\n", - prfid, idx, - AntX, AntY, AntZ, - x_mat[idx], y_mat[idx], z_mat[idx], - dR, - pos, index, weight, - rc_low.x, rc_low.y, - rc_high.x, rc_high.y, - rc_interp.x, rc_interp.y, - phase, - phCorr.x, phCorr.y, - im_final[idx].x, im_final[idx].y - ); + 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" + // " Pix: X=%-18.10e Y=%-18.10e Z=%-18.10e\n" + // " R=%-18.10e|dR=%-18.10e | pos=%-8.4f[%-6d+%-8.6f]\n" + // " RC: low=(%-18.10e,%-18.10e) high=(%-18.10e,%-18.10e)\n" + // " => interp=(%-18.10e,%-18.10e)\n" + // " Phase: val=%-18.10e | corr=(%-18.10e,%-18.10e)\n" + // " Final: im=(%-18.10e,%-18.10e)\n" + // "----------------------------------------\n", + // prfid, idx, + // AntX, AntY, AntZ, + // x_mat[idx], y_mat[idx], z_mat[idx], + // R,dR, + // pos, index, weight, + // rc_low.x, rc_low.y, + // rc_high.x, rc_high.y, + // rc_interp.x, rc_interp.y, + // phase, + // phCorr.x, phCorr.y, + // im_final[idx].x, im_final[idx].y + //); } } @@ -159,6 +159,9 @@ void bpBasic0CUDA(GPUDATA& data, int flag,double* h_R) { printf("fft finished!!\n"); // 图像重建 + + + double r_start = data.r_vec[0]; double dr = (data.r_vec[data.Nfft - 1] - r_start) / (data.Nfft - 1); printf("dr = %f\n",dr); @@ -166,8 +169,6 @@ void bpBasic0CUDA(GPUDATA& data, int flag,double* h_R) { long grid_size = (pixelcount + BLOCK_SIZE - 1) / BLOCK_SIZE; printf("grid finished!!\n"); - - //double* d_R = (double*)mallocCUDADevice(sizeof(double) * data.nx * data.ny); printf("r_start=%e;dr=%e;nR=%d\n", r_start, dr, data.Nfft); printf("BPimage .....\n"); @@ -207,20 +208,20 @@ void initGPUData(GPUDATA& h_data, GPUDATA& d_data) { d_data.z_mat = (double*)mallocCUDADevice(sizeof(double) * h_data.nx * h_data.ny); d_data.r_vec = h_data.r_vec;// (double*)mallocCUDADevice(sizeof(double) * h_data.Nfft); d_data.Freq = (double*)mallocCUDADevice(sizeof(double) * h_data.Nfft); - d_data.phdata = (cufftComplex*)mallocCUDADevice(sizeof(cufftComplex) * h_data.K * h_data.Np); - d_data.im_final = (cufftComplex*)mallocCUDADevice(sizeof(cufftComplex) * h_data.nx * h_data.ny); + d_data.phdata = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * h_data.Nfft * h_data.Np); + d_data.im_final = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * h_data.nx * h_data.ny); //HostToDevice(h_data.AntX, d_data.AntX,sizeof(double) * h_data.Np); //HostToDevice(h_data.AntY, d_data.AntY,sizeof(double) * h_data.Np); //HostToDevice(h_data.AntZ, d_data.AntZ,sizeof(double) * h_data.Np); //HostToDevice(h_data.minF, d_data.minF,sizeof(double) * h_data.Np); - HostToDevice(h_data.x_mat, d_data.x_mat,sizeof(double) * h_data.nx * h_data.ny); - HostToDevice(h_data.y_mat, d_data.y_mat,sizeof(double) * h_data.nx * h_data.ny); - HostToDevice(h_data.z_mat, d_data.z_mat,sizeof(double) * h_data.nx * h_data.ny); + HostToDevice(h_data.x_mat, d_data.x_mat,sizeof(double) * h_data.nx * h_data.ny); printf("image X Copy finished!!!\n"); + HostToDevice(h_data.y_mat, d_data.y_mat,sizeof(double) * h_data.nx * h_data.ny); printf("image Y Copy finished!!!\n"); + HostToDevice(h_data.z_mat, d_data.z_mat, sizeof(double) * h_data.nx * h_data.ny); printf("image Z Copy finished!!!\n"); HostToDevice(h_data.Freq, d_data.Freq, sizeof(double) * h_data.Nfft); //HostToDevice(h_data.r_vec, d_data.r_vec, sizeof(double) * h_data.Nfft); - HostToDevice(h_data.phdata, d_data.phdata, sizeof(cufftComplex) * h_data.K * h_data.Np); - HostToDevice(h_data.im_final, d_data.im_final, sizeof(cufftComplex) * h_data.nx * h_data.ny); + HostToDevice(h_data.phdata, d_data.phdata, sizeof(cuComplex) * h_data.Nfft * h_data.Np); printf("image echo Copy finished!!!\n"); + HostToDevice(h_data.im_final, d_data.im_final, sizeof(cuComplex) * h_data.nx * h_data.ny); printf("image data Copy finished!!!\n"); // 拷贝标量参数 d_data.Nfft = h_data.Nfft; diff --git a/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cuh b/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cuh index bfb7e35..d4d93f8 100644 --- a/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cuh +++ b/Toolbox/SimulationSARTool/SimulationSAR/BPBasic0_CUDA.cuh @@ -24,8 +24,8 @@ struct GPUDATA { double* x_mat, * y_mat, * z_mat;// 地面坐标 double* r_vec; // 坐标范围 double* Freq;// 频率 - cufftComplex* phdata;// 回波 - cufftComplex* im_final;// 图像 + cuComplex* phdata;// 回波 + cuComplex* im_final;// 图像 double R0; // 参考斜距 double deltaF; // 频点范围 }; diff --git a/Toolbox/SimulationSARTool/UnitTestMain.cpp b/Toolbox/SimulationSARTool/UnitTestMain.cpp index 35133e0..7a3f62c 100644 --- a/Toolbox/SimulationSARTool/UnitTestMain.cpp +++ b/Toolbox/SimulationSARTool/UnitTestMain.cpp @@ -15,91 +15,107 @@ #include "BPBasic0_CUDA.cuh" #include "ImageOperatorBase.h" #include "GPUBpSimulation.cuh" +#include -int main(int argc, char* argv[]) { + +void testSimualtionEchoPoint() { + GPUDATA h_data{}; /** 1. 轨道 **************************************************************************************************/ qDebug() << u8"1.轨道文件读取中。。。"; - QString inGPSPath = u8"C:\\Users\\30453\\Desktop\\script\\data\\GF3_Simulation.gpspos.data"; + QString inGPSPath = u8"C:\\Users\\30453\\Desktop\\script\\data\\GF3_Simulation.gpspos.data"; long gpspoints = gdalImage(inGPSPath).height; std::shared_ptr antpos = SatelliteAntPosOperator::readAntPosFile(inGPSPath, gpspoints); - std::shared_ptr antX((double*)mallocCUDAHost(sizeof(double) * gpspoints), FreeCUDAHost); - std::shared_ptr antY((double*)mallocCUDAHost(sizeof(double) * gpspoints), FreeCUDAHost); - std::shared_ptr antZ((double*)mallocCUDAHost(sizeof(double) * gpspoints), FreeCUDAHost); - for (long i = 0; i < gpspoints; i++) { - antX.get()[i] = antpos.get()[i].Px; - antY.get()[i] = antpos.get()[i].Py; - antZ.get()[i] = antpos.get()[i].Pz; + h_data.AntX = (double*)mallocCUDAHost(sizeof(double) * gpspoints); + h_data.AntY = (double*)mallocCUDAHost(sizeof(double) * gpspoints); + h_data.AntZ = (double*)mallocCUDAHost(sizeof(double) * gpspoints); + + for (long i = 0; i < gpspoints; i = i + 1) { + h_data.AntX[i] = antpos.get()[i].Px; + h_data.AntY[i] = antpos.get()[i].Py; + h_data.AntZ[i] = antpos.get()[i].Pz; } + //gpspoints = gpspoints / 2; + qDebug() << "GPS points Count:\t" << gpspoints; + qDebug() << "PRF 0:\t " << QString("%1,%2,%3").arg(h_data.AntX[0]).arg(h_data.AntY[0]).arg(h_data.AntZ[0]); + qDebug() << "PRF " << gpspoints - 1 << ":\t " << QString("%1,%2,%3") + .arg(h_data.AntX[gpspoints - 1]) + .arg(h_data.AntY[gpspoints - 1]) + .arg(h_data.AntZ[gpspoints - 1]); /** 2. 成像网格 **************************************************************************************************/ qDebug() << u8"轨道文件读取结束\n2.成像网格读取。。。"; - QString demxyzPath = u8"C:\\Users\\30453\\Desktop\\script\\data\\demxyz.bin"; + QString demxyzPath = u8"D:\\Programme\\vs2022\\RasterMergeTest\\simulationData\\demdataset\\demxyz.bin"; gdalImage demgridimg(demxyzPath); long dem_rowCount = demgridimg.height; long dem_ColCount = demgridimg.width; std::shared_ptr demX = readDataArr(demgridimg, 0, 0, dem_rowCount, dem_ColCount, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); std::shared_ptr demY = readDataArr(demgridimg, 0, 0, dem_rowCount, dem_ColCount, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); std::shared_ptr demZ = readDataArr(demgridimg, 0, 0, dem_rowCount, dem_ColCount, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); - - std::shared_ptr imgX((double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount), FreeCUDAHost); - std::shared_ptr imgY((double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount), FreeCUDAHost); - std::shared_ptr imgZ((double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount), FreeCUDAHost); + + h_data.x_mat = (double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount); + h_data.y_mat = (double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount); + h_data.z_mat = (double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount); + for (long i = 0; i < dem_rowCount; i++) { for (long j = 0; j < dem_ColCount; j++) { - imgX.get()[i * dem_ColCount + j] = demX.get()[i * dem_ColCount + j]; - imgY.get()[i * dem_ColCount + j] = demY.get()[i * dem_ColCount + j]; - imgZ.get()[i * dem_ColCount + j] = demZ.get()[i * dem_ColCount + j]; + h_data.x_mat[i * dem_ColCount + j] = demX.get()[i * dem_ColCount + j]; + h_data.y_mat[i * dem_ColCount + j] = demY.get()[i * dem_ColCount + j]; + h_data.z_mat[i * dem_ColCount + j] = demZ.get()[i * dem_ColCount + j]; } } + qDebug() << "dem row Count:\t" << dem_rowCount; + qDebug() << "dem col Count:\t" << dem_ColCount; + qDebug() << u8"成像网格读取结束"; /** 3. 频率网格 **************************************************************************************************/ qDebug() << u8"3.处理频率"; double centerFreq = 5.3e9; double bandwidth = 40e6; - long freqpoints = 8096; + long freqpoints = 2048; std::shared_ptr freqlist(getFreqPoints_mallocHost(centerFreq - bandwidth / 2, centerFreq + bandwidth / 2, freqpoints), FreeCUDAHost); + std::shared_ptr minF(new double[gpspoints], delArrPtr); + for (long i = 0; i < gpspoints; i++) { + minF.get()[i] = freqlist.get()[0]; + } + h_data.deltaF = bandwidth / (freqpoints - 1); qDebug() << "start Freq:\t" << centerFreq - bandwidth / 2; qDebug() << "end Freq:\t" << centerFreq + bandwidth / 2; qDebug() << "freq points:\t" << freqpoints; - qDebug() << "delta freq:\t" << freqlist.get()[0] - freqlist.get()[1]; + qDebug() << "delta freq:\t" << freqlist.get()[1] - freqlist.get()[0]; qDebug() << u8"频率结束"; /** 4. 初始化回波 **************************************************************************************************/ qDebug() << u8"4.初始化回波"; - std::shared_ptr phdata (createEchoPhase_mallocHost(gpspoints, freqpoints),FreeCUDAHost); - qDebug() << "Azimuth Points:\t"< phdata(createEchoPhase_mallocHost(gpspoints, freqpoints), FreeCUDAHost); + qDebug() << "Azimuth Points:\t" << gpspoints; + qDebug() << "Range Points:\t" << freqpoints; qDebug() << u8"初始化回波结束"; /** 5. 初始化图像 **************************************************************************************************/ qDebug() << u8"5.初始化图像"; - std::shared_ptr im_final(createEchoPhase_mallocHost(dem_rowCount, dem_ColCount), FreeCUDAHost); + h_data.im_final = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * dem_rowCount * dem_ColCount); qDebug() << "Azimuth Points:\t" << gpspoints; qDebug() << "Range Points:\t" << freqpoints; qDebug() << u8"初始化图像结束"; /** 6. 模型参数初始化 **************************************************************************************************/ qDebug() << u8"6.模型参数初始化"; - GPUDATA h_data; - h_data.AntX = antX.get(); - h_data.AntX = antY.get(); - h_data.AntX = antZ.get(); - h_data.x_mat = imgX.get(); - h_data.y_mat = imgY.get(); - h_data.z_mat = imgZ.get(); + h_data.nx = dem_ColCount; + h_data.ny = dem_rowCount; + h_data.Np = gpspoints; h_data.Freq = freqlist.get(); + h_data.minF = minF.get(); h_data.Nfft = freqpoints; h_data.K = h_data.Nfft; h_data.phdata = phdata.get(); - h_data.im_final = im_final.get(); + h_data.R0 = 900000; qDebug() << u8"模型参数结束"; - + /** 7. 目标 **************************************************************************************************/ - double Tx = -2028380.625000, Ty = 4139373.250000, Tz = 4393382.500000; - double Tslx = -2028380.625000, Tsly = 4139373.250000, Tslz = 4393382.500000; - - double p1 = 1, p2 = 0, p3 = 0, p4 = 0, p5 = 0, p6 = 0; + double Tx = -2029086.618142, Ty = 4139594.934504, Tz = 4392846.782027; + double Tslx = -2029086.618142, Tsly = 4139594.934504, Tslz = 4392846.782027; + double p1 = 33.3, p2 = 0, p3 = 0, p4 = 0, p5 = 0, p6 = 0; /** 7. 构建回波 **************************************************************************************************/ GPUDATA d_data; @@ -110,26 +126,295 @@ int main(int argc, char* argv[]) { p1, p2, p3, p4, p5, p6, d_data); - HostToDevice(h_data.phdata, d_data.phdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + /** 8. 展示回波 **************************************************************************************************/ - ImageShowDialogClass* dialog = new ImageShowDialogClass; - std::shared_ptr h_ifftphdata((cuComplex*)mallocCUDAHost(sizeof(cuComplex) * d_data.Np * d_data.Nfft), FreeCUDAHost); - std::shared_ptr d_ifftphdata((cuComplex*)mallocCUDADevice(sizeof(cuComplex) * d_data.Np * d_data.Nfft), FreeCUDADevice); - CUDAIFFT(d_data.phdata, d_ifftphdata.get(), d_data.Np, d_data.Nfft, d_data.Nfft); - FFTShift1D(d_ifftphdata.get(), d_data.Np, d_data.Nfft); - DeviceToHost(h_ifftphdata.get(), d_ifftphdata.get(), sizeof(cuComplex) * d_data.Np * d_data.Nfft); - std::shared_ptr echoAmpArr(new double[d_data.Np * d_data.Nfft], delArrPtr); { - for (long i = 0; i < d_data.Np * d_data.Nfft; i++) { - echoAmpArr.get()[i] = - 20 * std::log10(std::sqrt(std::pow(h_ifftphdata.get()[i].x, 2) + std::pow(h_ifftphdata.get()[i].y, 2))); + DeviceToHost(h_data.phdata, d_data.phdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + cuComplex* h_ifftphdata = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * d_data.Np * d_data.Nfft); + cuComplex* d_ifftphdata = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * d_data.Np * d_data.Nfft); + CUDAIFFT(d_data.phdata, d_ifftphdata, d_data.Np, d_data.Nfft, d_data.Nfft); + FFTShift1D(d_ifftphdata, d_data.Np, d_data.Nfft); + DeviceToHost(h_ifftphdata, d_ifftphdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + std::shared_ptr> ifftdata(new std::complex[d_data.Np * d_data.Nfft], delArrPtr); + { + for (long i = 0; i < d_data.Np; i++) { + for (long j = 0; j < d_data.Nfft; j++) { + ifftdata.get()[i * d_data.Nfft + j] = + std::complex( + h_ifftphdata[i * d_data.Nfft + j].x, + h_ifftphdata[i * d_data.Nfft + j].y); + } + } + } + testOutComplexDoubleArr(QString("echo_ifft.bin"), ifftdata.get(), d_data.Np, d_data.Nfft); + + FreeCUDADevice(d_ifftphdata); + FreeCUDAHost(h_ifftphdata); + + } + /** 9. 成像 **************************************************************************************************/ + + // 计算maxWr（需要先计算deltaF） + double deltaF = h_data.deltaF; // 从输入参数获取 + double maxWr = 299792458.0f / (2.0f * deltaF); + qDebug() << "maxWr :\t" << maxWr; + double* r_vec_host = (double*)mallocCUDAHost(sizeof(double) * h_data.Nfft);// new double[data.Nfft]; + const double step = maxWr / h_data.Nfft; + const double start = -1 * h_data.Nfft / 2.0f * step; + printf("nfft=%d\n", h_data.Nfft); + + for (int i = 0; i < h_data.Nfft; ++i) { + r_vec_host[i] = start + i * step; + } + + h_data.r_vec = r_vec_host; + d_data.r_vec = h_data.r_vec; + + qDebug() << "r_vec [0]:\t" << h_data.r_vec[0]; + qDebug() << "r_vec step:\t" << step; + + bpBasic0CUDA(d_data, 0); + + DeviceToHost(h_data.im_final, d_data.im_final, sizeof(cuComplex) * d_data.nx * d_data.ny); + + { + DeviceToHost(h_data.phdata, d_data.phdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + std::shared_ptr> ifftdata(new std::complex[d_data.Np * d_data.Nfft], delArrPtr); + { + for (long i = 0; i < d_data.Np; i++) { + for (long j = 0; j < d_data.Nfft; j++) { + ifftdata.get()[i * d_data.Nfft + j] = + std::complex( + h_data.phdata[i * d_data.Nfft + j].x, + h_data.phdata[i * d_data.Nfft + j].y); + } + } + } + testOutComplexDoubleArr(QString("echo_ifft_BPBasic.bin"), ifftdata.get(), d_data.Np, d_data.Nfft); + + + + std::shared_ptr> im_finals(new std::complex[d_data.nx * d_data.ny], delArrPtr); + { + for (long i = 0; i < d_data.ny; i++) { + for (long j = 0; j < d_data.nx; j++) { + im_finals.get()[i * d_data.nx + j] = std::complex( + h_data.im_final[i * d_data.nx + j].x, + h_data.im_final[i * d_data.nx + j].y); + } + } + } + testOutComplexDoubleArr(QString("im_finals.bin"), im_finals.get(), d_data.ny, d_data.nx); + } +} + + +void testBpImage() { + + GPUDATA h_data{}; + /** 0. 轨道 **************************************************************************************************/ + QString echoPath = "D:\\Programme\\vs2022\\RasterMergeTest\\LAMPCAE_SCANE-all-scane\\GF3_Simulation.xml"; + std::shared_ptr echoL0ds(new EchoL0Dataset); + echoL0ds->Open(echoPath); + qDebug() << u8"加载回拨文件：\t" << echoPath; + + /** 1. 轨道 **************************************************************************************************/ + qDebug() << u8"1.轨道文件读取中。。。"; + QString inGPSPath = echoL0ds->getGPSPointFilePath(); + long gpspoints = gdalImage(inGPSPath).height; + std::shared_ptr antpos = SatelliteAntPosOperator::readAntPosFile(inGPSPath, gpspoints); + h_data.AntX = (double*)mallocCUDAHost(sizeof(double) * gpspoints); + h_data.AntY = (double*)mallocCUDAHost(sizeof(double) * gpspoints); + h_data.AntZ = (double*)mallocCUDAHost(sizeof(double) * gpspoints); + + for (long i = 0; i < gpspoints; i = i + 1) { + h_data.AntX[i] = antpos.get()[i].Px; + h_data.AntY[i] = antpos.get()[i].Py; + h_data.AntZ[i] = antpos.get()[i].Pz; + } + //gpspoints = gpspoints / 2; + qDebug() << "GPS points Count:\t" << gpspoints; + qDebug() << "PRF 0:\t " << QString("%1,%2,%3").arg(h_data.AntX[0]).arg(h_data.AntY[0]).arg(h_data.AntZ[0]); + qDebug() << "PRF " << gpspoints - 1 << ":\t " << QString("%1,%2,%3") + .arg(h_data.AntX[gpspoints - 1]) + .arg(h_data.AntY[gpspoints - 1]) + .arg(h_data.AntZ[gpspoints - 1]); + /** 2. 成像网格 **************************************************************************************************/ + qDebug() << u8"轨道文件读取结束\n2.成像网格读取。。。"; + QString demxyzPath = u8"D:\\Programme\\vs2022\\RasterMergeTest\\simulationData\\demdataset\\demxyz.bin"; + gdalImage demgridimg(demxyzPath); + long dem_rowCount = demgridimg.height; + long dem_ColCount = demgridimg.width; + std::shared_ptr demX = readDataArr(demgridimg, 0, 0, dem_rowCount, dem_ColCount, 1, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr demY = readDataArr(demgridimg, 0, 0, dem_rowCount, dem_ColCount, 2, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + std::shared_ptr demZ = readDataArr(demgridimg, 0, 0, dem_rowCount, dem_ColCount, 3, GDALREADARRCOPYMETHOD::VARIABLEMETHOD); + + h_data.x_mat = (double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount); + h_data.y_mat = (double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount); + h_data.z_mat = (double*)mallocCUDAHost(sizeof(double) * dem_rowCount * dem_ColCount); + + for (long i = 0; i < dem_rowCount; i++) { + for (long j = 0; j < dem_ColCount; j++) { + h_data.x_mat[i * dem_ColCount + j] = demX.get()[i * dem_ColCount + j]; + h_data.y_mat[i * dem_ColCount + j] = demY.get()[i * dem_ColCount + j]; + h_data.z_mat[i * dem_ColCount + j] = demZ.get()[i * dem_ColCount + j]; } } - dialog->load_double_data(echoAmpArr.get(), d_data.Np, d_data.Nfft, QString("ifft")); - dialog->exec(); + qDebug() << "dem row Count:\t" << dem_rowCount; + qDebug() << "dem col Count:\t" << dem_ColCount; + + qDebug() << u8"成像网格读取结束"; + /** 3. 频率网格 **************************************************************************************************/ + qDebug() << u8"3.处理频率"; + double centerFreq = echoL0ds->getCenterFreq(); + double bandwidth = echoL0ds->getBandwidth(); + size_t freqpoints = echoL0ds->getPlusePoints(); + std::shared_ptr freqlist(getFreqPoints_mallocHost(centerFreq - bandwidth / 2, centerFreq + bandwidth / 2, freqpoints), FreeCUDAHost); + std::shared_ptr minF(new double[gpspoints], delArrPtr); + for (long i = 0; i < gpspoints; i++) { + minF.get()[i] = freqlist.get()[0]; + } + h_data.deltaF = bandwidth / (freqpoints - 1); + qDebug() << "start Freq:\t" << centerFreq - bandwidth / 2; + qDebug() << "end Freq:\t" << centerFreq + bandwidth / 2; + qDebug() << "freq points:\t" << freqpoints; + qDebug() << "delta freq:\t" << freqlist.get()[1] - freqlist.get()[0]; + qDebug() << u8"频率结束"; + /** 4. 初始化回波 **************************************************************************************************/ + qDebug() << u8"4.初始化回波"; + std::shared_ptr> echoData = echoL0ds->getEchoArr(); + size_t echosize = sizeof(cuComplex) * gpspoints * freqpoints; + qDebug() << "echo data size (byte) :\t" << echosize; + h_data.phdata = (cuComplex*)mallocCUDAHost(echosize); + shared_complexPtrToHostCuComplex(echoData.get(), h_data.phdata, gpspoints * freqpoints); + + qDebug() << "Azimuth Points:\t" << gpspoints; + qDebug() << "Range Points:\t" << freqpoints; + qDebug() << u8"初始化回波结束"; + /** 5. 初始化图像 **************************************************************************************************/ + qDebug() << u8"5.初始化图像"; + h_data.im_final = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * dem_rowCount * dem_ColCount); + qDebug() << "Azimuth Points:\t" << gpspoints; + qDebug() << "Range Points:\t" << freqpoints; + qDebug() << u8"初始化图像结束"; + + /** 6. 模型参数初始化 **************************************************************************************************/ + qDebug() << u8"6.模型参数初始化"; + + h_data.nx = dem_ColCount; + h_data.ny = dem_rowCount; + + h_data.Np = gpspoints; + h_data.Freq = freqlist.get(); + h_data.minF = minF.get(); + h_data.Nfft = freqpoints; + h_data.K = h_data.Nfft; + + h_data.R0 = echoL0ds->getRefPhaseRange(); + qDebug() << u8"模型参数结束"; + + + /** 7. 目标 **************************************************************************************************/ + double Tx = -2029086.618142, Ty = 4139594.934504, Tz = 4392846.782027; + double Tslx = -2029086.618142, Tsly = 4139594.934504, Tslz = 4392846.782027; + double p1 = 33.3, p2 = 0, p3 = 0, p4 = 0, p5 = 0, p6 = 0; + + /** 7. 构建回波 **************************************************************************************************/ + GPUDATA d_data; + initGPUData(h_data, d_data); + + + + /** 8. 展示回波 **************************************************************************************************/ + { + DeviceToHost(h_data.phdata, d_data.phdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + cuComplex* h_ifftphdata = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * d_data.Np * d_data.Nfft); + cuComplex* d_ifftphdata = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * d_data.Np * d_data.Nfft); + CUDAIFFT(d_data.phdata, d_ifftphdata, d_data.Np, d_data.Nfft, d_data.Nfft); + FFTShift1D(d_ifftphdata, d_data.Np, d_data.Nfft); + DeviceToHost(h_ifftphdata, d_ifftphdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + std::shared_ptr> ifftdata(new std::complex[d_data.Np * d_data.Nfft], delArrPtr); + { + for (long i = 0; i < d_data.Np; i++) { + for (long j = 0; j < d_data.Nfft; j++) { + ifftdata.get()[i * d_data.Nfft + j] = + std::complex( + h_ifftphdata[i * d_data.Nfft + j].x, + h_ifftphdata[i * d_data.Nfft + j].y); + } + } + } + testOutComplexDoubleArr(QString("echo_ifft.bin"), ifftdata.get(), d_data.Np, d_data.Nfft); + + FreeCUDADevice(d_ifftphdata); + FreeCUDAHost(h_ifftphdata); + + } + /** 9. 成像 **************************************************************************************************/ + + // 计算maxWr（需要先计算deltaF） + double deltaF = h_data.deltaF; // 从输入参数获取 + double maxWr = 299792458.0f / (2.0f * deltaF); + qDebug() << "maxWr :\t" << maxWr; + double* r_vec_host = (double*)mallocCUDAHost(sizeof(double) * h_data.Nfft);// new double[data.Nfft]; + const double step = maxWr / h_data.Nfft; + const double start = -1 * h_data.Nfft / 2.0f * step; + printf("nfft=%d\n", h_data.Nfft); + + for (int i = 0; i < h_data.Nfft; ++i) { + r_vec_host[i] = start + i * step; + } + + h_data.r_vec = r_vec_host; + d_data.r_vec = h_data.r_vec; + + qDebug() << "r_vec [0]:\t" << h_data.r_vec[0]; + qDebug() << "r_vec step:\t" << step; + + bpBasic0CUDA(d_data, 0); + + DeviceToHost(h_data.im_final, d_data.im_final, sizeof(cuComplex) * d_data.nx * d_data.ny); + + { + DeviceToHost(h_data.phdata, d_data.phdata, sizeof(cuComplex) * d_data.Np * d_data.Nfft); + std::shared_ptr> ifftdata(new std::complex[d_data.Np * d_data.Nfft], delArrPtr); + { + for (long i = 0; i < d_data.Np; i++) { + for (long j = 0; j < d_data.Nfft; j++) { + ifftdata.get()[i * d_data.Nfft + j] = + std::complex( + h_data.phdata[i * d_data.Nfft + j].x, + h_data.phdata[i * d_data.Nfft + j].y); + } + } + } + testOutComplexDoubleArr(QString("echo_ifft_BPBasic.bin"), ifftdata.get(), d_data.Np, d_data.Nfft); + std::shared_ptr> im_finals(new std::complex[d_data.nx * d_data.ny], delArrPtr); + { + for (long i = 0; i < d_data.ny; i++) { + for (long j = 0; j < d_data.nx; j++) { + im_finals.get()[i * d_data.nx + j] = std::complex( + h_data.im_final[i * d_data.nx + j].x, + h_data.im_final[i * d_data.nx + j].y); + } + } + } + testOutComplexDoubleArr(QString("im_finals.bin"), im_finals.get(), d_data.ny, d_data.nx); + } + + + +} + + +int main(int argc, char* argv[]) { + + QApplication a(argc, argv); + + testBpImage(); return 0; }