修改了 时域插值BP算法

pull/5/head
陈增辉 2025-02-25 13:18:19 +08:00
parent 3dd749c414
commit e68ccf57eb
21 changed files with 432 additions and 364 deletions

View File

@ -5,7 +5,7 @@
int main()
{
std::cout << "Hello World!\n";
std::cout << "Hello World!\n";
}
// 运行程序: Ctrl + F5 或调试 >“开始执行(不调试)”菜单

View File

@ -316,5 +316,11 @@ inline long getBlockRows(long sizeMB, long cols,long sizeMeta,long maxRows) {
}
inline long nextpow2(long n) {
long en = ceil(log2(n));
return pow(2,en);
}
#endif

View File

@ -360,14 +360,14 @@ double EchoL0Dataset::getRefPhaseRange()
// 打印信息的实现
void EchoL0Dataset::printInfo() {
std::cout << "Simulation Task Name: " << this->simulationTaskName.toStdString() << std::endl;
std::cout << "Pulse Count: " << this->PluseCount << std::endl;
std::cout << "Pulse Points: " << this->PlusePoints << std::endl;
std::cout << "Near Range: " << this->NearRange << std::endl;
std::cout << "Far Range: " << this->FarRange << std::endl;
std::cout << "Center Frequency: " << this->centerFreq << std::endl;
std::cout << "Sampling Frequency: " << this->Fs << std::endl;
std::cout << "Band width: " << this->bandwidth << std::endl;
qDebug() << "Simulation Task Name: " << this->simulationTaskName ;
qDebug() << "Pulse Count: " << this->PluseCount ;
qDebug() << "Pulse Points: " << this->PlusePoints;
qDebug() << "Near Range: " << this->NearRange ;
qDebug() << "Far Range: " << this->FarRange;
qDebug() << "Center Frequency: " << this->centerFreq ;
qDebug() << "Sampling Frequency: " << this->Fs ;
qDebug() << "Band width: " << this->bandwidth ;
}
// xml文件读写

View File

@ -385,7 +385,7 @@ CartesianCoordinates sphericalToCartesian(const SphericalCoordinates& spherica
double getlocalIncAngle(Landpoint& satepoint, Landpoint& landpoint, Landpoint& slopeVector) {
Landpoint Rsc = satepoint - landpoint; // AB=B-A
//double R = getlength(Rsc);
//std::cout << R << endl;
//qDebug() << R << endl;
double angle = getAngle(Rsc, slopeVector);
if (angle >= 180) {
return 360 - angle;
@ -399,7 +399,7 @@ double getlocalIncAngle(Landpoint& satepoint, Landpoint& landpoint, Landpoint& s
double getlocalIncAngle(Vector3D& satepoint, Vector3D& landpoint, Vector3D& slopeVector){
Vector3D Rsc = satepoint - landpoint; // AB=B-A
//double R = getlength(Rsc);
//std::cout << R << endl;
//qDebug() << R << endl;
double angle = getAngle(Rsc, slopeVector);
if (angle >= 180) {
return 360 - angle;

View File

@ -1756,13 +1756,13 @@ gdalImage CreategdalImage(const QString& img_path, int height, int width, int ba
OGRSpatialReference oSRS;
if (oSRS.importFromEPSG(epsgCode) != OGRERR_NONE) {
std::cerr << "Failed to import EPSG code " << epsgCode << std::endl;
qDebug() << "Failed to import EPSG code " << epsgCode ;
throw "Failed to import EPSG code ";
exit(1);
}
char* pszWKT = NULL;
oSRS.exportToWkt(&pszWKT);
std::cout << "WKT of EPSG:"<< epsgCode <<" :\n" << pszWKT << std::endl;
qDebug() << "WKT of EPSG:"<< epsgCode <<" :\n" << pszWKT ;
poDstDS->SetProjection(pszWKT);
double gt_ptr[6] = { 0 };
for (int i = 0; i < gt.rows(); i++) {
@ -2438,7 +2438,7 @@ ErrorCode MergeRasterInGeoCoding(QVector<gdalImage> imgdslist, gdalImage resulti
omp_set_lock(&lock);
processNumber = processNumber + blocklines;
std::cout << "\rprocess bar:\t" << processNumber * 100.0 / resultimg.height << " % " << "\t\t\t";
qDebug() << "\rprocess bar:\t" << processNumber * 100.0 / resultimg.height << " % " << "\t\t\t";
if (nullptr != dia) {
dia->showProcess(processNumber * 1.0 / resultimg.height, u8"合并图像");
}
@ -2449,7 +2449,7 @@ ErrorCode MergeRasterInGeoCoding(QVector<gdalImage> imgdslist, gdalImage resulti
omp_unset_lock(&lock);
}
omp_destroy_lock(&lock);
std::cout << std::endl;
progressDialog.setWindowTitle(u8"影像掩膜");
progressDialog.setLabelText(u8"影像掩膜");
for (starti = 0; starti < resultimg.height; starti = starti + resultline) {
@ -2846,7 +2846,7 @@ long GetEPSGFromRasterFile(QString filepath)
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES"); // 注册GDAL驱动
// std::cout<<filepath.toLocal8Bit().constData()<<std::endl;
// qDebug()<<filepath.toLocal8Bit().constData()<<std::endl;
// 打开影像文件
GDALDataset* poDataset;
poDataset = (GDALDataset*)GDALOpen(filepath.toUtf8().data(), GA_ReadOnly);
@ -2868,7 +2868,7 @@ long GetEPSGFromRasterFile(QString filepath)
return -1;
}
std::cout << pszProjection << std::endl;
qDebug() << pszProjection ;
const char* epscodestr = oSRS.GetAuthorityCode(nullptr);
if (NULL == epscodestr || nullptr == epscodestr) {
@ -2902,7 +2902,7 @@ std::shared_ptr<PointRaster> GetCenterPointInRaster(QString filepath)
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
// std::cout<<filepath.toLocal8Bit().constData()<<std::endl;
// qDebug()<<filepath.toLocal8Bit().constData()<<std::endl;
GDALDataset* poDataset = (GDALDataset*)GDALOpen(filepath.toUtf8().data(), GA_ReadOnly);
if (nullptr == poDataset || NULL == poDataset) {
qDebug() << "Could not open dataset";
@ -3074,7 +3074,7 @@ void resampleRaster(const char* inputRaster, const char* outputRaster, double ta
GDALClose(poDataset);
GDALClose(poOutDataset);
std::cout << "Resampling completed." << std::endl;
qDebug() << "Resampling completed." ;
}
void transformRaster(const char* inputFile, const char* outputFile, int sourceEPSG, int targetEPSG) {
@ -3157,7 +3157,7 @@ ErrorCode transformCoordinate(double x, double y, int sourceEPSG, int targetEPSG
// 创建坐标变换对象
OGRCoordinateTransformation* transform = OGRCreateCoordinateTransformation(&sourceSRS, &targetSRS);
if (transform == nullptr) {
std::cerr << "Failed to create coordinate transformation." << std::endl;
qDebug() << "Failed to create coordinate transformation." ;
return ErrorCode::FAIL;
}
@ -3165,11 +3165,11 @@ ErrorCode transformCoordinate(double x, double y, int sourceEPSG, int targetEPSG
double transformedX = x;
double transformedY = y;
if (transform->Transform(1, &transformedX, &transformedY)) {
std::cout << "Original Coordinates: (" << x << ", " << y << ")" << std::endl;
std::cout << "Transformed Coordinates (EPSG:" << targetEPSG << "): (" << transformedX << ", " << transformedY << ")" << std::endl;
qDebug() << "Original Coordinates: (" << x << ", " << y << ")" ;
qDebug() << "Transformed Coordinates (EPSG:" << targetEPSG << "): (" << transformedX << ", " << transformedY << ")" ;
}
else {
std::cerr << "Coordinate transformation failed." << std::endl;
qDebug() << "Coordinate transformation failed." ;
}
// 清理
@ -3246,7 +3246,7 @@ void cropRasterByLatLon(const char* inputFile, const char* outputFile,double min
delete[] pData;
}
std::cout << "Raster cropped and saved to: " << outputFile << std::endl;
qDebug() << "Raster cropped and saved to: " << outputFile ;
// 清理
GDALClose(poDataset);
@ -3370,9 +3370,9 @@ void testOutAmpArr(QString filename, float* amp, long rowcount, long colcount)
}
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;
qDebug() << filename.toLocal8Bit().constData() ;
qDebug() << "max:\t" << h_amp_img.maxCoeff() ;
qDebug() << "min:\t" << h_amp_img.minCoeff() ;
}
void testOutAmpArr(QString filename, double* amp, long rowcount, long colcount)
@ -3388,9 +3388,9 @@ void testOutAmpArr(QString filename, double* amp, long rowcount, long colcount)
}
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;
qDebug() << filename.toLocal8Bit().constData() ;
qDebug() << "max:\t" << h_amp_img.maxCoeff() ;
qDebug() << "min:\t" << h_amp_img.minCoeff() ;
}
@ -3405,9 +3405,9 @@ void testOutClsArr(QString filename, long* amp, long rowcount, long colcount) {
}
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;
qDebug() << filename.toLocal8Bit().constData() ;
qDebug() << "max:\t" << h_amp_img.maxCoeff() ;
qDebug() << "min:\t" << h_amp_img.minCoeff() ;
}

View File

@ -131,7 +131,7 @@ namespace RasterToolBase {
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES"); // 注册GDAL驱动
// std::cout<<filepath.toLocal8Bit().constData()<<std::endl;
// qDebug()<<filepath.toLocal8Bit().constData()<<std::endl;
// 打开影像文件
GDALDataset* poDataset;
poDataset = (GDALDataset*)GDALOpen(filepath.toUtf8().data(), GA_ReadOnly);
@ -176,7 +176,7 @@ namespace RasterToolBase {
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
// std::cout<<filepath.toLocal8Bit().constData()<<std::endl;
// qDebug()<<filepath.toLocal8Bit().constData()<<std::endl;
GDALDataset* poDataset = (GDALDataset*)GDALOpen(filepath.toUtf8().data(), GA_ReadOnly);
if(nullptr==poDataset||NULL==poDataset) {
qDebug() << "Could not open dataset";

View File

@ -188,6 +188,7 @@
<EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalDependencies>cufft.lib;%(AdditionalDependencies)</AdditionalDependencies>
</Link>
<CudaCompile>
<GenerateRelocatableDeviceCode>true</GenerateRelocatableDeviceCode>

View File

@ -6,6 +6,9 @@
#include <complex>
#include <device_launch_parameters.h>
#include <cuda_runtime.h>
#include <cufft.h>
#include <cufftw.h>
#include <cufftXt.h>
#include <cublas_v2.h>
#include <cuComplex.h>
#include <chrono>
@ -17,13 +20,21 @@
#define BLOCK_DIM 1024
#define REDUCE_SCALE 4
// CUDA核函数用于缩放数据
__global__ void scaleKernel(cuComplex* data, int size, float scale) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
data[idx].x *= scale;
data[idx].y *= scale;
}
}
// ´òÓ¡GPU²ÎÊý
void printDeviceInfo(int deviceId) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, deviceId);
std::cout << "Device " << deviceId << ": " << deviceProp.name << std::endl;
std::cout << " Compute Capability: " << deviceProp.major << "." << deviceProp.minor << std::endl;
std::cout << " Total Global Memory: " << deviceProp.totalGlobalMem / (1024 * 1024) << " MB" << std::endl;
@ -35,12 +46,12 @@ void printDeviceInfo(int deviceId) {
<< deviceProp.maxThreadsDim[1] << ", " << deviceProp.maxThreadsDim[2] << ")" << std::endl;
std::cout << " Max Grid Size: (" << deviceProp.maxGridSize[0] << ", "
<< deviceProp.maxGridSize[1] << ", " << deviceProp.maxGridSize[2] << ")" << std::endl;
std::cout << " Multiprocessor Count: " << deviceProp.multiProcessorCount << std::endl;
std::cout << " Clock Rate: " << deviceProp.clockRate / 1000 << " MHz" << std::endl;
std::cout << " Memory Clock Rate: " << deviceProp.memoryClockRate / 1000 << " MHz" << std::endl;
std::cout << " Memory Bus Width: " << deviceProp.memoryBusWidth << " bits" << std::endl;
std::cout << " L2 Cache Size: " << deviceProp.l2CacheSize / 1024 << " KB" << std::endl;
std::cout << " Max Texture Dimensions: (" << deviceProp.maxTexture1D << ", "
std::cout << " Multiprocessor Count: " << deviceProp.multiProcessorCount << std::endl;
std::cout << " Clock Rate: " << deviceProp.clockRate / 1000 << " MHz" << std::endl;
std::cout << " Memory Clock Rate: " << deviceProp.memoryClockRate / 1000 << " MHz" << std::endl;
std::cout << " Memory Bus Width: " << deviceProp.memoryBusWidth << " bits" << std::endl;
std::cout << " L2 Cache Size: " << deviceProp.l2CacheSize / 1024 << " KB" << std::endl;
std::cout << " Max Texture Dimensions: (" << deviceProp.maxTexture1D << ", "
<< deviceProp.maxTexture2D[0] << "x" << deviceProp.maxTexture2D[1] << ", "
<< deviceProp.maxTexture3D[0] << "x" << deviceProp.maxTexture3D[1] << "x" << deviceProp.maxTexture3D[2] << ")" << std::endl;
std::cout << " Unified Addressing: " << (deviceProp.unifiedAddressing ? "Yes" : "No") << std::endl;
@ -495,6 +506,68 @@ void PrintLasterError(const char* s)
extern "C" void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long InRowCount, long InColCount, long outColCount) {
cufftHandle plan;
cufftResult result;
// 创建批量IFFT计划
int rank = 1;
int n[] = { InColCount }; // 每个IFFT处理freqcount点
int inembed[] = { InColCount };
int onembed[] = { InColCount };
int istride = 1;
int ostride = 1;
int idist = InColCount; // 输入批次间距
int odist = InColCount; // 输出批次间距
int batch = InRowCount; // 批处理数量
result = cufftPlanMany(&plan, rank, n,
inembed, istride, idist,
onembed, ostride, odist,
CUFFT_C2C, batch);
if (result != CUFFT_SUCCESS) {
PrintLasterError("CUDAIFFT");
return;
}
// 执行IFFT
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;
}
// 等待IFFT完成并缩放数据
cudaDeviceSynchronize();
cufftDestroy(plan);
}
extern "C" void CUDAIFFTScale(cuComplex* inArr, cuComplex* outArr, long InRowCount, long InColCount, long outColCount)
{
CUDAIFFT(inArr, outArr, InRowCount, InColCount, outColCount);
float scale = 1.0f / InColCount;
int totalElements = InRowCount * InColCount;
dim3 block(256);
dim3 grid((totalElements + block.x - 1) / block.x);
scaleKernel << <grid, block >> > (outArr, totalElements, scale);
cudaDeviceSynchronize();
return ;
}
#endif

View File

@ -104,5 +104,10 @@ extern "C" GPUBASELIBAPI long NextBlockPad(long num, long blocksize);
extern "C" GPUBASELIBAPI void PrintLasterError(const char* s);
extern "C" GPUBASELIBAPI void CUDAIFFTScale(cuComplex* inArr, cuComplex* outArr,long InRowCount,long InColCount,long outColCount);
extern "C" GPUBASELIBAPI void CUDAIFFT(cuComplex* inArr, cuComplex* outArr, long InRowCount, long InColCount, long outColCount);
#endif
#endif

View File

@ -493,7 +493,7 @@ ErrorCode GF3RDCreateLookTable(QString inxmlPath, QString indemPath, QString out
}
else {}
processNumber = processNumber + blockrows;
std::cout << "\rprocess bar:\t" << processNumber * 100.0 / rowcount << " % " << "\t\t\t";
qDebug() << "\rprocess bar:\t" << processNumber * 100.0 / rowcount << " % " << "\t\t\t";
if (progressDialog.maximum() <= processNumber) {
processNumber = progressDialog.maximum() - 1;
}

View File

@ -233,7 +233,7 @@ ErrorCode GF3PolyfitSatelliteOribtModel::getAntnnaDirection(SatelliteOribtNode&
double nexttime = node.time + 1e-6;
SatelliteOribtNode node1 = this->getSatelliteOribtNode(nexttime, flag);
//std::cout << "getAntnnaDirection corrdination " << std::endl;
//qDebug() << "getAntnnaDirection corrdination " << std::endl;
double Vx = (node1.Px - node.Px);
double Vy = (node1.Py - node.Py);
@ -252,13 +252,13 @@ ErrorCode GF3PolyfitSatelliteOribtModel::getAntnnaDirection(SatelliteOribtNode&
Eigen::Vector3d axisX0 = { Vx,Vy,Vz }; // x 轴 --飞行方向
Eigen::Vector3d axisY0 = axisZ0.cross(axisX0); // y 轴 --右手定则 -- 初始坐标系
//std::cout << "axis_X0=[ " << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//std::cout << "axis_Y0=[ " << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//std::cout << "axis_Z0=[ " << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
//qDebug() << "axis_X0=[ " << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//qDebug() << "axis_Y0=[ " << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//qDebug() << "axis_Z0=[ " << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
double rotateAngle = this->RightLook ? -this->beamAngle : this->beamAngle; // 旋转角度 左逆时针theta , 右(顺时针): -theta
//std::cout << "rotateAngle=" << rotateAngle << std::endl;
//std::cout << "Look side:\t" << (this->RightLook ? "right" : "left") << std::endl;
//qDebug() << "rotateAngle=" << rotateAngle << std::endl;
//qDebug() << "Look side:\t" << (this->RightLook ? "right" : "left") << std::endl;
// 1.2. 根据波位角确定卫星绕X轴-飞行轴
Eigen::Matrix3d rotateMatrixBeam = rotationMatrix(axisX0, rotateAngle * d2r); // 旋转矩阵
axisZ0 = rotateMatrixBeam * axisZ0; // 旋转矩阵
@ -291,10 +291,10 @@ ErrorCode GF3PolyfitSatelliteOribtModel::getAntnnaDirection(SatelliteOribtNode&
node.AntZaxisZ = axisZ0[2];
//std::cout << "axis_X=[" << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//std::cout << "axis_Y=[" << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//std::cout << "axis_Z=[" << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
//std::cout << "------------------------------------" << std::endl;
//qDebug() << "axis_X=[" << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//qDebug() << "axis_Y=[" << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//qDebug() << "axis_Z=[" << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
//qDebug() << "------------------------------------" << std::endl;
return ErrorCode::SUCCESS;
}

View File

@ -127,7 +127,15 @@ PSTNAlgorithm::PSTNAlgorithm(QString infile_path)
OrbitPoly orbit(polynum, polySatellitePara, SatelliteModelStartTime);
this->orbit = orbit;
qDebug() << "sate polynum\t" << polynum ;
std::cout << "sate polySatellitePara\n" << polySatellitePara ;
qDebug() << "sate polySatellitePara\n" ;
for (long i = 0; i < polynum; i++) {
qDebug() << polySatellitePara(i, 0) << "\t"
<< polySatellitePara(i, 1) << "\t"
<< polySatellitePara(i, 2) << "\t"
<< polySatellitePara(i, 3) << "\t"
<< polySatellitePara(i, 4) << "\t"
<< polySatellitePara(i, 5) << "\t";
}
qDebug() << "sate SatelliteModelStartTime\t" << SatelliteModelStartTime ;
@ -142,7 +150,7 @@ PSTNAlgorithm::PSTNAlgorithm(QString infile_path)
getline(infile, buf); dd = stod(buf);
this->UTC = Eigen::Matrix<double, 1, 3>{ yy,mm,dd };
std::cout << "UTC\t" << this->UTC << std::endl;
qDebug() << "\nWGS84 to J2000 Params:\t" ;
std::cout << "\nWGS84 to J2000 Params:\t" ;
getline(infile, buf); this->Xp = stod(buf); qDebug() << "Xp\t" << this->Xp ;
getline(infile, buf); this->Yp = stod(buf); qDebug() << "Yp\t" << this->Yp ;
getline(infile, buf); this->Dut1 = stod(buf); qDebug() << "dut1\t" << this->Dut1 ;

View File

@ -437,7 +437,7 @@ QMap<QString, QString> LAMPScatterS1BDataset::ReadPolAttribution(int ncid, int p
break;
}
default:
std::cout << "Attribute Name: " << att_name << ", Type: Unknown" << std::endl;
qDebug() << "Attribute Name: " << att_name << ", Type: Unknown";
break;
}
}

View File

@ -15,94 +15,109 @@
#ifdef __CUDANVCC___
// 定义参数
__device__ cuComplex cuCexpf(cuComplex d)
{
float factor = exp(d.x);
return make_cuComplex(factor * cos(d.y), factor * sin(d.y));
}
__global__ void CUDA_TBPImage(
float* antPx, float* antPy, float* antPz,
float* imgx, float* imgy, float* imgz,
float* RArr,
long* Cids,
cuComplex* echoArr,
cuComplex* imgArr,
cuComplex* imgEchoArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount
__global__ void kernel_pixelTimeBP(
double* antPx, double* antPy, double* antPz,
double* imgx, double* imgy, double* imgz,
cuComplex* TimeEchoArr, long prfcount, long pointCount,
cuComplex* imgArr, long imH, long imW,
double startLamda, double Rnear, double dx,double RefRange
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
//printf("\nidx:\t %d %d %d\n", idx, linecount, plusepoint);
if (idx < rowcount * colcount) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
long pixelcount = imH * imW;
if (idx < pixelcount) {
double Tx = imgx[idx], Ty = imgy[idx], Tz = imgz[idx], PR = 0;
double Px = 0, Py = 0, Pz = 0;
double Rid = -1;
long maxPointIdx = pointCount - 1;
long pid = 0;
long pid0 = 0;
long pid1 = 0;
double phi = 0;
cuComplex s0=make_cuComplex(0,0), s1=make_cuComplex(0,0);
cuComplex s = make_cuComplex(0, 0);
cuComplex phiCorr = make_cuComplex(0, 0);
for (long spid = 0; spid < prfcount; spid = spid + 8) {
float px = antPx[prfid];
float py = antPy[prfid];
float pz = antPz[prfid];
#pragma unroll
for (long sid = 0; sid < 8; sid++)
{
pid = spid + sid;
if (pid < prfcount) {}
else {
continue;
}
float tx = imgx[idx];
float ty = imgy[idx];
float tz = imgz[idx];
Px = antPx[pid];
Py = antPy[pid];
Pz = antPz[pid];
float R = sqrtf((px-tx) * (px - tx) + (py-ty) * (py-ty) + (pz-tz) * (pz-tz));
float Cidf= 2 * (R - Rnear) / LIGHTSPEED * fs;
long Cid = floorf(Cidf);
RArr[idx] = R;
Cids[idx] = Cid;
if(Cid <0|| Cid >= freqcount){}
else {
float factorj = freq * 4 * PI / LIGHTSPEED;
cuComplex Rfactorj = make_cuComplex(0, factorj * R);
cuComplex Rphi =cuCexpf(Rfactorj);// 校正项
cuComplex echotemp = echoArr[prfid * freqcount + Cid];
imgEchoArr[idx] = echotemp;
imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echotemp, Rphi));// 矫正
//printf("R=%f;Rid=%d;factorj=%f;Rfactorj=complex(%f,%f);Rphi=complex(%f,%f);\n", R, Rid, factorj,
// Rfactorj.x, Rfactorj.y,
// Rphi.x, Rphi.y);
PR = sqrt((Px - Tx) * (Px - Tx) + (Py - Ty) * (Py - Ty) + (Pz - Tz) * (Pz - Tz));
Rid = (PR - Rnear) / dx; // 行数
if (Rid<0 || Rid>maxPointIdx) {
continue;
}
else {}
pid0 = floor(Rid);
pid1 = ceil(Rid);
if (pid1<0 || pid0<0 || pid0>maxPointIdx || pid1>maxPointIdx) {
continue;
}
else {}
// 线性插值
s0 = TimeEchoArr[pid0];
s1 = TimeEchoArr[pid1];
s.x = s0.x * (Rid - pid0) + s1.x * (pid1 - Rid);
s.y = s0.y * (Rid - pid0) + s1.y * (pid1 - Rid);
// 相位校正
phi = 4 * LIGHTSPEED/startLamda* (PR - RefRange) ; // 4PI/lamda * R
// exp(ix)=cos(x)+isin(x)
phiCorr.x = cos(phi);
phiCorr.y = sin(phi);
s = cuCmulf(s, phiCorr); // 校正后
imgArr[idx] = cuCaddf(imgArr[idx], s);
}
}
}
}
extern "C" void CUDATBPImage(float* antPx, float* antPy, float* antPz,
float* imgx, float* imgy, float* imgz,
float* R,
long* Cids,
cuComplex* echoArr,
cuComplex* imgArr,
cuComplex* imgEchoArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount)
{
int blockSize = 256; // 每个块的线程数
int numBlocks = (rowcount * colcount + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
//printf("\nCUDA_RFPC_SiglePRF blockSize:%d ,numBlock:%d\n",blockSize,numBlocks);
// 调用 CUDA 核函数 CUDA_RFPC_Kernel
CUDA_TBPImage << <numBlocks, blockSize >> > (
extern "C" void TimeBPImage(double* antPx, double* antPy, double* antPz,
double* imgx, double* imgy, double* imgz,
cuComplex* TimeEchoArr, long prfcount, long pointCount,
cuComplex* imgArr, long imH, long imW,
double startLamda, double Rnear, double dx, double RefRange
)
{
long pixelcount = imH * imW;
int grid_size = (pixelcount + BLOCK_SIZE - 1) / BLOCK_SIZE;
kernel_pixelTimeBP << <grid_size, BLOCK_SIZE >> > (
antPx, antPy, antPz,
imgx, imgy, imgz,
R, Cids,
echoArr, imgArr, imgEchoArr,
freq, fs, Rnear, Rfar,
rowcount, colcount,
prfid, freqcount
TimeEchoArr, prfcount, pointCount,
imgArr, imH, imW,
startLamda, Rnear, dx, RefRange
);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDATBPImage CUDA Error: %s\n", cudaGetErrorString(err));
exit(2);
}
#endif // __CUDADEBUG__
PrintLasterError("TimeBPImage");
cudaDeviceSynchronize();
}

View File

@ -11,30 +11,15 @@
extern __global__ void CUDA_TBPImage(
float* antPx, float* antPy, float* antPz,
float* imgx, float* imgy, float* imgz, float* R,
cuComplex* echoArr, cuComplex* imgArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount
);
extern "C" void CUDATBPImage(
float* antPx,
float* antPy,
float* antPz,
float* imgx,
float* imgy,
float* imgz,
float* R,
long* Cids,
cuComplex* echoArr,
cuComplex* imgArr,
cuComplex* imgEchoArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount
extern "C" void TimeBPImage(
double* antPx, double* antPy, double* antPz,
double* imgx, double* imgy, double* imgz,
cuComplex* TimeEchoArr, long prfcount, long pointCount,
cuComplex* imgArr, long imH, long imW,
double startLamda, double Rnear, double dx, double RefRange
);
#endif

View File

@ -177,15 +177,15 @@ void QImageSARRFPC::onBtnaccept()
// 打印参数
// 打印解析的参数
std::cout << "GPS XML Path: " << GPSXmlPath.toStdString() << "\n"
<< "Task XML Path: " << TaskXmlPath.toStdString() << "\n"
<< "DEM TIFF Path: " << demTiffPath.toStdString() << "\n"
<< "Land Cover Path: " << landConverPath.toStdString() << "\n"
qDebug() << "GPS XML Path: " << GPSXmlPath << "\n"
<< "Task XML Path: " << TaskXmlPath << "\n"
<< "DEM TIFF Path: " << demTiffPath<< "\n"
<< "Land Cover Path: " << landConverPath << "\n"
<< "Trans AntPattern Path: " << TransAntPatternFilePath.toStdString() << "\n"
<< "Reception AntPattern Path: " << ReceiveAntPatternFilePath.toStdString() << "\n"
<< "Output Path: " << OutEchoPath.toStdString() << "\n"
<< "Simulation Task Name: " << simulationtaskName.toStdString() << "\n";
<< "Trans AntPattern Path: " << TransAntPatternFilePath << "\n"
<< "Reception AntPattern Path: " << ReceiveAntPatternFilePath << "\n"
<< "Output Path: " << OutEchoPath << "\n"
<< "Simulation Task Name: " << simulationtaskName << "\n";
long cpucore_num = std::thread::hardware_concurrency();

View File

@ -330,8 +330,13 @@ ErrorCode RFPCProcessCls::InitParams()
double imgStart_end = this->TaskSetting->getSARImageEndTime() - this->TaskSetting->getSARImageStartTime();
this->PluseCount = ceil(imgStart_end * this->TaskSetting->getPRF());
double rangeTimeSample = (this->TaskSetting->getFarRange() - this->TaskSetting->getNearRange()) * 2.0 / LIGHTSPEED;
this->PlusePoint = ceil(rangeTimeSample * this->TaskSetting->getFs());
// 远斜距
//double rangeTimeSample = * 2.0 / LIGHTSPEED;
double drange = LIGHTSPEED / 2.0 / this->TaskSetting->getBandWidth();
this->PlusePoint = ceil((this->TaskSetting->getFarRange() - this->TaskSetting->getNearRange()) / LIGHTSPEED * 2 * this->TaskSetting->getBandWidth());
this->TaskSetting->setFarRange(this->TaskSetting->getNearRange() + (this->PlusePoint-1) * drange);
//ceil(rangeTimeSample * this->TaskSetting->getFs());
// 初始化回波存放位置
qDebug() << "--------------Echo Data Setting ---------------------------------------";
@ -696,9 +701,9 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
}
}
std::cout << "class id recoding" << std::endl;
qDebug() << "class id recoding" ;
for (long id : clamap.keys()) {
std::cout << id << " -> " << clamap[id] << std::endl;
qDebug() << id << " -> " << clamap[id] ;
}
}
@ -718,17 +723,17 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
}
// 打印日志
std::cout << "sigma params:" << std::endl;
std::cout << "classid:\tp1\tp2\tp3\tp4\tp5\tp6" << std::endl;
qDebug() << "sigma params:" ;
qDebug() << "classid:\tp1\tp2\tp3\tp4\tp5\tp6" ;
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;
qDebug() << ii << ":\t" << h_clsSigmaParam[ii].p1;
qDebug() << "\t" << h_clsSigmaParam[ii].p2;
qDebug() << "\t" << h_clsSigmaParam[ii].p3;
qDebug() << "\t" << h_clsSigmaParam[ii].p4;
qDebug() << "\t" << h_clsSigmaParam[ii].p5;
qDebug() << "\t" << h_clsSigmaParam[ii].p6 ;
}
std::cout << "";
qDebug() << "";
}
HostToDevice(h_clsSigmaParam, d_clsSigmaParam, sizeof(CUDASigmaParam) * clamapid);

View File

@ -501,7 +501,7 @@ std::vector<RadiationPatternGainPoint> ReadGainFile(QString antPatternFilePath)
return dataPoints;
}
std::string filepath = antPatternFilePath.toLocal8Bit().constData();
std::cout << "ant file path:\t" << filepath << std::endl;
qDebug() << "ant file path:\t" << QString::fromStdString(filepath);
std::ifstream file(filepath);
@ -520,18 +520,17 @@ std::vector<RadiationPatternGainPoint> ReadGainFile(QString antPatternFilePath)
headers.push_back(header);
}
std::cout << "Headers:";
qDebug() << "Headers:";
for (const auto& h : headers) {
std::cout << " " << h;
qDebug() << " " << QString::fromStdString(h);
}
std::cout << std::endl;
if (headers.size() < 3) {
file.close();
return dataPoints; // ·µ»Ø¿ÕÏòÁ¿
}
std::cout << "Parse ant radiation pattern contains " << std::endl;
qDebug() << "Parse ant radiation pattern contains " ;
long theta_id = -1;
long phi_id = -1;
@ -569,7 +568,7 @@ std::vector<RadiationPatternGainPoint> ReadGainFile(QString antPatternFilePath)
lines.push_back(line);
}
std::cout << "Read file over" << std::endl;
qDebug() << "Read file over" ;
file.close();
@ -776,7 +775,7 @@ ErrorCode AbstractRadiationPattern::RecontructGainMatrix(double threshold)
Eigen::VectorXd phipoints = linspace(minphi, maxphi, phinum);
Eigen::MatrixXd gaintemp = Eigen::MatrixXd::Zero(thetanum, phinum);
std::cout << "gain init" << std::endl;
qDebug() << "gain init" ;
for (long i = 0; i < thetanum; i++) {
for (long j = 0; j < phinum; j++) {
gaintemp(i, j) = this->getGain(thetapoints[i], phipoints[j]);

View File

@ -232,7 +232,7 @@ ErrorCode PolyfitSatelliteOribtModel::getAntnnaDirection(SatelliteOribtNode& nod
double nexttime = node.time + 1e-6;
SatelliteOribtNode node1 = this->getSatelliteOribtNode(nexttime, flag);
//std::cout << "getAntnnaDirection corrdination " << std::endl;
//qDebug() << "getAntnnaDirection corrdination " << std::endl;
double Vx = (node1.Px - node.Px);
double Vy = (node1.Py - node.Py);
@ -251,13 +251,13 @@ ErrorCode PolyfitSatelliteOribtModel::getAntnnaDirection(SatelliteOribtNode& nod
Eigen::Vector3d axisX0 = { Vx,Vy,Vz }; // x 轴 --飞行方向
Eigen::Vector3d axisY0 = axisZ0.cross(axisX0); // y 轴 --右手定则 -- 初始坐标系
//std::cout << "axis_X0=[ " << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//std::cout << "axis_Y0=[ " << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//std::cout << "axis_Z0=[ " << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
//qDebug() << "axis_X0=[ " << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//qDebug() << "axis_Y0=[ " << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//qDebug() << "axis_Z0=[ " << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
double rotateAngle = this->RightLook ? -this->beamAngle : this->beamAngle; // 旋转角度 左逆时针theta , 右(顺时针): -theta
//std::cout << "rotateAngle=" << rotateAngle << std::endl;
//std::cout << "Look side:\t" << (this->RightLook ? "right" : "left") << std::endl;
//qDebug() << "rotateAngle=" << rotateAngle << std::endl;
//qDebug() << "Look side:\t" << (this->RightLook ? "right" : "left") << std::endl;
// 1.2. 根据波位角确定卫星绕X轴-飞行轴
Eigen::Matrix3d rotateMatrixBeam = rotationMatrix(axisX0, rotateAngle*d2r); // 旋转矩阵
axisZ0=rotateMatrixBeam*axisZ0; // 旋转矩阵
@ -290,10 +290,10 @@ ErrorCode PolyfitSatelliteOribtModel::getAntnnaDirection(SatelliteOribtNode& nod
node.AntZaxisZ =axisZ0[2];
//std::cout << "axis_X=[" << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//std::cout << "axis_Y=[" << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//std::cout << "axis_Z=[" << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
//std::cout << "------------------------------------" << std::endl;
//qDebug() << "axis_X=[" << axisX0.x() << "," << axisX0.y() << "," << axisX0.z() << "]" << std::endl;
//qDebug() << "axis_Y=[" << axisY0.x() << "," << axisY0.y() << "," << axisY0.z() << "]" << std::endl;
//qDebug() << "axis_Z=[" << axisZ0.x() << "," << axisZ0.y() << "," << axisZ0.z() << "]" << std::endl;
//qDebug() << "------------------------------------" << std::endl;
return ErrorCode::SUCCESS;
}

View File

@ -29,7 +29,7 @@ void CreatePixelXYZ(std::shared_ptr<EchoL0Dataset> echoL0ds, QString outPixelXYZ
double Rref = echoL0ds->getRefPhaseRange();
double centerInc = echoL0ds->getCenterAngle()*d2r;
long echocol = 1073741824 / 8 / 4 / prfcount*8;
std::cout << "echocol:\t " << echocol << std::endl;
qDebug() << "echocol:\t " << echocol ;
echocol = echocol < 3000 ? 3000 : echocol;
long startcolidx = 0;
for (startcolidx = 0; startcolidx < freqcount; startcolidx = startcolidx + echocol) {
@ -38,7 +38,7 @@ void CreatePixelXYZ(std::shared_ptr<EchoL0Dataset> echoL0ds, QString outPixelXYZ
if (startcolidx + tempechocol >= freqcount) {
tempechocol = freqcount - startcolidx;
}
std::cout << "\r[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] imgxyz :\t" << startcolidx << "\t-\t" << startcolidx + tempechocol << " / " << freqcount << std::endl;
qDebug() << "\r[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz") << "] imgxyz :\t" << startcolidx << "\t-\t" << startcolidx + tempechocol << " / " << freqcount ;
Eigen::MatrixXd demx = xyzRaster.getData(0, startcolidx, prfcount, tempechocol, 1);
Eigen::MatrixXd demy = xyzRaster.getData(0, startcolidx, prfcount, tempechocol, 2);
@ -191,7 +191,7 @@ ErrorCode TBPImageAlgCls::Process(long num_thread)
if (startline + templine >= imageheight) {
templine = imageheight - startline;
}
std::cout << "\r[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] imgxyz :\t" << startline << "\t-\t" << startline + templine << " / " << imageheight << std::endl;
qDebug() << "\r[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz") << "] imgxyz :\t" << startline << "\t-\t" << startline + templine << " / " << imageheight ;
std::shared_ptr<std::complex<double>> imageRaster = this->L1ds->getImageRaster(startline, templine);
@ -220,12 +220,15 @@ ErrorCode TBPImageAlgCls::ProcessGPU()
long pixelCount = rowCount * colCount;
long PRFCount = this->L0ds->getPluseCount();
long PlusePoints = this->L0ds->getPlusePoints();
long bandwidth = this->L0ds->getBandwidth();
double Rnear = this->L1ds->getNearRange();
double Rfar = this->L1ds->getFarRange();
double refRange = this->L0ds->getRefPhaseRange();
double dx = LIGHTSPEED / 2.0 / bandwidth; // c/2b
Rfar = Rnear + dx * (PlusePoints - 1); // 更新斜距距离
float Rnear = this->L1ds->getNearRange();
float Rfar = this->L1ds->getFarRange();
float fs = this->L1ds->getFs();
double dx = LIGHTSPEED / 2 / fs;
float freq = this->L1ds->getCenterFreq();
double factorj = freq * 4 * M_PI / LIGHTSPEED ;
@ -234,17 +237,25 @@ ErrorCode TBPImageAlgCls::ProcessGPU()
qDebug() << "Rnear:\t" << Rnear;
qDebug() << "Rfar:\t" << Rfar;
qDebug() << "refRange:\t" << this->getEchoL1()->getRefPhaseRange();
qDebug() << "fs:\t" << fs;
qDebug() << "dx:\t" << dx;
qDebug() << "freq:\t" << freq;
qDebug() << "rowCount:\t" << rowCount;
qDebug() << "colCount:\t" << colCount;
qDebug() << "PRFCount:\t" << PRFCount;
qDebug() << "PlusePoints:\t" << PlusePoints;
// 反方向计算起始相位
double deltaF = bandwidth / (PlusePoints - 1);
double startfreq = freq - bandwidth / 2;
double startlamda = LIGHTSPEED / startfreq;
std::shared_ptr<float> Pxs (new float[this->L0ds->getPluseCount()]);
std::shared_ptr<float> Pys (new float[this->L0ds->getPluseCount()]);
std::shared_ptr<float> Pzs (new float[this->L0ds->getPluseCount()]);
std::shared_ptr<double> Pxs (new double[this->L0ds->getPluseCount()]);
std::shared_ptr<double> Pys (new double[this->L0ds->getPluseCount()]);
std::shared_ptr<double> Pzs (new double[this->L0ds->getPluseCount()]);
{
std::shared_ptr<double> antpos = this->L0ds->getAntPos();
@ -263,13 +274,29 @@ ErrorCode TBPImageAlgCls::ProcessGPU()
}
antpos.reset();
}
// 计算成像的基本参数
// 距离向分辨率
double dr = LIGHTSPEED / 2.0 / (PRFCount * deltaF);
qDebug() << "------- resolution ----------------------------------";
qDebug() << "Range Resolution (m):\t" << dx ;
qDebug() << "Cross Resolution (m):\t" << dr;
qDebug() << "start Freq (Hz):\t" << startfreq;
qDebug() << "start lamda (m):\t" << startlamda;
qDebug() << "rowCount:\t" << rowCount;
qDebug() << "colCount:\t" << colCount;
qDebug() << "PRFCount:\t" << PRFCount;
qDebug() << "PlusePoints:\t" << PlusePoints;
qDebug() << "Rnear:\t" << Rnear;
qDebug() << "Rfar:\t" << Rfar;
qDebug() << "refRange:\t" << refRange;
// 方位向分辨率
// 按照回波分块,图像分块
long echoBlockline = Memory1GB / 8 / 2 / PlusePoints * 2;
long echoBlockline = Memory1GB / 8 / 2 / PlusePoints * 2; //2GB
echoBlockline = echoBlockline < 1 ? 1 : echoBlockline;
long imageBlockline = Memory1GB / 8 / 2 / colCount * 2;
long imageBlockline = Memory1GB / 8 / 2 / colCount * 2; //2GB
imageBlockline = imageBlockline < 1 ? 1 : imageBlockline;
gdalImage imageXYZ(this->outRasterXYZPath);
@ -280,11 +307,11 @@ ErrorCode TBPImageAlgCls::ProcessGPU()
if (startimgrowid + imageBlockline >= rowCount) {
tempimgBlockline = rowCount - startimgrowid;
}
std::cout << "\r[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] image create :\t" << startimgrowid << "\t-\t" << startimgrowid + tempimgBlockline<<"\t/\t"<< rowCount << std::endl;
qDebug() << "\r[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz") << "] image create :\t" << startimgrowid << "\t-\t" << startimgrowid + tempimgBlockline<<"\t/\t"<< rowCount ;
// 提取局部pixel x,y,z
std::shared_ptr<float> img_x = readDataArr<float>(imageXYZ,startimgrowid,0,tempimgBlockline,colCount,1,GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<float> img_y = readDataArr<float>(imageXYZ,startimgrowid,0,tempimgBlockline,colCount,2,GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<float> img_z = readDataArr<float>(imageXYZ,startimgrowid,0,tempimgBlockline,colCount,3,GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> img_x = readDataArr<double>(imageXYZ,startimgrowid,0,tempimgBlockline,colCount,1,GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> img_y = readDataArr<double>(imageXYZ,startimgrowid,0,tempimgBlockline,colCount,2,GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<double> img_z = readDataArr<double>(imageXYZ,startimgrowid,0,tempimgBlockline,colCount,3,GDALREADARRCOPYMETHOD::VARIABLEMETHOD);
std::shared_ptr<std::complex<double>> imgArr = this->L1ds->getImageRaster(startimgrowid, tempimgBlockline);
// 获取回波
@ -295,21 +322,21 @@ ErrorCode TBPImageAlgCls::ProcessGPU()
tempechoBlockline = PRFCount - startechoid;
}
std::shared_ptr<std::complex<double>> echoArr = this->L0ds->getEchoArr(startechoid, tempechoBlockline);
std::shared_ptr<float> antpx(new float[tempechoBlockline*PlusePoints]);
std::shared_ptr<float> antpy(new float[tempechoBlockline* PlusePoints]);
std::shared_ptr<float> antpz(new float[tempechoBlockline* PlusePoints]);
std::shared_ptr<double> antpx(new double[tempechoBlockline*PlusePoints],delArrPtr);
std::shared_ptr<double> antpy(new double[tempechoBlockline* PlusePoints], delArrPtr);
std::shared_ptr<double> antpz(new double[tempechoBlockline* PlusePoints], delArrPtr);
// 复制
for (long anti = 0; anti < tempechoBlockline; anti++) {
antpx.get()[anti] = Pxs.get()[anti + startechoid];
antpy.get()[anti] = Pys.get()[anti + startechoid];
antpz.get()[anti] = Pzs.get()[anti + startechoid];
}
TBPImageGPUAlg(
TBPImageGPUAlg2(
antpx, antpy, antpz,
img_x, img_y, img_z,
echoArr, imgArr,
freq, fs,
Rnear, Rfar,
startfreq, dx,
Rnear, Rfar, refRange,
tempimgBlockline, colCount,
tempechoBlockline, PlusePoints,
startechoid,startimgrowid
@ -324,192 +351,135 @@ ErrorCode TBPImageAlgCls::ProcessGPU()
return ErrorCode::SUCCESS;
}
void TBPImageGPUAlg(std::shared_ptr<float> antPx, std::shared_ptr<float> antPy, std::shared_ptr<float> antPz,
std::shared_ptr<float> imgx, std::shared_ptr<float> imgy, std::shared_ptr<float> imgz,
void TBPImageGPUAlg2(std::shared_ptr<double> antPx, std::shared_ptr<double> antPy, std::shared_ptr<double> antPz,
std::shared_ptr<double> img_x, std::shared_ptr<double> img_y, std::shared_ptr<double> img_z,
std::shared_ptr<std::complex<double>> echoArr,
std::shared_ptr<std::complex<double>> imgArr,
float freq, float fs, float Rnear, float Rfar,
std::shared_ptr<std::complex<double>> img_arr,
double freq, double dx, double Rnear, double Rfar, double refRange,
long rowcount, long colcount,
long prfcount, long freqcount,
long startPRFId, long startRowID
)
{
// 声明GPU变量
float* h_antPx = (float*)mallocCUDAHost(sizeof(float) * prfcount);
float* h_antPy = (float*)mallocCUDAHost(sizeof(float) * prfcount);
float* h_antPz = (float*)mallocCUDAHost(sizeof(float) * prfcount);
float* d_antPx = (float*)mallocCUDADevice(sizeof(float) * prfcount);
float* d_antPy = (float*)mallocCUDADevice(sizeof(float) * prfcount);
float* d_antPz = (float*)mallocCUDADevice(sizeof(float) * prfcount);
float* h_imgx = (float*)mallocCUDAHost(sizeof(float) * rowcount * colcount);
float* h_imgy = (float*)mallocCUDAHost(sizeof(float) * rowcount * colcount);
float* h_imgz = (float*)mallocCUDAHost(sizeof(float) * rowcount * colcount);
float* d_imgx = (float*)mallocCUDADevice(sizeof(float) * rowcount * colcount);
float* d_imgy = (float*)mallocCUDADevice(sizeof(float) * rowcount * colcount);
float* d_imgz = (float*)mallocCUDADevice(sizeof(float) * rowcount * colcount);
long IFFTPadNum = nextpow2(freqcount);
// 先处理脉冲傅里叶变换
cuComplex* h_echoArr = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * prfcount * freqcount);
cuComplex* d_echoArr = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * prfcount * freqcount);
std::shared_ptr<cuComplex> d_echoArrIFFT((cuComplex*)mallocCUDADevice(sizeof(cuComplex) * prfcount * IFFTPadNum), FreeCUDADevice);
cuComplex* h_imgArr = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * rowcount * colcount);
cuComplex* d_imgArr = (cuComplex*)mallocCUDADevice( sizeof(cuComplex) * rowcount * colcount);
cuComplex* h_imgEchoArr = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * rowcount * colcount);
cuComplex* d_imgEchoArr = (cuComplex*)mallocCUDADevice(sizeof(cuComplex) * rowcount * colcount);
float* h_R=(float*)mallocCUDAHost(sizeof(float) * rowcount * colcount);
float* d_R=(float*)mallocCUDADevice(sizeof(float) * rowcount * colcount);
long* h_CIdx = (long*)mallocCUDAHost(sizeof(long) * rowcount * colcount);
long* d_CIdx = (long*)mallocCUDADevice(sizeof(long) * rowcount * colcount);
// 回波赋值
for (long i = 0; i < prfcount; i++) {
for (long j = 0; j < freqcount; j++) {
h_echoArr[i * freqcount + j] = make_cuComplex(echoArr.get()[i * freqcount + j].real(),
echoArr.get()[i * freqcount + j].imag());
}
}
HostToDevice(h_echoArr, d_echoArr, sizeof(cuComplex) * prfcount * freqcount);
CUDAIFFT(d_echoArr, d_echoArrIFFT.get(), prfcount, freqcount, IFFTPadNum);
// 结束傅里叶变换
FreeCUDAHost(h_echoArr);
FreeCUDADevice(d_echoArr);
// 初始化
std::shared_ptr<double> h_antPx ((double*)mallocCUDAHost(sizeof(double) * prfcount),FreeCUDAHost);
std::shared_ptr<double> h_antPy ((double*)mallocCUDAHost(sizeof(double) * prfcount),FreeCUDAHost);
std::shared_ptr<double> h_antPz ((double*)mallocCUDAHost(sizeof(double) * prfcount),FreeCUDAHost);
std::shared_ptr<double> d_antPx ((double*)mallocCUDADevice(sizeof(double) * prfcount),FreeCUDADevice);
std::shared_ptr<double> d_antPy ((double*)mallocCUDADevice(sizeof(double) * prfcount),FreeCUDADevice);
std::shared_ptr<double> d_antPz ((double*)mallocCUDADevice(sizeof(double) * prfcount),FreeCUDADevice);
std::shared_ptr<double> h_imgx((double*)mallocCUDAHost(sizeof(double) * rowcount * colcount),FreeCUDAHost);
std::shared_ptr<double> h_imgy((double*)mallocCUDAHost(sizeof(double) * rowcount * colcount),FreeCUDAHost);
std::shared_ptr<double> h_imgz((double*)mallocCUDAHost(sizeof(double) * rowcount * colcount),FreeCUDAHost);
std::shared_ptr<double> d_imgx ((double*)mallocCUDADevice(sizeof(double) * rowcount * colcount),FreeCUDADevice);
std::shared_ptr<double> d_imgy ((double*)mallocCUDADevice(sizeof(double) * rowcount * colcount),FreeCUDADevice);
std::shared_ptr<double> d_imgz ((double*)mallocCUDADevice(sizeof(double) * rowcount * colcount),FreeCUDADevice);
// 天线位置
for (long i = 0; i < prfcount; i++) {
h_antPx[i] = antPx.get()[i];
h_antPy[i] = antPy.get()[i];
h_antPz[i] = antPz.get()[i];
h_antPx.get()[i] = antPx.get()[i];
h_antPy.get()[i] = antPy.get()[i];
h_antPz.get()[i] = antPz.get()[i];
}
// 成像平面
for (long i = 0; i < rowcount; i++) {
for (long j = 0; j < colcount; j++) {
h_imgx[i * colcount + j]=imgx.get()[i * colcount + j];
h_imgy[i * colcount + j]=imgy.get()[i * colcount + j];
h_imgz[i * colcount + j]=imgz.get()[i * colcount + j];
h_imgx.get()[i * colcount + j] = imgx.get()[i * colcount + j];
h_imgy.get()[i * colcount + j] = imgy.get()[i * colcount + j];
h_imgz.get()[i * colcount + j] = imgz.get()[i * colcount + j];
}
}
// 回波
for (long i = 0; i < prfcount; i++) {
for (long j = 0; j < freqcount; j++) {
h_echoArr[i * freqcount + j] = make_cuComplex(echoArr.get()[i * freqcount + j].real(),
echoArr.get()[i * freqcount + j].imag());
}
}
HostToDevice(h_antPx.get(), d_antPx.get(), sizeof(double) * prfcount);
HostToDevice(h_antPy.get(), d_antPy.get(), sizeof(double) * prfcount);
HostToDevice(h_antPz.get(), d_antPz.get(), sizeof(double) * prfcount);
HostToDevice(h_imgx.get(), d_imgx.get(), sizeof(double) * rowcount * colcount);
HostToDevice(h_imgy.get(), d_imgy.get(), sizeof(double) * rowcount * colcount);
HostToDevice(h_imgz.get(), d_imgz.get(), sizeof(double) * rowcount * colcount);
std::shared_ptr<cuComplex> h_imgArr((cuComplex*)mallocCUDAHost(sizeof(cuComplex) * rowcount * colcount), FreeCUDAHost);
std::shared_ptr<cuComplex> d_imgArr((cuComplex*)mallocCUDADevice(sizeof(cuComplex) * rowcount * colcount), FreeCUDADevice);
// 图像
for (long i = 0; i < rowcount; i++) {
for (long j = 0; j < colcount; j++) {
h_imgArr[i * colcount + j].x = imgArr.get()[i * colcount + j].real();
h_imgArr[i * colcount + j].y = imgArr.get()[i * colcount + j].imag();
h_imgArr.get()[i * colcount + j].x = imgArr.get()[i * colcount + j].real();
h_imgArr.get()[i * colcount + j].y = imgArr.get()[i * colcount + j].imag();
}
}
HostToDevice(h_imgArr.get(), d_imgArr.get(), sizeof(cuComplex) * rowcount * colcount);
// Host -> GPU
HostToDevice(h_antPx, d_antPx, sizeof(float) * prfcount);
HostToDevice(h_antPy, d_antPy, sizeof(float) * prfcount);
HostToDevice(h_antPz, d_antPz, sizeof(float) * prfcount);
// 直接使用
double startlamda = LIGHTSPEED / freq;
TimeBPImage(
d_antPx.get(), d_antPy.get(), d_antPz.get(),
d_imgx.get(), d_imgy.get(), d_imgz.get(),
d_echoArrIFFT.get(), prfcount, IFFTPadNum,
d_imgArr.get(), rowcount, colcount,
startlamda, Rnear, dx, refRange
);
HostToDevice(h_imgx, d_imgx, sizeof(float) * rowcount * colcount);
HostToDevice(h_imgy, d_imgy, sizeof(float) * rowcount * colcount);
HostToDevice(h_imgz, d_imgz, sizeof(float) * rowcount * colcount);
HostToDevice(h_R, d_R, sizeof(float) * rowcount * colcount);
HostToDevice(h_CIdx, d_CIdx, sizeof(long) * rowcount * colcount);
HostToDevice(h_echoArr, d_echoArr, sizeof(cuComplex) * prfcount * freqcount);
HostToDevice(h_imgArr, d_imgArr, sizeof(cuComplex) * rowcount * colcount);
HostToDevice(h_imgEchoArr, d_imgEchoArr, sizeof(cuComplex) * rowcount * colcount);
#ifdef __TBPIMAGEDEBUG__
// 定义采样点
long tc[4] = { 6956,6542,7003,6840};
long tr[4] = { 1100,9324,9415,11137};
for (long iii = 0; iii < 4; iii) {
tr[iii] = tr[iii] - startRowID;
}
std::shared_ptr<float> Rlist(new float[4*prfcount], delArrPtr);
std::shared_ptr<long> CIdslist(new long[4*prfcount], delArrPtr);
std::shared_ptr<float> imgchoReal (new float[4 * prfcount], delArrPtr);
std::shared_ptr<float> imgchoImag (new float[4 * prfcount], delArrPtr);
std::shared_ptr<float> imgdataReal(new float[4 * prfcount], delArrPtr);
std::shared_ptr<float> imgdataImag(new float[4 * prfcount], delArrPtr);
#endif
for (long prfid = 0; prfid < prfcount; prfid++) {
CUDATBPImage(
d_antPx,d_antPy,d_antPz,
d_imgx,d_imgy,d_imgz,
d_R,
d_CIdx,
d_echoArr,
d_imgArr,
d_imgEchoArr,
freq, fs, Rnear, Rfar,
rowcount, colcount,
prfid, freqcount
);
#ifdef __TBPIMAGEDEBUG__
DeviceToHost(h_R, d_R, sizeof(float) * rowcount * colcount);
DeviceToHost(h_CIdx, d_CIdx, sizeof(float) * rowcount * colcount);
DeviceToHost(h_imgEchoArr, d_imgEchoArr, sizeof(cuComplex) * rowcount * colcount);
DeviceToHost(h_imgArr, d_imgArr, sizeof(cuComplex) * rowcount * colcount);
for (long iii = 0; iii < 4; iii++) {
Rlist.get()[prfid * 4 + iii] = h_R[tr[iii] * colcount + tc[iii]];
CIdslist.get()[prfid * 4 + iii] = h_CIdx[tr[iii] * colcount + tc[iii]];
imgchoReal.get()[prfid * 4 + iii] = h_imgEchoArr[tr[iii] * colcount + tc[iii]].x;
imgchoImag.get()[prfid * 4 + iii] = h_imgEchoArr[tr[iii] * colcount + tc[iii]].y;
imgdataReal.get()[prfid * 4 + iii] = h_imgArr[tr[iii] * colcount + tc[iii]].x;
imgdataImag.get()[prfid * 4 + iii] = h_imgArr[tr[iii] * colcount + tc[iii]].y;
}
#endif
}
// Device -> Host
DeviceToHost(h_imgArr, d_imgArr, sizeof(cuComplex) * rowcount * colcount);
#ifdef __TBPIMAGEDEBUG__
testOutAmpArr(QString("Rlist_%1.bin").arg(startPRFId), Rlist.get(), prfcount, 4);
testOutAmpArr(QString("imgchoReal_%1.bin").arg(startPRFId), imgchoReal.get(), prfcount, 4);
testOutAmpArr(QString("imgchoImag_%1.bin").arg(startPRFId), imgchoImag.get(), prfcount, 4);
testOutAmpArr(QString("imgdataReal_%1.bin").arg(startPRFId), imgdataReal.get(), prfcount, 4);
testOutAmpArr(QString("imgdataImag_%1.bin").arg(startPRFId), imgdataImag.get(), prfcount, 4);
testOutClsArr(QString("CIdslist_%1.bin").arg(startPRFId), CIdslist.get(), prfcount, 4);
#endif
DeviceToHost(h_imgArr.get(), d_imgArr.get(), sizeof(cuComplex)* rowcount* colcount);
for (long i = 0; i < rowcount; i++) {
for (long j = 0; j < colcount; j++) {
imgArr.get()[i * colcount + j] = std::complex<double>(h_imgArr[i * colcount + j].x,h_imgArr[i * colcount + j].y);
imgArr.get()[i * colcount + j] = std::complex<double>(h_imgArr.get()[i * colcount + j].x, h_imgArr.get()[i * colcount + j].y);
}
}
FreeCUDAHost(h_antPx); FreeCUDADevice(d_antPx);
FreeCUDAHost(h_antPy); FreeCUDADevice(d_antPy);
FreeCUDAHost(h_antPz); FreeCUDADevice(d_antPz);
FreeCUDAHost(h_imgx); FreeCUDADevice(d_imgx);
FreeCUDAHost(h_imgy); FreeCUDADevice(d_imgy);
FreeCUDAHost(h_imgz); FreeCUDADevice(d_imgz);
FreeCUDAHost(h_echoArr); FreeCUDADevice(d_echoArr);
FreeCUDAHost(h_imgArr); FreeCUDADevice(d_imgArr);
FreeCUDAHost(h_R); FreeCUDADevice(d_R);
}
void TBPImageAlgCls::setGPU(bool flag)
{
this->GPURUN = flag;

View File

@ -65,12 +65,13 @@ void CreatePixelXYZ(std::shared_ptr<EchoL0Dataset> echoL0ds,QString outPixelXYZP
void TBPImageProcess(QString echofile,QString outImageFolder,QString imagePlanePath,long num_thread);
void TBPImageGPUAlg(std::shared_ptr<float> antPx, std::shared_ptr<float> antPy, std::shared_ptr<float> antPz,
std::shared_ptr<float> img_x, std::shared_ptr<float> img_y, std::shared_ptr<float> img_z,
std::shared_ptr<std::complex<double>> echoArr,
void TBPImageGPUAlg2(std::shared_ptr<double> antPx, std::shared_ptr<double> antPy, std::shared_ptr<double> antPz,
std::shared_ptr<double> img_x, std::shared_ptr<double> img_y, std::shared_ptr<double> img_z,
std::shared_ptr<std::complex<double>> echoArr,
std::shared_ptr<std::complex<double>> img_arr,
float freq, float fs, float Rnear, float Rfar,
double freq, double dx, double Rnear, double Rfar,double refRange,
long rowcount, long colcount,
long prfcount,long freqcount,
long startPRFId,long startRowID
);
long prfcount, long freqcount,
long startPRFId, long startRowID
);