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

1. 调整了RFPC的逻辑 

2. 增加了影像按行列裁剪工作
pull/3/head
陈增辉 2025-01-27 11:37:46 +08:00
parent 9cf05eae73
commit fb8f0409e1
25 changed files with 1148 additions and 971 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

33
BaseTool/BaseConstVariable.h
View File

@ -15,7 +15,7 @@
#define __CUDANVCC___ // 定义CUDA函数
#define __PRFDEBUG__
#define __PRFDEBUG_PRFINF__
//#define __PRFDEBUG_PRFINF__
//#define __ECHOTIMEDEBUG__
#define __TBPIMAGEDEBUG__
@ -171,6 +171,30 @@ struct PatternImageDesc {
};
struct CUDA_AntSate_PtrList {
long PRF_len = 0;
double* h_antpx = nullptr, * d_antpx = nullptr;
double* h_antpy = nullptr, * d_antpy = nullptr;
double* h_antpz = nullptr, * d_antpz = nullptr;
double* h_antvx = nullptr, * d_antvx = nullptr;
double* h_antvy = nullptr, * d_antvy = nullptr;
double* h_antvz = nullptr, * d_antvz = nullptr;
double* h_antdirectx = nullptr, * d_antdirectx = nullptr;
double* h_antdirecty = nullptr, * d_antdirecty = nullptr;
double* h_antdirectz = nullptr, * d_antdirectz = nullptr;
double* h_antXaxisX = nullptr, * d_antXaxisX = nullptr;
double* h_antXaxisY = nullptr, * d_antXaxisY = nullptr;
double* h_antXaxisZ = nullptr, * d_antXaxisZ = nullptr;
double* h_antYaxisX = nullptr, * d_antYaxisX = nullptr;
double* h_antYaxisY = nullptr, * d_antYaxisY = nullptr;
double* h_antYaxisZ = nullptr, * d_antYaxisZ = nullptr;
double* h_antZaxisX = nullptr, * d_antZaxisX = nullptr;
double* h_antZaxisY = nullptr, * d_antZaxisY = nullptr;
double* h_antZaxisZ = nullptr, * d_antZaxisZ = nullptr;
};
/*********************************************** 指针回收区域 ********************************************************************/
inline void delArrPtr(void* p)
@ -193,8 +217,11 @@ inline void PrintTime() {
/** 计算分块 ******************************************************************/
inline long getBlockRows(long sizeMB, long cols,long sizeMeta) {
return (round(Memory1MB * 1.0 / sizeMeta * sizeMB) + cols - 1) / cols;
inline long getBlockRows(long sizeMB, long cols,long sizeMeta,long maxRows) {
long rownum= (round(Memory1MB * 1.0 / sizeMeta / cols * sizeMB) + cols - 1);
rownum = rownum < 0 ? 1 : rownum;
rownum =rownum < maxRows ? rownum : maxRows;
return rownum;
}

2
BaseTool/EchoDataFormat.cpp
View File

@ -253,7 +253,7 @@ QString EchoL0Dataset::getEchoDataFilename()
void EchoL0Dataset::initEchoArr(std::complex<double> init0)
{
long blockline = Memory1MB * 2000 / 8 / 2 / this->PlusePoints;
long blockline = Memory1MB / 8 / 2 / this->PlusePoints * 8000;
long start = 0;
for (start = 0; start < this->PluseCount; start = start + blockline) {

25
BaseTool/ImageOperatorBase.cpp
View File

@ -1725,6 +1725,31 @@ int saveMatrixXcd2TiFF(Eigen::MatrixXcd data, QString out_tiff_path)
return -1;
}
void clipRaster(QString inRasterPath, QString outRasterPath, long minRow, long maxRow, long minCol, long maxCol)
{
long rownum = maxRow - minRow + 1;
long colnum = maxCol - minCol + 1;
gdalImage inimg(inRasterPath);
Eigen::MatrixXd gt = inimg.gt;
Landpoint lp = inimg.getLandPoint(minRow, minCol, 0);
gt(0, 0) = lp.lon;
gt(1, 0) = lp.lat;
gdalImage outimg= CreategdalImageDouble(outRasterPath, rownum, colnum, inimg.band_num, gt, inimg.projection, true, true, true);
for (long bi = 1; bi < inimg.band_num + 1; bi++) {
Eigen::MatrixXd brasterData = inimg.getData(minRow, minCol, rownum, colnum, bi);
outimg.saveImage(brasterData, 0, 0, bi);
qDebug() << "writer raster band : " << bi;
}
qDebug() << "writer raster overring";
}
ErrorCode MergeRasterProcess(QVector<QString> filepaths, QString outfileptah, QString mainString, MERGEMODE mergecode, bool isENVI, ShowProessAbstract* dia )
{

7
BaseTool/ImageOperatorBase.h
View File

@ -258,6 +258,11 @@ int saveMatrixXcd2TiFF(Eigen::MatrixXcd data, QString out_tiff_path);
//----------------------------------------------------
void clipRaster(QString inRasterPath, QString outRasterPath, long minRow, long maxRow, long minCol, long maxCol);
//--------------------- 图像合并流程 ------------------------------
enum MERGEMODE
{
@ -285,6 +290,8 @@ void testOutClsArr(QString filename, long* amp, long rowcount, long colcount);
//--------------------- 图像文件读写 ------------------------------
template<typename T>
std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, int rows_count, int cols_count, int band_ids, GDALREADARRCOPYMETHOD method)
{

110
BaseToolbox/QClipRasterByRowCols.cpp Normal file
View File

@ -0,0 +1,110 @@
#include "QClipRasterByRowCols.h"
#include <QFileDialog>
#include <QMessageBox>
#include "ImageOperatorBase.h"
QClipRasterByRowCols::QClipRasterByRowCols(QWidget *parent)
: QDialog(parent)
{
ui.setupUi(this);
connect(this->ui.accepBtn, SIGNAL(accepted()), this, SLOT(accepBtnaccept()));
connect(this->ui.accepBtn, SIGNAL(rejected()), this, SLOT(accepBtnreject()));
connect(this->ui.InRasterBtn, SIGNAL(clicked(bool)), this, SLOT(onInRasterBtnClicked(bool)));
connect(this->ui.OutRasterBtn, SIGNAL(clicked(bool)), this, SLOT(onOutRasterBtnClicked(bool)));
}
QClipRasterByRowCols::~QClipRasterByRowCols()
{
}
void QClipRasterByRowCols::accepBtnaccept()
{
QString inRasterPath = this->ui.lineEdit_InRaster->text();
QString outRasterPath = this->ui.lineEdit_OutRaster->text();
long minRow = this->ui.lineEdit_topRow->value();
long maxRow = this->ui.lineEdit_bottomRow->value();
long minCol = this->ui.lineEdit_LeftCol->value();
long maxCol = this->ui.lineEdit_RightCol->value();
if (maxCol < minCol || maxRow < minRow) {
QMessageBox::warning(nullptr, u8"警告", u8"裁剪行列范围填写错误");
return;
}
else {}
// 图像裁剪
this->ui.progressBar->setValue(10);
clipRaster(inRasterPath, outRasterPath, minRow, maxRow, minCol, maxCol);
this->ui.progressBar->setValue(100);
QMessageBox::information(nullptr, u8"信息", u8"影像处理完成");
this->ui.progressBar->setValue(0);
}
void QClipRasterByRowCols::accepBtnreject()
{
this->close();
}
void QClipRasterByRowCols::onInRasterBtnClicked(bool)
{
QString fileName = QFileDialog::getOpenFileName(
this, // 父窗口
tr(u8"选择xml文件"), // 标题
QString(), // 默认路径
tr(u8"tiff Files (*.tiff);;tif Files (*.tif);;dat Files (*.dat);;bin Files (*.bin);;All Files (*)") // 文件过滤器
);
// 如果用户选择了文件
if (!fileName.isEmpty()) {
QString message = "选择的文件有：\n";
this->ui.lineEdit_InRaster->setText(fileName);
}
else {
QMessageBox::information(this, tr(u8"没有选择文件"), tr(u8"没有选择任何文件"));
}
if (!fileName.isEmpty()) {
gdalImage inRaster(this->ui.lineEdit_InRaster->text());
this->ui.lineEdit_topRow->setMinimum(0);
this->ui.lineEdit_bottomRow->setMinimum(0);
this->ui.lineEdit_LeftCol->setMinimum(0);
this->ui.lineEdit_RightCol->setMinimum(0);
this->ui.lineEdit_topRow->setMaximum(inRaster.height);
this->ui.lineEdit_bottomRow->setMaximum(inRaster.height);
this->ui.lineEdit_LeftCol->setMaximum(inRaster.width);
this->ui.lineEdit_RightCol->setMaximum(inRaster.width);
}
}
void QClipRasterByRowCols::onOutRasterBtnClicked(bool)
{
QString fileName = QFileDialog::getSaveFileName(
this, // 父窗口
tr(u8"选择xml文件"), // 标题
QString(), // 默认路径
tr(u8"dat Files (*.dat);;bin Files (*.bin);;All Files (*)") // 文件过滤器
);
// 如果用户选择了文件
if (!fileName.isEmpty()) {
QString message = "选择的文件有：\n";
this->ui.lineEdit_OutRaster->setText(fileName);
}
else {
QMessageBox::information(this, tr(u8"没有选择文件"), tr(u8"没有选择任何文件。"));
}
}

26
BaseToolbox/QClipRasterByRowCols.h Normal file
View File

@ -0,0 +1,26 @@
#pragma once
// 根据行列范围参见影像
#include <QDialog>
#include "ui_QClipRasterByRowCols.h"
class QClipRasterByRowCols : public QDialog
{
Q_OBJECT
public:
QClipRasterByRowCols(QWidget *parent = nullptr);
~QClipRasterByRowCols();
public slots:
void accepBtnaccept();
void accepBtnreject();
void onInRasterBtnClicked(bool);
void onOutRasterBtnClicked(bool);
private:
Ui::QClipRasterByRowColsClass ui;
};

261
BaseToolbox/QClipRasterByRowCols.ui Normal file
View File

@ -0,0 +1,261 @@
<?xml version="1.0" encoding="UTF-8"?>
<ui version="4.0">
<class>QClipRasterByRowColsClass</class>
<widget class="QDialog" name="QClipRasterByRowColsClass">
<property name="geometry">
<rect>
<x>0</x>
<y>0</y>
<width>600</width>
<height>400</height>
</rect>
</property>
<property name="windowTitle">
<string>根据行列数裁剪影像</string>
</property>
<layout class="QVBoxLayout" name="verticalLayout">
<item>
<widget class="QFrame" name="frame">
<property name="frameShape">
<enum>QFrame::StyledPanel</enum>
</property>
<property name="frameShadow">
<enum>QFrame::Raised</enum>
</property>
<layout class="QGridLayout" name="gridLayout">
<item row="0" column="0">
<widget class="QLabel" name="label">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="maximumSize">
<size>
<width>16777215</width>
<height>30</height>
</size>
</property>
<property name="text">
<string>输入影像：</string>
</property>
</widget>
</item>
<item row="0" column="1">
<widget class="QLineEdit" name="lineEdit_InRaster">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="maximumSize">
<size>
<width>16777215</width>
<height>30</height>
</size>
</property>
</widget>
</item>
<item row="0" column="2">
<widget class="QPushButton" name="InRasterBtn">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="maximumSize">
<size>
<width>16777215</width>
<height>30</height>
</size>
</property>
<property name="text">
<string>选择</string>
</property>
</widget>
</item>
<item row="1" column="0">
<widget class="QLabel" name="label_2">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="maximumSize">
<size>
<width>16777215</width>
<height>30</height>
</size>
</property>
<property name="text">
<string>裁剪结果:</string>
</property>
</widget>
</item>
<item row="1" column="1">
<widget class="QLineEdit" name="lineEdit_OutRaster">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="maximumSize">
<size>
<width>16777215</width>
<height>30</height>
</size>
</property>
</widget>
</item>
<item row="1" column="2">
<widget class="QPushButton" name="OutRasterBtn">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="maximumSize">
<size>
<width>16777215</width>
<height>30</height>
</size>
</property>
<property name="text">
<string>选择</string>
</property>
</widget>
</item>
</layout>
</widget>
</item>
<item>
<widget class="QFrame" name="frame_2">
<property name="frameShape">
<enum>QFrame::StyledPanel</enum>
</property>
<property name="frameShadow">
<enum>QFrame::Raised</enum>
</property>
<layout class="QGridLayout" name="gridLayout_2">
<item row="3" column="3">
<widget class="QSpinBox" name="lineEdit_bottomRow">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="minimum">
<number>0</number>
</property>
<property name="maximum">
<number>999999999</number>
</property>
</widget>
</item>
<item row="2" column="2">
<widget class="QSpinBox" name="lineEdit_LeftCol">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="minimum">
<number>0</number>
</property>
<property name="maximum">
<number>999999999</number>
</property>
<property name="decimals" stdset="0">
<number>0</number>
</property>
</widget>
</item>
<item row="1" column="3">
<widget class="QSpinBox" name="lineEdit_topRow">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="minimum">
<number>0</number>
</property>
<property name="maximum">
<number>999999999</number>
</property>
</widget>
</item>
<item row="2" column="4">
<widget class="QSpinBox" name="lineEdit_RightCol">
<property name="minimumSize">
<size>
<width>0</width>
<height>30</height>
</size>
</property>
<property name="minimum">
<number>0</number>
</property>
<property name="maximum">
<number>999999999</number>
</property>
</widget>
</item>
<item row="0" column="2">
<widget class="QLabel" name="label_3">
<property name="maximumSize">
<size>
<width>16777215</width>
<height>25</height>
</size>
</property>
<property name="text">
<string>参见范围</string>
</property>
</widget>
</item>
</layout>
</widget>
</item>
<item>
<spacer name="verticalSpacer">
<property name="orientation">
<enum>Qt::Vertical</enum>
</property>
<property name="sizeHint" stdset="0">
<size>
<width>20</width>
<height>40</height>
</size>
</property>
</spacer>
</item>
<item>
<widget class="QProgressBar" name="progressBar">
<property name="value">
<number>0</number>
</property>
</widget>
</item>
<item>
<widget class="QDialogButtonBox" name="accepBtn">
<property name="standardButtons">
<set>QDialogButtonBox::Cancel|QDialogButtonBox::Ok</set>
</property>
</widget>
</item>
</layout>
</widget>
<layoutdefault spacing="6" margin="11"/>
<resources/>
<connections/>
</ui>

813
GPUTool/GPURFPC.cu
View File

@ -11,12 +11,12 @@
#include "BaseConstVariable.h"
#include "GPURFPC.cuh"
#ifdef __CUDANVCC___
/* 机器函数 ****************************************************************************************************************************/
__device__ double GPU_getSigma0dB(CUDASigmaParam param, double theta) {//线性值
@ -167,386 +167,10 @@ __device__ double GPU_BillerInterpAntPattern(double* antpattern,
return GainValue;
}
}
__device__ cuComplex GPU_calculationEcho(double sigma0, double TransAnt, double ReciveAnt,
double localangle, double R, double slopeangle, double Pt, double lamda) {
double amp = Pt * TransAnt * ReciveAnt;
amp = amp * sigma0;
amp = amp / (powf(4 * LAMP_CUDA_PI, 2) * powf(R, 4)); // 反射强度
double phi = (-4 * LAMP_CUDA_PI / lamda) * R;
cuComplex echophi = make_cuComplex(0, phi);
cuComplex echophiexp = cuCexpf(echophi);
cuComplex echo = make_cuComplex(echophiexp.x * amp, echophiexp.y * amp);
return echo;
}
__global__ void CUDA_SatelliteAntDirectNormal(double* RstX, double* RstY, double* RstZ,
double antXaxisX, double antXaxisY, double antXaxisZ,
double antYaxisX, double antYaxisY, double antYaxisZ,
double antZaxisX, double antZaxisY, double antZaxisZ,
double antDirectX, double antDirectY, double antDirectZ,
double* thetaAnt, double* phiAnt
, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
double Xst = -1 * RstX[idx]; // 卫星 --> 地面
double Yst = -1 * RstY[idx];
double Zst = -1 * RstZ[idx];
double AntXaxisX = antXaxisX;
double AntXaxisY = antXaxisY;
double AntXaxisZ = antXaxisZ;
double AntYaxisX = antYaxisX;
double AntYaxisY = antYaxisY;
double AntYaxisZ = antYaxisZ;
double AntZaxisX = antZaxisX;
double AntZaxisY = antZaxisY;
double AntZaxisZ = antZaxisZ;
// 归一化
double RstNorm = sqrtf(Xst * Xst + Yst * Yst + Zst * Zst);
double AntXaxisNorm = sqrtf(AntXaxisX * AntXaxisX + AntXaxisY * AntXaxisY + AntXaxisZ * AntXaxisZ);
double AntYaxisNorm = sqrtf(AntYaxisX * AntYaxisX + AntYaxisY * AntYaxisY + AntYaxisZ * AntYaxisZ);
double AntZaxisNorm = sqrtf(AntZaxisX * AntZaxisX + AntZaxisY * AntZaxisY + AntZaxisZ * AntZaxisZ);
double Rx = Xst / RstNorm;
double Ry = Yst / RstNorm;
double Rz = Zst / RstNorm;
double Xx = AntXaxisX / AntXaxisNorm;
double Xy = AntXaxisY / AntXaxisNorm;
double Xz = AntXaxisZ / AntXaxisNorm;
double Yx = AntYaxisX / AntYaxisNorm;
double Yy = AntYaxisY / AntYaxisNorm;
double Yz = AntYaxisZ / AntYaxisNorm;
double Zx = AntZaxisX / AntZaxisNorm;
double Zy = AntZaxisY / AntZaxisNorm;
double Zz = AntZaxisZ / AntZaxisNorm;
double Xant = (Rx * Yy * Zz - Rx * Yz * Zy - Ry * Yx * Zz + Ry * Yz * Zx + Rz * Yx * Zy - Rz * Yy * Zx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
double Yant = -(Rx * Xy * Zz - Rx * Xz * Zy - Ry * Xx * Zz + Ry * Xz * Zx + Rz * Xx * Zy - Rz * Xy * Zx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
double Zant = (Rx * Xy * Yz - Rx * Xz * Yy - Ry * Xx * Yz + Ry * Xz * Yx + Rz * Xx * Yy - Rz * Xy * Yx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
// 计算theta 与 phi
double Norm = sqrtf(Xant * Xant + Yant * Yant + Zant * Zant); // 计算 pho
double ThetaAnt = acosf(Zant / Norm); // theta 与 Z轴的夹角
double PhiAnt = atanf(Yant / Xant); // -pi/2 ~pi/2
if (abs(Yant) < PRECISIONTOLERANCE) { // X轴上
PhiAnt = 0;
}
else if (abs(Xant) < PRECISIONTOLERANCE) { // Y轴上，原点
if (Yant > 0) {
PhiAnt = PI / 2;
}
else {
PhiAnt = -PI / 2;
}
}
else if (Xant < 0) {
if (Yant > 0) {
PhiAnt = PI + PhiAnt;
}
else {
PhiAnt = -PI + PhiAnt;
}
}
else { // Xant>0 X 正轴
}
if (isnan(PhiAnt)) {
printf("V=[%f,%f,%f];norm=%f;thetaAnt=%f;phiAnt=%f;\n", Xant, Yant, Zant, Norm, ThetaAnt, PhiAnt);
}
//if (abs(ThetaAnt - 0) < PRECISIONTOLERANCE) {
// PhiAnt = 0;
//}
//else {}
thetaAnt[idx] = ThetaAnt * r2d;
phiAnt[idx] = PhiAnt * r2d;
//printf("Rst=[%f,%f,%f];AntXaxis = [%f, %f, %f];AntYaxis=[%f,%f,%f];AntZaxis=[%f,%f,%f];phiAnt=%f;thetaAnt=%f;\n", Xst, Yst, Zst
// , AntXaxisX, AntXaxisY, AntXaxisZ
// , AntYaxisX, AntYaxisY, AntYaxisZ
// , AntZaxisX, AntZaxisY, AntZaxisZ
// , phiAnt[idx]
// , thetaAnt[idx]
//);
}
}
__global__ void CUDA_BillerInterpAntPattern(double* antpattern,
double starttheta, double startphi, double dtheta, double dphi,
long thetapoints, long phipoints,
double* searththeta, double* searchphi, double* searchantpattern,
long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
double stheta = searththeta[idx];
double sphi = searchphi[idx];
double pthetaid = (stheta - starttheta) / dtheta;//
double pphiid = (sphi - startphi) / dphi;
long lasttheta = floorf(pthetaid);
long nextTheta = lasttheta + 1;
long lastphi = floorf(pphiid);
long nextPhi = lastphi + 1;
if (lasttheta < 0 || nextTheta < 0 || lastphi < 0 || nextPhi < 0 ||
lasttheta >= thetapoints || nextTheta >= thetapoints || lastphi >= phipoints || nextPhi >= phipoints)
{
searchantpattern[idx] = 0;
}
else {
double x = stheta;
double y = sphi;
double x1 = lasttheta * dtheta + starttheta;
double x2 = nextTheta * dtheta + starttheta;
double y1 = lastphi * dphi + startphi;
double y2 = nextPhi * dphi + startphi;
double z11 = antpattern[lasttheta * phipoints + lastphi];
double z12 = antpattern[lasttheta * phipoints + nextPhi];
double z21 = antpattern[nextTheta * phipoints + lastphi];
double z22 = antpattern[nextTheta * phipoints + nextPhi];
z11 = powf(10, z11 / 10);
z12 = powf(10, z12 / 10);
z21 = powf(10, z21 / 10);
z22 = powf(10, z22 / 10);
double GainValue = (z11 * (x2 - x) * (y2 - y)
+ z21 * (x - x1) * (y2 - y)
+ z12 * (x2 - x) * (y - y1)
+ z22 * (x - x1) * (y - y1));
GainValue = GainValue / ((x2 - x1) * (y2 - y1));
searchantpattern[idx] = GainValue;
}
}
}
__global__ void CUDA_AntPatternInterpGain(double* anttheta, double* antphi, double* gain,
double* antpattern, double starttheta, double startphi, double dtheta, double dphi, int thetapoints, int phipoints, long len) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
double temptheta = anttheta[idx];
double tempphi = antphi[idx];
double antPatternGain = GPU_BillerInterpAntPattern(antpattern,
starttheta, startphi, dtheta, dphi, thetapoints, phipoints,
temptheta, tempphi);
gain[idx] = antPatternGain;
}
}
__global__ void CUDA_InterpSigma(
long* demcls, double* sigmaAmp, double* localanglearr, long len,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
long clsid = demcls[idx];
double localangle = localanglearr[idx];
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2) {
sigmaAmp[idx] = 0;
}
else {}
if (abs(tempsigma.p1) < PRECISIONTOLERANCE &&
abs(tempsigma.p2) < PRECISIONTOLERANCE &&
abs(tempsigma.p3) < PRECISIONTOLERANCE &&
abs(tempsigma.p4) < PRECISIONTOLERANCE &&
abs(tempsigma.p5) < PRECISIONTOLERANCE &&
abs(tempsigma.p6) < PRECISIONTOLERANCE
) {
sigmaAmp[idx] = 0;
}
else {
double sigma = GPU_getSigma0dB(tempsigma, localangle);
sigma = powf(10.0, sigma / 10.0);// 后向散射系数
//printf("cls:%d;localangle=%f;sigma0=%f;\n", clsid, localangle, sigma);
sigmaAmp[idx] = sigma;
}
}
}
__global__ void CUDAKernel_RFPC_Computer_R_Gain(
double antX, double antY, double antZ, // 天线的坐标
double* targetX, double* targetY, double* targetZ, long len, // 地面坐标
long* demCls,
double* demSlopeX, double* demSlopeY, double* demSlopeZ, // 地表坡度矢量
double antXaxisX, double antXaxisY, double antXaxisZ, // 天线坐标系的X轴
double antYaxisX, double antYaxisY, double antYaxisZ,// 天线坐标系的Y轴
double antZaxisX, double antZaxisY, double antZaxisZ,// 天线坐标系的Z轴
double antDirectX, double antDirectY, double antDirectZ,// 天线的指向
double Pt,// 发射能量
double refPhaseRange,
double* TransAntpattern, double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图
double* ReceiveAntpattern, double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
double NearR, double FarR, // 距离范围
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
float* outR, // 输出距离
float* outAmp
) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
double tx = targetX[idx];
double ty = targetY[idx];
double tz = targetZ[idx];
double RstX = antX - tx; // 计算坐标矢量
double RstY = antY - ty;
double RstZ = antZ - tz;
double slopeX = demSlopeX[idx];
double slopeY = demSlopeY[idx];
double slopeZ = demSlopeZ[idx];
double RstR2 = RstX * RstX + RstY * RstY + RstZ * RstZ;
double RstR = sqrt(RstR2); // 矢量距离
//printf("idx=%d;antX=%f;antY=%f;antZ=%f;targetX=%f;targetY=%f;targetZ=%f;RstR=%.6f;diffR=%.6f;\n", idx,antX,antY,antZ,targetX,targetY,targetZ,RstR, RstR - 9.010858499003178e+05);
if (RstR<NearR || RstR>FarR) {
outAmp[idx] = 0;
outR[idx] = 0;
}
else {
// 求解坡度
double slopR = sqrtf(slopeX * slopeX + slopeY * slopeY + slopeZ * slopeZ); //
double dotAB = RstX * slopeX + RstY * slopeY + RstZ * slopeZ;
double localangle = acosf(dotAB / (RstR * slopR)); // 局地入射角
double ampGain = 0;
// 求解天线方向图指向
CUDAVectorEllipsoidal antVector = GPU_SatelliteAntDirectNormal(
RstX, RstY, RstZ,
antXaxisX, antXaxisY, antXaxisZ,
antYaxisX, antYaxisY, antYaxisZ,
antZaxisX, antZaxisY, antZaxisZ,
antDirectX, antDirectY, antDirectZ
);
if (antVector.Rho > 0) {
// 发射方向图
double temptheta = antVector.theta * r2d;
double tempphi = antVector.phi * r2d;
double TansantPatternGain =
GPU_BillerInterpAntPattern(
TransAntpattern,
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
temptheta, tempphi);
// 接收方向图
double antPatternGain = GPU_BillerInterpAntPattern(
ReceiveAntpattern,
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
temptheta, tempphi);
// 计算
double sigma0 = 0;
{
long clsid = demCls[idx];
//printf("clsid=%d\n", clsid);
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2) {
sigma0 = 0;
}
else {}
if (abs(tempsigma.p1) < PRECISIONTOLERANCE &&
abs(tempsigma.p2) < PRECISIONTOLERANCE &&
abs(tempsigma.p3) < PRECISIONTOLERANCE &&
abs(tempsigma.p4) < PRECISIONTOLERANCE &&
abs(tempsigma.p5) < PRECISIONTOLERANCE &&
abs(tempsigma.p6) < PRECISIONTOLERANCE
) {
sigma0 = 0;
}
else {
double sigma = GPU_getSigma0dB(tempsigma, localangle);
sigma0 = powf(10.0, sigma / 10.0);// 后向散射系数
}
}
ampGain = TansantPatternGain * antPatternGain;
ampGain = ampGain / (powf(4 * LAMP_CUDA_PI, 2) * powf(RstR, 4)); // 反射强度
outAmp[idx] = float(ampGain * Pt * sigma0);
outR[idx] = float(RstR - refPhaseRange);
//printf("%f-%f=%f\n", RstR , refPhaseRange, outR[idx]);
}
else {
}
}
}
}
__global__ void CUDAKernel_PRF_GeneratorEcho(float* Rarr, float* ampArr,
long pixelcount,
float f0, float dfreq,long freqnum,
float* echo_real, float* echo_imag, long prfid)
{
//// 假定共享内存大小为49152 byte
//// 假定每个Block 线程数大小为 32
__shared__ float s_R[GPU_SHARE_MEMORY]; // 距离 32*12 * 8= 49.2kb
__shared__ float s_Amp[GPU_SHARE_MEMORY]; // 振幅 3072 * 8= 49.2kb 49.2*2 = 98.4 < 100 KB
int idx = blockIdx.x * blockDim.x + threadIdx.x;; // 获取当前的线程编码
int tid = threadIdx.x;// 获取 单个 block 中的线程ID
const long startPIX = idx * GPU_SHARE_STEP; // 计算偏移
int curthreadidx = 0;
for (long i = 0; i < GPU_SHARE_STEP; i++) {
curthreadidx = i * BLOCK_SIZE + tid; // 计算分块
s_R[curthreadidx] = (startPIX + i) < pixelcount ? Rarr[startPIX + i] : 0.0;
s_Amp[curthreadidx] = (startPIX + i) < pixelcount ? ampArr[startPIX + i] : 0.0;
}
//__syncthreads(); // 确定所有待处理数据都已经进入程序中
if (startPIX < pixelcount) { // 存在可能处理的计算
float temp_real = 0;
float temp_imag = 0;
float factorjTemp = 0;
float temp_phi = 0;
float temp_amp = 0;
long dataid = 0;
curthreadidx = 0;
for (long fid = 0; fid < freqnum; fid++) {
factorjTemp = RFPCPIDIVLIGHT *(f0+ fid* dfreq);
//printf("factorj : %f , %f\n", factorjTemp, f0 + fid * dfreq);
temp_real = 0;
temp_imag = 0;
for (long j = 0; j < GPU_SHARE_STEP; j++) {
dataid = j * BLOCK_SIZE + tid;
temp_phi = s_R[dataid] * factorjTemp;
temp_amp = s_Amp[dataid];
temp_real += temp_amp* cosf(temp_phi);
temp_imag += temp_amp* sinf(temp_phi);
}
atomicAdd(&echo_real[prfid * freqnum + fid], temp_real); // 更新实部
atomicAdd(&echo_imag[prfid * freqnum + fid], temp_imag); // 更新虚部
}
}
}
/* 核函数 ****************************************************************************************************************************/
// 计算每块
__global__ void CUDA_Kernel_Computer_R_amp(
double* antX, double* antY, double* antZ,
@ -554,10 +178,10 @@ __global__ void CUDA_Kernel_Computer_R_amp(
double* antYaxisX, double* antYaxisY, double* antYaxisZ,
double* antZaxisX, double* antZaxisY, double* antZaxisZ,
double* antDirectX, double* antDirectY, double* antDirectZ,
long sPid, long PRFCount,
double* targetX, double* targetY, double* targetZ, long* demCls, long TargetNumber,
double* demSlopeX, double* demSlopeY, double* demSlopeZ,
long sPosId,long pixelcount,
long PRFCount, // 整体的脉冲数，
double* targetX, double* targetY, double* targetZ, long* demCls,
double* demSlopeX, double* demSlopeY, double* demSlopeZ ,
long startPosId, long pixelcount,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
double Pt,
double refPhaseRange,
@ -565,320 +189,263 @@ __global__ void CUDA_Kernel_Computer_R_amp(
double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints,
double* ReceiveAntpattern,
double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,
double maxTransAntPatternValue, double maxReceiveAntPatternValue,
double NearR, double FarR,
long BlockPRFCount,
long BlockPostions, // 模块
float* d_temp_R, float* d_temp_amps// 计算输出
) {
long idx = blockIdx.x * blockDim.x + threadIdx.x; // 获取当前的线程编码
long prfId = idx / BlockPostions;
long posId = idx % BlockPostions;
long aprfId = sPid + prfId;
long aposId = posId;
if (prfId< BlockPRFCount&& posId < BlockPostions &&(sPid + prfId) < PRFCount) {
double RstX = antX[aprfId] - targetX[aposId]; // 计算坐标矢量
double RstY = antY[aprfId] - targetY[aposId];
double RstZ = antZ[aprfId] - targetZ[aposId];
long prfId = idx / SHAREMEMORY_FLOAT_HALF;
long posId = idx % SHAREMEMORY_FLOAT_HALF+ startPosId; // 当前线程对应的影像点
if (prfId < PRFCount && posId < pixelcount) {
double RstX = antX[prfId] - targetX[posId]; // 计算坐标矢量
double RstY = antY[prfId] - targetY[posId];
double RstZ = antZ[prfId] - targetZ[posId];
double RstR = sqrt(RstX * RstX + RstY * RstY + RstZ * RstZ); // 矢量距离
if (RstR<NearR || RstR>FarR) {
d_temp_R[idx] = 0;
d_temp_amps[idx] = 0;
return;
}
else {
double slopeX = demSlopeX[aposId];
double slopeY = demSlopeY[aposId];
double slopeZ = demSlopeZ[aposId];
double slopeX = demSlopeX[posId];
double slopeY = demSlopeY[posId];
double slopeZ = demSlopeZ[posId];
double slopR = sqrtf(slopeX * slopeX + slopeY * slopeY + slopeZ * slopeZ); //
double dotAB = RstX * slopeX + RstY * slopeY + RstZ * slopeZ;
double localangle = acosf(dotAB / (RstR * slopR)); // 局地入射角
if (abs(slopR - 0) > 1e-3) {
double dotAB = RstX * slopeX + RstY * slopeY + RstZ * slopeZ;
double localangle = acos(dotAB / (RstR * slopR));
double ampGain = 0;
// 求解天线方向图指向
CUDAVectorEllipsoidal antVector = GPU_SatelliteAntDirectNormal(
RstX, RstY, RstZ,
antXaxisX[aprfId], antXaxisY[aprfId], antXaxisZ[aprfId],
antYaxisX[aprfId], antYaxisY[aprfId], antYaxisZ[aprfId],
antZaxisX[aprfId], antZaxisY[aprfId], antZaxisZ[aprfId],
antDirectX[aprfId], antDirectY[aprfId], antDirectZ[aprfId]
);
antVector.theta = antVector.theta * r2d;
antVector.phi = antVector.phi * r2d;
if (antVector.Rho > 0) {
double TansantPatternGain = GPU_BillerInterpAntPattern(
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2|| isnan(localangle)) {
d_temp_R[idx] = 0;
d_temp_amps[idx] = 0;
return;
}
else {}
double ampGain = 0;
// 求解天线方向图指向
CUDAVectorEllipsoidal antVector = GPU_SatelliteAntDirectNormal(
RstX, RstY, RstZ,
antXaxisX[prfId], antXaxisY[prfId], antXaxisZ[prfId],
antYaxisX[prfId], antYaxisY[prfId], antYaxisZ[prfId],
antZaxisX[prfId], antZaxisY[prfId], antZaxisZ[prfId],
antDirectX[prfId], antDirectY[prfId], antDirectZ[prfId]
);
antVector.theta = antVector.theta * r2d;
antVector.phi = antVector.phi * r2d;
//printf("theta: %f , phi: %f \n", antVector.theta, antVector.phi);
if (antVector.Rho > 0) {
double TansantPatternGain = GPU_BillerInterpAntPattern(
TransAntpattern,
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
antVector.theta, antVector.phi);
double antPatternGain = GPU_BillerInterpAntPattern(
ReceiveAntpattern,
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
antVector.theta, antVector.phi);
double antPatternGain = GPU_BillerInterpAntPattern(
ReceiveAntpattern,
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
antVector.theta, antVector.phi);
double sigma0 = 0;
{
long clsid = demCls[idx];
//printf("clsid=%d\n", clsid);
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2) {
sigma0 = 0;
double sigma0 = 0;
{
long clsid = demCls[posId];
//printf("clsid=%d\n", clsid);
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
if (abs(tempsigma.p1) < PRECISIONTOLERANCE &&
abs(tempsigma.p2) < PRECISIONTOLERANCE &&
abs(tempsigma.p3) < PRECISIONTOLERANCE &&
abs(tempsigma.p4) < PRECISIONTOLERANCE &&
abs(tempsigma.p5) < PRECISIONTOLERANCE &&
abs(tempsigma.p6) < PRECISIONTOLERANCE
) {
sigma0 = 0;
}
else {
double sigma = GPU_getSigma0dB(tempsigma, localangle);
sigma0 = powf(10.0, sigma / 10.0);
}
}
ampGain = TansantPatternGain * antPatternGain;
if (10 * log10(ampGain / maxReceiveAntPatternValue / maxTransAntPatternValue) < -3) { // 小于-3dB
d_temp_R[idx] = 0;
d_temp_amps[idx] = 0;
return;
}
else {}
if (abs(tempsigma.p1) < PRECISIONTOLERANCE &&
abs(tempsigma.p2) < PRECISIONTOLERANCE &&
abs(tempsigma.p3) < PRECISIONTOLERANCE &&
abs(tempsigma.p4) < PRECISIONTOLERANCE &&
abs(tempsigma.p5) < PRECISIONTOLERANCE &&
abs(tempsigma.p6) < PRECISIONTOLERANCE
) {
sigma0 = 0;
}
else {
double sigma = GPU_getSigma0dB(tempsigma, localangle);
sigma0 = powf(10.0, sigma / 10.0);// 后向散射系数
}
ampGain = ampGain / (powf(4 * LAMP_CUDA_PI, 2) * powf(RstR, 4)); // 反射强度
d_temp_amps[idx] = float(ampGain * Pt * sigma0);
d_temp_R[idx] = float(RstR - refPhaseRange);
return;
}
else {
d_temp_R[idx] = 0;
d_temp_amps[idx] = 0;
return;
}
ampGain = TansantPatternGain * antPatternGain;
ampGain = ampGain / (powf(4 * LAMP_CUDA_PI, 2) * powf(RstR, 4)); // 反射强度
d_temp_amps[idx] = float(ampGain * Pt * sigma0);
d_temp_R[idx] = float(RstR - refPhaseRange);
}
else {
d_temp_R[idx] = 0;
d_temp_amps[idx] = 0;
return;
}
}
}
}
__global__ void CUDA_Kernel_Computer_echo(
float* d_temp_R, float* d_temp_amps,long posNum,
float f0, float dfreq, long FreqPoints,long maxfreqnum,
float* d_temp_R, float* d_temp_amps, long posNum,
float f0, float dfreq,
long FreqPoints, // 当前频率的分块
long maxfreqnum, // 最大脉冲值
float* d_temp_echo_real, float* d_temp_echo_imag,
long temp_PRF_Count
) {// * blockDim.x + threadIdx.x;
__shared__ float s_R[SHAREMEMORY_FLOAT_HALF] ;
__shared__ float s_amp[SHAREMEMORY_FLOAT_HALF] ;
) {
__shared__ float s_R[SHAREMEMORY_FLOAT_HALF]; // 注意一个完整的block_size 共享相同内存
__shared__ float s_amp[SHAREMEMORY_FLOAT_HALF];
long tid = threadIdx.x;
long tid = threadIdx.x;
long bid = blockIdx.x;
long idx= bid * blockDim.x + tid;
long idx = bid * blockDim.x + tid;
long prfId = idx / FreqPoints; // 脉冲ID
long fId = idx % FreqPoints;//频率ID
long psid = 0;
for (long ii = 0; ii < BLOCK_SIZE; ii++) {
psid = tid * BLOCK_SIZE + ii;
s_R[psid] = d_temp_R[psid];
s_amp[psid] = d_temp_amps[psid];
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(); // 确定所有待处理数据都已经进入程序中
long prfId = idx / FreqPoints; // 脉冲
long fId = idx % FreqPoints;// 频率
if (fId < maxfreqnum&& prfId< temp_PRF_Count) {
if (fId < maxfreqnum && prfId < temp_PRF_Count) {
long echo_ID = prfId * maxfreqnum + fId; // 计算对应的回波位置
float factorjTemp = RFPCPIDIVLIGHT * (f0 + fId * dfreq);
float temp_real = 0;
float temp_imag = 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];
temp_real += temp_amp * cosf(temp_phi);
temp_imag += temp_amp * sinf(temp_phi);
temp_real += (temp_amp * cosf(temp_phi));
temp_imag += (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]);
//}
}
d_temp_echo_real[idx] += temp_real;
d_temp_echo_imag[idx] += temp_imag;
//printf("echo_ID=%d; ehodata=(%f,%f)\n", echo_ID, temp_real, temp_imag);
//printf("(%f %f %f) ", factorjTemp, s_amp[0], s_R[0]);
d_temp_echo_real[echo_ID] += /*d_temp_echo_real[echo_ID] + */temp_real;
d_temp_echo_imag[echo_ID] += /*d_temp_echo_imag[echo_ID] +*/ temp_imag;
}
}
/**
*
* 分块计算主流程
*/
void CUDA_RFPC_MainProcess(
double* antX, double* antY, double* antZ,
double* antXaxisX, double* antXaxisY, double* antXaxisZ,
double* antYaxisX, double* antYaxisY, double* antYaxisZ,
double* antZaxisX, double* antZaxisY, double* antZaxisZ,
double* antDirectX, double* antDirectY, double* antDirectZ,
long PRFCount, long FreqNum,
float f0, float dfreq,
double Pt,
double refPhaseRange,
double* TransAntpattern,
double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints,
double* ReceiveAntpattern,
double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,
double NearR, double FarR,
double* targetX, double* targetY, double* targetZ, long* demCls, long TargetNumber,
double* antX, double* antY, double* antZ,
double* antXaxisX, double* antXaxisY, double* antXaxisZ,
double* antYaxisX, double* antYaxisY, double* antYaxisZ,
double* antZaxisX, double* antZaxisY, double* antZaxisZ,
double* antDirectX, double* antDirectY, double* antDirectZ,
long PRFCount, long FreqNum,
float f0, float dfreq,
double Pt,
double refPhaseRange,
double* TransAntpattern,
double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints,
double* ReceiveAntpattern,
double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,
double maxTransAntPatternValue, double maxReceiveAntPatternValue,
double NearR, double FarR,
double* targetX, double* targetY, double* targetZ, long* demCls, long TargetNumber,
double* demSlopeX, double* demSlopeY, double* demSlopeZ,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
float* out_echoReal, float* out_echoImag)
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
float* out_echoReal, float* out_echoImag,
float* d_temp_R, float* d_temp_amp
)
{
long TargetNumberPerIter = 1024;
long maxPositionNumber = (SHAREMEMORY_BYTE / 2 / sizeof(double));
long freqpoints = NextBlockPad(FreqNum, BLOCK_SIZE); // 内存分布情况
long BlockPRFCount = getBlockRows(2000, freqpoints, sizeof(double));
long BlockTarlist = getBlockRows(2000, BlockPRFCount, sizeof(double));//1GB
BlockTarlist = BlockTarlist > SHAREMEMORY_FLOAT_HALF ? SHAREMEMORY_FLOAT_HALF : BlockTarlist;
double* h_tX = (double*)mallocCUDAHost(sizeof(double) * BlockTarlist);
double* h_tY = (double*)mallocCUDAHost(sizeof(double) * BlockTarlist);
double* h_tZ = (double*)mallocCUDAHost(sizeof(double) * BlockTarlist);
double* h_sloperX = (double*)mallocCUDAHost(sizeof(double) * BlockTarlist);
double* h_sloperY = (double*)mallocCUDAHost(sizeof(double) * BlockTarlist);
double* h_sloperZ = (double*)mallocCUDAHost(sizeof(double) * BlockTarlist);
long* h_cls = (long*)mallocCUDAHost(sizeof(long) * BlockTarlist);
double* d_tX = (double*)mallocCUDADevice(sizeof(double) * BlockTarlist);
double* d_tY = (double*)mallocCUDADevice(sizeof(double) * BlockTarlist);
double* d_tZ = (double*)mallocCUDADevice(sizeof(double) * BlockTarlist);
double* d_sloperX = (double*)mallocCUDADevice(sizeof(double) * BlockTarlist);
double* d_sloperY = (double*)mallocCUDADevice(sizeof(double) * BlockTarlist);
double* d_sloperZ = (double*)mallocCUDADevice(sizeof(double) * BlockTarlist);
long* d_cls = (long*)mallocCUDADevice(sizeof(long) * BlockTarlist);
float* d_temp_R = (float*)mallocCUDADevice(sizeof(float) * BlockPRFCount * BlockTarlist); //2GB 距离
float* d_temp_amp = (float*)mallocCUDADevice(sizeof(float) * BlockPRFCount * BlockTarlist);//2GB 强度
float* d_temp_echo_real = (float*)mallocCUDADevice(sizeof(float) * BlockPRFCount * freqpoints);//2GB
float* d_temp_echo_imag = (float*)mallocCUDADevice(sizeof(float) * BlockPRFCount * freqpoints);//2GB
float* h_temp_echo_real = (float*)mallocCUDAHost(sizeof(float) * BlockPRFCount * freqpoints);//2GB
float* h_temp_echo_imag = (float*)mallocCUDAHost(sizeof(float) * BlockPRFCount * freqpoints);//2GB
long BLOCK_FREQNUM = NextBlockPad(FreqNum, BLOCK_SIZE); // 256*freqBlockID
long cudaBlocknum = 0;
for (long spid = 0; spid < PRFCount; spid = spid + BlockPRFCount) {
// step 0 ,初始化
{
cudaBlocknum = (BlockPRFCount * freqpoints + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDAKernel_MemsetBlock << < cudaBlocknum, BLOCK_SIZE >> > (d_temp_echo_real, 0, BlockPRFCount * freqpoints);
CUDAKernel_MemsetBlock << < cudaBlocknum, BLOCK_SIZE >> > (d_temp_echo_imag, 0, BlockPRFCount * freqpoints);
}
long freqpoints = BLOCK_FREQNUM;
printf("freqpoints:%d\n", freqpoints);
long process = 0;
for (long sTi = 0; sTi < TargetNumber; sTi = sTi + SHAREMEMORY_FLOAT_HALF) {
cudaBlocknum = (PRFCount * SHAREMEMORY_FLOAT_HALF + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDA_Kernel_Computer_R_amp << <cudaBlocknum, BLOCK_SIZE >> > (
antX, antY, antZ,
antXaxisX, antXaxisY, antXaxisZ,
antYaxisX, antYaxisY, antYaxisZ,
antZaxisX, antZaxisY, antZaxisZ,
antDirectX, antDirectY, antDirectZ,
PRFCount,
targetX, targetY, targetZ, demCls,
demSlopeX, demSlopeY, demSlopeZ,
sTi, TargetNumber,
sigma0Paramslist, sigmaparamslistlen,
Pt,
refPhaseRange,
TransAntpattern,
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
ReceiveAntpattern,
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
maxTransAntPatternValue, maxReceiveAntPatternValue,
NearR, FarR,
d_temp_R, d_temp_amp// 计算输出
);
for (long sTi = 0; sTi < TargetNumber; sTi = sTi + BlockTarlist) {
// step 1,地面参数-> GPU内存
{
for (long ii = 0; ii < BlockTarlist && (sTi + ii) < TargetNumber; ii++) {
h_tX[sTi + ii] = targetX[sTi + ii];
h_tY[sTi + ii] = targetY[sTi + ii];
h_tZ[sTi + ii] = targetZ[sTi + ii];
h_sloperX[sTi + ii] = demSlopeX[sTi + ii];
h_sloperY[sTi + ii] = demSlopeY[sTi + ii];
h_sloperZ[sTi + ii] = demSlopeZ[sTi + ii];
h_cls[sTi + ii] = demCls[sTi + ii];
}
PRINT("Host -> Device start ,BlockTarlist %d \n", BlockTarlist);
HostToDevice(h_tX, d_tX, sizeof(double) * BlockTarlist);
HostToDevice(h_tY, d_tY, sizeof(double) * BlockTarlist);
HostToDevice(h_tZ, d_tZ, sizeof(double) * BlockTarlist);
HostToDevice(h_sloperX, d_sloperX, sizeof(double) * BlockTarlist);
HostToDevice(h_sloperY, d_sloperY, sizeof(double) * BlockTarlist);
HostToDevice(h_sloperZ, d_sloperZ, sizeof(double) * BlockTarlist);
HostToDevice(h_cls, d_cls, sizeof(long) * BlockTarlist);
PRINT("Host -> Device finished \n");
}
// step 2 计算距离
{
cudaBlocknum = (BlockPRFCount * BlockTarlist + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDA_Kernel_Computer_R_amp << <cudaBlocknum, BLOCK_SIZE >> > (
antX, antY, antZ,
antXaxisX, antXaxisY, antXaxisZ,
antYaxisX, antYaxisY, antYaxisZ,
antZaxisX, antZaxisY, antZaxisZ,
antDirectX, antDirectY, antDirectZ,
spid, PRFCount,
d_tX, d_tY, d_tZ, d_cls, BlockTarlist,
d_sloperX, d_sloperY, d_sloperZ,
sTi, TargetNumber,
sigma0Paramslist, sigmaparamslistlen,
Pt,
refPhaseRange,
TransAntpattern,
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
ReceiveAntpattern,
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
NearR, FarR,
BlockPRFCount,
BlockTarlist, // 模块
d_temp_R, d_temp_amp// 计算输出
);
}
// step 3 计算回波
{
cudaBlocknum = (BlockPRFCount * freqpoints + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDA_Kernel_Computer_echo << <cudaBlocknum, BLOCK_SIZE >> > (
d_temp_R, d_temp_amp, BlockTarlist,
f0, dfreq, freqpoints, FreqNum,
d_temp_echo_real, d_temp_echo_imag,
BlockPRFCount
);
}
PRINT("PRF %d / %d , TargetID: %d / %d \n", spid, PRFCount, sTi, sTi+ BlockTarlist);
PrintLasterError("CUDA_Kernel_Computer_R_amp");
cudaBlocknum = (PRFCount * BLOCK_FREQNUM + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDA_Kernel_Computer_echo << <cudaBlocknum, BLOCK_SIZE >> > (
d_temp_R, d_temp_amp, SHAREMEMORY_FLOAT_HALF,
f0, dfreq,
freqpoints, FreqNum,
out_echoReal, out_echoImag,
PRFCount
);
PrintLasterError("CUDA_Kernel_Computer_echo");
if ((sTi * 100.0 / TargetNumber ) - process >= 1) {
process = sTi * 100.0 / TargetNumber;
PRINT("TargetID [%f]: %d / %d finished\n", sTi*100.0/ TargetNumber,sTi, TargetNumber);
}
DeviceToDevice(h_temp_echo_real, d_temp_echo_real, sizeof(float) * BlockPRFCount * freqpoints);
DeviceToDevice(h_temp_echo_imag, d_temp_echo_imag, sizeof(float) * BlockPRFCount * freqpoints);
for (long ii = 0; ii < BlockPRFCount ; ii++) {
for (long jj = 0; jj < FreqNum; ii++) {
out_echoReal[(ii+spid) * FreqNum + jj] += h_temp_echo_real[ii * FreqNum + jj];
out_echoImag[(ii+spid) * FreqNum + jj] += h_temp_echo_imag[ii * FreqNum + jj];
}
}
//PRINT("");
}
// 显卡内存释放
FreeCUDAHost(h_tX);
FreeCUDAHost(h_tY);
FreeCUDAHost(h_tZ);
FreeCUDAHost(h_sloperX);
FreeCUDAHost(h_sloperY);
FreeCUDAHost(h_sloperZ);
FreeCUDAHost(h_cls);
FreeCUDADevice(d_tX);
FreeCUDADevice(d_tY);
FreeCUDADevice(d_tZ);
FreeCUDADevice(d_sloperX);
FreeCUDADevice(d_sloperY);
FreeCUDADevice(d_sloperZ);
FreeCUDADevice(d_cls);
FreeCUDADevice(d_temp_R);
FreeCUDADevice(d_temp_amp);
FreeCUDAHost(h_temp_echo_real);
FreeCUDAHost(h_temp_echo_imag);
FreeCUDADevice(d_temp_echo_real);
FreeCUDADevice(d_temp_echo_imag);
cudaDeviceSynchronize();
}
#endif

94
GPUTool/GPURFPC.cuh
View File

@ -11,10 +11,6 @@
#define RFPCPIDIVLIGHT -4*PI/(LIGHTSPEED/1e9)
#define GPU_SHARE_MEMORY 5888
#define GPU_SHARE_STEP 23
extern "C" struct CUDASigmaParam {
double p1;
double p2;
@ -25,65 +21,6 @@ extern "C" struct CUDASigmaParam {
};
extern __device__ double GPU_getSigma0dB(CUDASigmaParam param, double theta);
extern __device__ CUDAVectorEllipsoidal GPU_SatelliteAntDirectNormal(
double RstX, double RstY, double RstZ,
double antXaxisX, double antXaxisY, double antXaxisZ,
double antYaxisX, double antYaxisY, double antYaxisZ,
double antZaxisX, double antZaxisY, double antZaxisZ,
double antDirectX, double antDirectY, double antDirectZ
);
extern __device__ double GPU_BillerInterpAntPattern(double* antpattern,
double starttheta, double startphi, double dtheta, double dphi,
long thetapoints, long phipoints,
double searththeta, double searchphi);
extern __global__ void CUDA_AntPatternInterpGain(double* anttheta, double* antphi, double* gain,
double* antpattern, double starttheta, double startphi, double dtheta, double dphi, int thetapoints, int phipoints, long len);
extern __global__ void CUDA_InterpSigma(
long* demcls, double* sigmaAmp, double* localanglearr, long len,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen);
extern __global__ void CUDA_BillerInterpAntPattern(double* antpattern,
double starttheta, double startphi, double dtheta, double dphi,
long thetapoints, long phipoints,
double* searththeta, double* searchphi, double* searchantpattern,
long len);
extern __global__ void CUDAKernel_RFPC_Computer_R_Gain(
double antX, double antY, double antZ, // 天线的坐标
double* targetX, double* targetY, double* targetZ, long len, // 地面坐标
long* demCls,
double* demSlopeX, double* demSlopeY, double* demSlopeZ, // 地表坡度矢量
double antXaxisX, double antXaxisY, double antXaxisZ, // 天线坐标系的X轴
double antYaxisX, double antYaxisY, double antYaxisZ,// 天线坐标系的Y轴
double antZaxisX, double antZaxisY, double antZaxisZ,// 天线坐标系的Z轴
double antDirectX, double antDirectY, double antDirectZ,// 天线的指向
double Pt,// 发射能量
double refPhaseRange,
double* TransAntpattern, double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图
double* ReceiveAntpattern, double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
double NearR, double FarR, // 距离范围
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
float* outR, // 输出距离
float* outAmp
);
extern __global__ void CUDAKernel_PRF_GeneratorEcho(float* Rarr, float* ampArr,
long pixelcount,
float f0, float dfreq, long freqnum,
float* echo_real, float* echo_imag, long prfid);
@ -97,28 +34,6 @@ extern __global__ void CUDAKernel_PRF_GeneratorEcho(float* Rarr, float* ampArr,
//
//extern "C" void CUDA_RFPC_MainBlock(
// double* antX, double* antY, double* antZ, // 天线的坐标
// double* antXaxisX, double* antXaxisY, double* antXaxisZ, // 天线坐标系的X轴
// double* antYaxisX, double* antYaxisY, double* antYaxisZ,// 天线坐标系的Y轴
// double* antZaxisX, double* antZaxisY, double* antZaxisZ,// 天线坐标系的Z轴
// double* antDirectX, double* antDirectY, double* antDirectZ,// 天线的指向
// long startpid, long PRFCount, // 脉冲数
// float f0, float dfreq, long freqnum, // 频率数
// double* targetX, double* targetY, double* targetZ, long TargetPixelNumber, // 地面坐标
// long* demCls, // 地表类别
// double* demSlopeX, double* demSlopeY, double* demSlopeZ, // 地表坡度矢量
// double NearR, double FarR, // 距离范围
//
// float* out_echoReal, float* out_echoImag,// 输出回波
// float* temp_R, float* temp_amp
// //,double* temp_phi ,double* temp_real, double* tmep_imag// 临时变量
//);
extern "C" void CUDA_RFPC_MainProcess(
// 天线
double* antX, double* antY, double* antZ, // 天线坐标
@ -126,22 +41,25 @@ extern "C" void CUDA_RFPC_MainProcess(
double* antYaxisX, double* antYaxisY, double* antYaxisZ,// 天线坐标系的Y轴
double* antZaxisX, double* antZaxisY, double* antZaxisZ,// 天线坐标系的Z轴
double* antDirectX, double* antDirectY, double* antDirectZ,// 天线的指向
long PRFCount, long FreqNum, // 脉冲数量，频率数量
long PRFCount, long FreqNum, // 脉冲数量，频率数量
float f0, float dfreq,// 起始频率，终止频率
double Pt,// 发射能量
double refPhaseRange,
// 天线方向图
double* TransAntpattern, double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图
double* ReceiveAntpattern, double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
double maxTransAntPatternValue,double maxReceiveAntPatternValue,
double NearR, double FarR, // 距离范围
// 地面
double* targetX, double* targetY, double* targetZ, long* demCls, long TargetPixelNumber, // 地面坐标、地表覆盖类型，像素数
double* demSlopeX, double* demSlopeY, double* demSlopeZ,// 地表坡度矢量
double* demSlopeX, double* demSlopeY, double* demSlopeZ, // 地表坡度矢量
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图像
float* out_echoReal, float* out_echoImag// 输出回波
float* out_echoReal, float* out_echoImag,// 输出回波
float* d_temp_R, float* d_temp_amp
);

2
GPUTool/GPUTBPImage.cu
View File

@ -94,7 +94,7 @@ extern "C" void CUDATBPImage(float* antPx, float* antPy, float* antPz,
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDATBPImage CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();

47
GPUTool/GPUTool.cu
View File

@ -225,7 +225,7 @@ extern "C" void CUDA_MemsetBlock(cuComplex* data, cuComplex init0, long len) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDAmake_VectorA_B CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -250,7 +250,7 @@ extern "C" void* mallocCUDAHost(long memsize) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("mallocCUDAHost CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -264,7 +264,7 @@ extern "C" void FreeCUDAHost(void* ptr) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("FreeCUDAHost CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -278,7 +278,7 @@ extern "C" void* mallocCUDADevice(long memsize) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("mallocCUDADevice CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -292,7 +292,7 @@ extern "C" void FreeCUDADevice(void* ptr) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("FreeCUDADevice CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -306,10 +306,10 @@ extern "C" void HostToDevice(void* hostptr, void* deviceptr, long memsize) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("HostToDevice CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
@ -319,7 +319,7 @@ extern "C" void DeviceToHost(void* hostptr, void* deviceptr, long memsize) {
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("DeviceToHost CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -331,8 +331,8 @@ void DeviceToDevice(void* s_deviceptr, void* t_deviceptr, long memsize)
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("DeviceToHost CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
printf("DeviceToDevice CUDA Error: %s\n", cudaGetErrorString(err));
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -351,7 +351,7 @@ extern "C" void CUDAdistanceAB(float* Ax, float* Ay, float* Az, float* Bx, float
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDAdistanceAB CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -368,7 +368,7 @@ extern "C" void CUDABdistanceAs(float* Ax, float* Ay, float* Az, float Bx, float
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDABdistanceAs CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -384,7 +384,7 @@ extern "C" void CUDAmake_VectorA_B(float sX, float sY, float sZ, float* tX, floa
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDAmake_VectorA_B CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -401,7 +401,7 @@ extern "C" void CUDANorm_Vector(float* Vx, float* Vy, float* Vz, float* R, long
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDANorm_Vector CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -417,7 +417,7 @@ extern "C" void CUDAcosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, f
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDAcosAngle_VA_AB CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -435,7 +435,7 @@ extern "C" void CUDAGridPointLinearInterp1(float* v, float* q, float* qv, long
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDALinearInterp1 CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -452,7 +452,7 @@ extern "C" void CUDADSin(double* y, double* X, int n)
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("sin CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -470,7 +470,7 @@ extern "C" void CUDADCos(double* y, double* X, int n)
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("sin CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
@ -482,6 +482,15 @@ long NextBlockPad(long num, long blocksize)
return ((num + blocksize - 1) / blocksize) * blocksize;
}
void PrintLasterError(const char* s)
{
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("%s: %s\n", s, cudaGetErrorString(err));
exit(2);
}
}
#endif
@ -503,7 +512,7 @@ extern "C" float CUDA_SUM(float* d_x, long N)
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDALinearInterp1 CUDA Error: %s\n", cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
exit(2);
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();

10
GPUTool/GPUTool.cuh
View File

@ -15,10 +15,11 @@
#define LAMP_CUDA_PI 3.141592653589793238462643383279
// SHAREMEMORY_FLOAT_HALF_STEP * BLOCK_SIZE = SHAREMEMORY_FLOAT_HALF
#define BLOCK_SIZE 256
#define SHAREMEMORY_BYTE 49152
#define SHAREMEMORY_FLOAT_HALF 6144
#define SHAREMEMORY_FLOAT_HALF_STEP 24
// ´ňÓĄGPU˛ÎĘý
void printDeviceInfo(int deviceId);
@ -55,11 +56,6 @@ extern __device__ float GPU_VectorNorm2(CUDAVector A);
extern __device__ float GPU_dotVector(CUDAVector A, CUDAVector B);
extern __device__ float GPU_CosAngle_VectorA_VectorB(CUDAVector A, CUDAVector B);
// ś¨ŇĺČŤžÖşŻĘý
extern __global__ void CUDA_DistanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* R, long len);
extern __global__ void CUDA_B_DistanceA(float* Ax, float* Ay, float* Az, float Bx, float By, float Bz, float* R, long len);
@ -105,7 +101,7 @@ extern "C" void CUDADCos(double* y, double* X, int n);
extern "C" long NextBlockPad(long num,long blocksize);
extern "C" void PrintLasterError(const char* s);
#endif

7
RasterProcessTool.cpp
View File

@ -29,8 +29,12 @@ void RasterProcessTool::addBoxToolItemSLOT(QToolAbstract* item)
if (parentItem && ui.treeWidgetToolBox->itemWidget(parentItem, 0) == nullptr) {
QTreeWidgetItem* actionItem = new QTreeWidgetItem(parentItem);
parentItem->addChild(actionItem);
QIcon icon(QString::fromUtf8(":/RasterProcessTool/toolicon"));
QPushButton* button = new QPushButton(ui.treeWidgetToolBox);
button->setIcon(icon);
button->setText(toolName);
button->setLayoutDirection(Qt::LeftToRight);
button->setStyleSheet("QPushButton { text-align: left; }");
ui.treeWidgetToolBox->setItemWidget(actionItem, 0, button);
connect(button, SIGNAL(clicked()), item, SLOT(excute()));
item->setParent(ui.treeWidgetToolBox);
@ -77,6 +81,9 @@ QTreeWidgetItem* RasterProcessTool::findOrCreateTopLevelItem( QString& name) {
// 如果没有找到，创建新的顶级节点
QTreeWidgetItem* newItem = new QTreeWidgetItem(ui.treeWidgetToolBox);
QIcon icon(QString::fromUtf8(":/RasterProcessTool/toolboxIcon"));
newItem->setIcon(0,icon);
newItem->setTextAlignment(0, Qt::AlignLeft);
newItem->setText(0, name);
return newItem;
}

4
RasterProcessTool.qrc
View File

@ -1,4 +1,6 @@
<RCC>
<qresource prefix="RasterProcessTool">
<qresource prefix="/RasterProcessTool">
<file alias="toolicon.png">resource/toolicon.png</file>
<file alias="toolboxIcon">resource/toolboxIcon.png</file>
</qresource>
</RCC>

12
RasterProcessTool.ui
View File

@ -6,8 +6,8 @@
<rect>
<x>0</x>
<y>0</y>
<width>761</width>
<height>404</height>
<width>1920</width>
<height>1080</height>
</rect>
</property>
<property name="windowTitle">
@ -21,8 +21,8 @@
<rect>
<x>0</x>
<y>0</y>
<width>761</width>
<height>23</height>
<width>1920</width>
<height>22</height>
</rect>
</property>
</widget>
@ -47,6 +47,10 @@
<property name="text">
<string>工具箱</string>
</property>
<property name="icon">
<iconset resource="RasterProcessTool.qrc">
<normaloff>:/RasterProcessTool/toolboxIcon</normaloff>:/RasterProcessTool/toolboxIcon</iconset>
</property>
</column>
</widget>
</item>

3
RasterProcessTool.vcxproj
View File

@ -123,6 +123,7 @@
<ClCompile Include="BaseToolbox\DEMLLA2XYZTool.cpp" />
<ClCompile Include="BaseToolbox\GF3CalibrationAndGeocodingClass.cpp" />
<ClCompile Include="BaseToolbox\GF3PSTNClass.cpp" />
<ClCompile Include="BaseToolbox\QClipRasterByRowCols.cpp" />
<ClCompile Include="BaseToolbox\QComplex2AmpPhase.cpp" />
<ClCompile Include="BaseToolbox\QImportGF3StripL1ADataset.cpp" />
<ClCompile Include="BaseToolbox\QOrthSlrRaster.cpp" />
@ -164,6 +165,7 @@
<QtRcc Include="Imageshow\qcustomplot.qrc" />
<QtRcc Include="RasterProcessTool.qrc" />
<QtUic Include="BaseToolbox\DEMLLA2XYZTool.ui" />
<QtUic Include="BaseToolbox\QClipRasterByRowCols.ui" />
<QtUic Include="BaseToolbox\QComplex2AmpPhase.ui" />
<QtUic Include="BaseToolbox\QImportGF3StripL1ADataset.ui" />
<QtUic Include="BaseToolbox\QOrthSlrRaster.ui" />
@ -195,6 +197,7 @@
<QtMoc Include="BaseToolbox\QImportGF3StripL1ADataset.h" />
<QtMoc Include="BaseToolbox\QOrthSlrRaster.h" />
<QtMoc Include="BaseToolbox\QRDOrthProcessClass.h" />
<QtMoc Include="BaseToolbox\QClipRasterByRowCols.h" />
<ClInclude Include="BaseToolbox\SatelliteGF3xmlParser.h" />
<ClInclude Include="BaseToolbox\SateOrbit.h" />
<ClInclude Include="BaseToolbox\simptsn.h" />

11
RasterProcessTool.vcxproj.filters
View File

@ -169,6 +169,9 @@
<ClCompile Include="BaseToolbox\WGS84_J2000.cpp">
<Filter>BaseToolbox</Filter>
</ClCompile>
<ClCompile Include="BaseToolbox\QClipRasterByRowCols.cpp">
<Filter>BaseToolbox</Filter>
</ClCompile>
</ItemGroup>
<ItemGroup>
<ClInclude Include="SimulationSAR\RFPCProcessCls.h">
@ -249,7 +252,6 @@
<ClInclude Include="BaseToolbox\WGS84_J2000.h">
<Filter>BaseToolbox</Filter>
</ClInclude>
<ClInclude Include="GPUTool\GPUGarbage.cuh" />
</ItemGroup>
<ItemGroup>
<QtMoc Include="QMergeRasterProcessDialog.h">
@ -294,6 +296,9 @@
<QtMoc Include="BaseTool\QToolProcessBarDialog.h">
<Filter>BaseTool</Filter>
</QtMoc>
<QtMoc Include="BaseToolbox\QClipRasterByRowCols.h">
<Filter>BaseToolbox</Filter>
</QtMoc>
</ItemGroup>
<ItemGroup>
<QtUic Include="QMergeRasterProcessDialog.ui">
@ -332,6 +337,9 @@
<QtUic Include="BaseToolbox\QRDOrthProcessClass.ui">
<Filter>BaseToolbox</Filter>
</QtUic>
<QtUic Include="BaseToolbox\QClipRasterByRowCols.ui">
<Filter>BaseToolbox</Filter>
</QtUic>
</ItemGroup>
<ItemGroup>
<CudaCompile Include="GPUTool\GPURFPC.cu">
@ -343,7 +351,6 @@
<CudaCompile Include="GPUTool\GPUTool.cu">
<Filter>GPUTool</Filter>
</CudaCompile>
<CudaCompile Include="GPUTool\GPUGarbage.cu" />
</ItemGroup>
<ItemGroup>
<None Include="cpp.hint" />

25
RegisterToolbox.cpp
View File

@ -8,6 +8,7 @@
#include "QImageSARRFPC.h"
#include "QSimulationBPImage.h"
#include "DEMLLA2XYZTool.h"
#include "QClipRasterByRowCols.h"
GF3ImportDataToolButton::GF3ImportDataToolButton(QWidget* parent) :QToolAbstract(parent)
{
@ -139,6 +140,8 @@ void RegisterPreToolBox(RasterProcessTool* mainWindows)
MergeRasterProcessToolButton* items5 = new MergeRasterProcessToolButton(nullptr);
SARSimlulationRFPCToolButton* items6 = new SARSimlulationRFPCToolButton(nullptr);
SARSimulationTBPImageToolButton* items7 = new SARSimulationTBPImageToolButton(nullptr);
DEMLLA2XYZToolButton* items8 = new DEMLLA2XYZToolButton(nullptr);
ClipRasterByRowCols* items9 = new ClipRasterByRowCols(nullptr);
emit mainWindows->addBoxToolItemSIGNAL(items1);
emit mainWindows->addBoxToolItemSIGNAL(items2);
@ -147,6 +150,8 @@ void RegisterPreToolBox(RasterProcessTool* mainWindows)
emit mainWindows->addBoxToolItemSIGNAL(items5);
emit mainWindows->addBoxToolItemSIGNAL(items6);
emit mainWindows->addBoxToolItemSIGNAL(items7);
emit mainWindows->addBoxToolItemSIGNAL(items8);
emit mainWindows->addBoxToolItemSIGNAL(items9);
}
@ -164,4 +169,22 @@ void DEMLLA2XYZToolButton::excute()
{
DEMLLA2XYZTool* dialog = new DEMLLA2XYZTool;
dialog->show();
}
}
ClipRasterByRowCols::ClipRasterByRowCols(QWidget* parent)
{
this->toolPath = QVector<QString>(0);
this->toolPath.push_back(u8"基础处理");
this->toolname = QString(u8"裁剪影像根据行列号");
}
ClipRasterByRowCols::~ClipRasterByRowCols()
{
}
void ClipRasterByRowCols::excute()
{
QClipRasterByRowCols* dialog = new QClipRasterByRowCols;
dialog->show();
}

18
RegisterToolbox.h
View File

@ -87,5 +87,23 @@ public slots:
};
class ClipRasterByRowCols :public QToolAbstract {
Q_OBJECT
public:
ClipRasterByRowCols(QWidget* parent = nullptr);
~ClipRasterByRowCols();
public slots:
virtual void excute() override;
};
void RegisterPreToolBox(RasterProcessTool* mainWindows);

8
SimulationSAR/QImageSARRFPC.ui
View File

@ -25,7 +25,7 @@
<x>0</x>
<y>0</y>
<width>853</width>
<height>634</height>
<height>637</height>
</rect>
</property>
<layout class="QGridLayout" name="gridLayout">
@ -103,7 +103,7 @@
</size>
</property>
<property name="text">
<string>D:/Programme/vs2022/RasterMergeTest/simulationData/landcover_aligned2.dat</string>
<string>D:/Programme/vs2022/RasterMergeTest/simulationData/demdataset/landcover_center_int32.dat</string>
</property>
</widget>
</item>
@ -233,7 +233,7 @@
</size>
</property>
<property name="text">
<string>D:/Programme/vs2022/RasterMergeTest/simulationData/demdataset/demxyz.bin</string>
<string>D:/Programme/vs2022/RasterMergeTest/simulationData/demdataset/demxyz_center.bin</string>
</property>
</widget>
</item>
@ -350,7 +350,7 @@
</size>
</property>
<property name="text">
<string>D:/Programme/vs2022/RasterMergeTest/simulationData/demdataset/demsloper.bin</string>
<string>D:/Programme/vs2022/RasterMergeTest/simulationData/demdataset/demsloper_center.bin</string>
</property>
</widget>
</item>

578
SimulationSAR/RFPCProcessCls.cpp
View File

@ -1,4 +1,4 @@

#include "stdafx.h"
#include "RFPCProcessCls.h"
#include "BaseConstVariable.h"
@ -32,6 +32,192 @@
CUDA_AntSate_PtrList* malloc_AntSate_PtrList(long PRFCount)
{
CUDA_AntSate_PtrList* antlist = (CUDA_AntSate_PtrList*)malloc(sizeof(CUDA_AntSate_PtrList));
antlist->h_antpx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antpy = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antpz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antvx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antvy = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antvz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antdirectx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antdirecty = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antdirectz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antXaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antXaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antXaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antYaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antYaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antYaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antZaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antZaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->h_antZaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antlist->d_antpx = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antpy = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antpz = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antvx = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antvy = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antvz = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antdirectx = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antdirecty = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antdirectz = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antXaxisX = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antXaxisY = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antXaxisZ = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antYaxisX = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antYaxisY = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antYaxisZ = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antZaxisX = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antZaxisY = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->d_antZaxisZ = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
antlist->PRF_len = PRFCount;
return antlist;
}
void Free_AntSate_PtrList(CUDA_AntSate_PtrList* antlist)
{
FreeCUDAHost(antlist->h_antpx);
FreeCUDAHost(antlist->h_antpy);
FreeCUDAHost(antlist->h_antpz);
FreeCUDAHost(antlist->h_antvx);
FreeCUDAHost(antlist->h_antvy);
FreeCUDAHost(antlist->h_antvz);
FreeCUDAHost(antlist->h_antdirectx);
FreeCUDAHost(antlist->h_antdirecty);
FreeCUDAHost(antlist->h_antdirectz);
FreeCUDAHost(antlist->h_antXaxisX);
FreeCUDAHost(antlist->h_antXaxisY);
FreeCUDAHost(antlist->h_antXaxisZ);
FreeCUDAHost(antlist->h_antYaxisX);
FreeCUDAHost(antlist->h_antYaxisY);
FreeCUDAHost(antlist->h_antYaxisZ);
FreeCUDAHost(antlist->h_antZaxisX);
FreeCUDAHost(antlist->h_antZaxisY);
FreeCUDAHost(antlist->h_antZaxisZ);
FreeCUDADevice(antlist->d_antpx);
FreeCUDADevice(antlist->d_antpy);
FreeCUDADevice(antlist->d_antpz);
FreeCUDADevice(antlist->d_antvx);
FreeCUDADevice(antlist->d_antvy);
FreeCUDADevice(antlist->d_antvz);
FreeCUDADevice(antlist->d_antdirectx);
FreeCUDADevice(antlist->d_antdirecty);
FreeCUDADevice(antlist->d_antdirectz);
FreeCUDADevice(antlist->d_antXaxisX);
FreeCUDADevice(antlist->d_antXaxisY);
FreeCUDADevice(antlist->d_antXaxisZ);
FreeCUDADevice(antlist->d_antYaxisX);
FreeCUDADevice(antlist->d_antYaxisY);
FreeCUDADevice(antlist->d_antYaxisZ);
FreeCUDADevice(antlist->d_antZaxisX);
FreeCUDADevice(antlist->d_antZaxisY);
FreeCUDADevice(antlist->d_antZaxisZ);
antlist->h_antpx = nullptr;
antlist->h_antpy = nullptr;
antlist->h_antpz = nullptr;
antlist->h_antvx = nullptr;
antlist->h_antvy = nullptr;
antlist->h_antvz = nullptr;
antlist->h_antdirectx = nullptr;
antlist->h_antdirecty = nullptr;
antlist->h_antdirectz = nullptr;
antlist->h_antXaxisX = nullptr;
antlist->h_antXaxisY = nullptr;
antlist->h_antXaxisZ = nullptr;
antlist->h_antYaxisX = nullptr;
antlist->h_antYaxisY = nullptr;
antlist->h_antYaxisZ = nullptr;
antlist->h_antZaxisX = nullptr;
antlist->h_antZaxisY = nullptr;
antlist->h_antZaxisZ = nullptr;
antlist->d_antpx = nullptr;
antlist->d_antpy = nullptr;
antlist->d_antpz = nullptr;
antlist->d_antvx = nullptr;
antlist->d_antvy = nullptr;
antlist->d_antvz = nullptr;
antlist->d_antdirectx = nullptr;
antlist->d_antdirecty = nullptr;
antlist->d_antdirectz = nullptr;
antlist->d_antXaxisX = nullptr;
antlist->d_antXaxisY = nullptr;
antlist->d_antXaxisZ = nullptr;
antlist->d_antYaxisX = nullptr;
antlist->d_antYaxisY = nullptr;
antlist->d_antYaxisZ = nullptr;
antlist->d_antZaxisX = nullptr;
antlist->d_antZaxisY = nullptr;
antlist->d_antZaxisZ = nullptr;
free(antlist);
antlist = nullptr;
}
void COPY_AntStation_FROM_HOST_GPU(std::shared_ptr<SatelliteOribtNode[]> sateOirbtNodes,
std::shared_ptr<CUDA_AntSate_PtrList> gpupptr,
long startPID,
long PRF_len)
{
assert(gpupptr->PRF_len <= PRF_len);
long prfid = 0;
for (long tempprfid = 0; tempprfid < PRF_len; tempprfid++) {
prfid = tempprfid + startPID;
gpupptr->h_antpx[tempprfid] = sateOirbtNodes[prfid].Px;
gpupptr->h_antpy[tempprfid] = sateOirbtNodes[prfid].Py;
gpupptr->h_antpz[tempprfid] = sateOirbtNodes[prfid].Pz;
gpupptr->h_antvx[tempprfid] = sateOirbtNodes[prfid].Vx;
gpupptr->h_antvy[tempprfid] = sateOirbtNodes[prfid].Vy;
gpupptr->h_antvz[tempprfid] = sateOirbtNodes[prfid].Vz; //6
gpupptr->h_antdirectx[tempprfid] = sateOirbtNodes[prfid].AntDirecX;
gpupptr->h_antdirecty[tempprfid] = sateOirbtNodes[prfid].AntDirecY;
gpupptr->h_antdirectz[tempprfid] = sateOirbtNodes[prfid].AntDirecZ;
gpupptr->h_antXaxisX[tempprfid] = sateOirbtNodes[prfid].AntXaxisX;
gpupptr->h_antXaxisY[tempprfid] = sateOirbtNodes[prfid].AntXaxisY;
gpupptr->h_antXaxisZ[tempprfid] = sateOirbtNodes[prfid].AntXaxisZ;//12
gpupptr->h_antYaxisX[tempprfid] = sateOirbtNodes[prfid].AntYaxisX;
gpupptr->h_antYaxisY[tempprfid] = sateOirbtNodes[prfid].AntYaxisY;
gpupptr->h_antYaxisZ[tempprfid] = sateOirbtNodes[prfid].AntYaxisZ;//15
gpupptr->h_antZaxisX[tempprfid] = sateOirbtNodes[prfid].AntZaxisX;
gpupptr->h_antZaxisY[tempprfid] = sateOirbtNodes[prfid].AntZaxisY;
gpupptr->h_antZaxisZ[tempprfid] = sateOirbtNodes[prfid].AntZaxisZ;//18
}
HostToDevice(gpupptr->h_antpx, gpupptr->d_antpx, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antpy, gpupptr->d_antpy, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antpz, gpupptr->d_antpz, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antvx, gpupptr->d_antvx, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antvy, gpupptr->d_antvy, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antvz, gpupptr->d_antvz, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antdirectx, gpupptr->d_antdirectx, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antdirecty, gpupptr->d_antdirecty, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antdirectz, gpupptr->d_antdirectz, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antXaxisX, gpupptr->d_antXaxisX, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antXaxisY, gpupptr->d_antXaxisY, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antXaxisZ, gpupptr->d_antXaxisZ, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antYaxisX, gpupptr->d_antYaxisX, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antYaxisY, gpupptr->d_antYaxisY, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antYaxisZ, gpupptr->d_antYaxisZ, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antZaxisX, gpupptr->d_antZaxisX, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antZaxisY, gpupptr->d_antZaxisY, sizeof(double) * PRF_len);
HostToDevice(gpupptr->h_antZaxisZ, gpupptr->d_antZaxisZ, sizeof(double) * PRF_len);
}
RFPCProcessCls::RFPCProcessCls()
{
@ -95,7 +281,7 @@ void RFPCProcessCls::setOutEchoPath(QString OutEchoPath)
ErrorCode RFPCProcessCls::Process(long num_thread)
{
// RFPC 算法
// RFPC 算法
qDebug() << u8"params init ....";
ErrorCode stateCode = this->InitParams();
if (stateCode != ErrorCode::SUCCESS) {
@ -113,7 +299,7 @@ ErrorCode RFPCProcessCls::Process(long num_thread)
qDebug() << "InitEchoMaskArray";
//stateCode = this->RFPCMainProcess(num_thread);
// 初始化回波
// 初始化回波
this->EchoSimulationData->initEchoArr(std::complex<double>(0, 0));
stateCode = this->RFPCMainProcess_GPU();
@ -136,17 +322,17 @@ ErrorCode RFPCProcessCls::InitParams()
}
// 归一化绝对路径
// 归一化绝对路径
this->OutEchoPath = QDir(this->OutEchoPath).absolutePath();
// 回波大小
// 回波大小
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());
// 初始化回波存放位置
// 初始化回波存放位置
qDebug() << "--------------Echo Data Setting ---------------------------------------";
this->EchoSimulationData = std::shared_ptr<EchoL0Dataset>(new EchoL0Dataset);
this->EchoSimulationData->setCenterFreq(this->TaskSetting->getCenterFreq());
@ -200,9 +386,9 @@ std::shared_ptr<SatelliteOribtNode[]> RFPCProcessCls::getSatelliteOribtNodes(dou
std::shared_ptr<SatelliteOribtNode[]> sateOirbtNodes(new SatelliteOribtNode[this->PluseCount], delArrPtr);
{ // 姿态计算不同
{ // 姿态计算不同
qDebug() << "Ant position finished started !!!";
// 计算姿态
// 计算姿态
std::shared_ptr<double> antpos = this->EchoSimulationData->getAntPos();
double dAt = 1e-6;
double prf_time_dt = 0;
@ -215,7 +401,7 @@ std::shared_ptr<SatelliteOribtNode[]> RFPCProcessCls::getSatelliteOribtNodes(dou
this->TaskSetting->getSatelliteOribtNode(prf_time, sateOirbtNode, antflag);
this->TaskSetting->getSatelliteOribtNode(prf_time_dt, sateOirbtNode_dAt, antflag);
sateOirbtNode.AVx = (sateOirbtNode_dAt.Vx - sateOirbtNode.Vx) / dAt; // 加速度
sateOirbtNode.AVx = (sateOirbtNode_dAt.Vx - sateOirbtNode.Vx) / dAt; // 加速度
sateOirbtNode.AVy = (sateOirbtNode_dAt.Vy - sateOirbtNode.Vy) / dAt;
sateOirbtNode.AVz = (sateOirbtNode_dAt.Vz - sateOirbtNode.Vz) / dAt;
@ -268,7 +454,7 @@ void RFPCProcessMain(long num_thread,
return;
}
else {
// 打印参数
// 打印参数
qDebug() << "--------------Task Seting ---------------------------------------";
qDebug() << "SARImageStartTime: " << task->getSARImageStartTime();
qDebug() << "SARImageEndTime: " << task->getSARImageEndTime();
@ -283,7 +469,7 @@ void RFPCProcessMain(long num_thread,
qDebug() << (task->getFarRange() - task->getNearRange()) * 2 / LIGHTSPEED * task->getFs();
qDebug() << "\n\n";
}
// 1.2 设置天线方向图
// 1.2 设置天线方向图
std::vector<RadiationPatternGainPoint> TansformPatternGainpoints = ReadGainFile(TansformPatternFilePath);
std::shared_ptr<AbstractRadiationPattern> TansformPatternGainPtr = CreateAbstractRadiationPattern(TansformPatternGainpoints);
@ -293,7 +479,7 @@ void RFPCProcessMain(long num_thread,
task->setTransformRadiationPattern(TansformPatternGainPtr);
task->setReceiveRadiationPattern(ReceivePatternGainPtr);
//2. 读取GPS节点
//2. 读取GPS节点
std::vector<SatelliteOribtNode> nodes;
ErrorCode stateCode = ReadSateGPSPointsXML(GPSXmlPath, nodes);
@ -304,8 +490,8 @@ void RFPCProcessMain(long num_thread,
}
else {}
std::shared_ptr<AbstractSatelliteOribtModel> SatelliteOribtModel = CreataPolyfitSatelliteOribtModel(nodes, task->getSARImageStartTime(), 3); // 以成像开始时间作为 时间参考起点
SatelliteOribtModel->setbeamAngle(task->getCenterLookAngle(), task->getIsRightLook()); // 设置天线方向图
std::shared_ptr<AbstractSatelliteOribtModel> SatelliteOribtModel = CreataPolyfitSatelliteOribtModel(nodes, task->getSARImageStartTime(), 3); // 以成像开始时间作为 时间参考起点
SatelliteOribtModel->setbeamAngle(task->getCenterLookAngle(), task->getIsRightLook()); // 设置天线方向图
if (nullptr == SatelliteOribtModel)
{
@ -324,156 +510,61 @@ void RFPCProcessMain(long num_thread,
RFPC.setLandCoverPath(LandCoverPath); //qDebug() << "setLandCoverPath";
RFPC.setOutEchoPath(OutEchoPath); //qDebug() << "setOutEchoPath";
qDebug() << "-------------- RFPC start---------------------------------------";
RFPC.Process(num_thread); // 处理程序
RFPC.Process(num_thread); // 处理程序
qDebug() << "-------------- RFPC end---------------------------------------";
}
ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
/** 内存分配***************************************************/
/** 内存分配***************************************************/
long TargetMemoryMB = 500;
/** 参数区域***************************************************/
/** 参数区域***************************************************/
QVector<double> freqlist = this->TaskSetting->getFreqList();
long freqnum = freqlist.count();
float f0 = float(freqlist[0] / 1e9);
float dfreq = float((freqlist[1] - freqlist[0]) / 1e9);
#if (defined __PRFDEBUG__) && (defined __PRFDEBUG_PRFINF__)
double* h_freqPtr = (double*)mallocCUDAHost(sizeof(double) * freqnum);
for (long fid = 0; fid < freqnum; fid++) {
h_freqPtr[fid] = (f0 + dfreq * fid) * 1e9;
}
testOutAmpArr("freqlist.bin", h_freqPtr, freqnum, 1);
#endif
long PRFCount = this->EchoSimulationData->getPluseCount();
double NearRange = this->EchoSimulationData->getNearRange(); // 近斜距
double NearRange = this->EchoSimulationData->getNearRange(); // 近斜距
double FarRange = this->EchoSimulationData->getFarRange();
double Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 发射电压 1v
double Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 发射电压 1v
double lamda = this->TaskSetting->getCenterLamda(); // 波长
double refphaseRange = this->TaskSetting->getRefphaseRange(); // 参考相位斜距
double lamda = this->TaskSetting->getCenterLamda(); // 波长
double refphaseRange = this->TaskSetting->getRefphaseRange(); // 参考相位斜距
double prf_time = 0;
double dt = 1 / this->TaskSetting->getPRF();// 获取每次脉冲的时间间隔
bool antflag = true; // 计算天线方向图
double dt = 1 / this->TaskSetting->getPRF();// 获取每次脉冲的时间间隔
bool antflag = true; // 计算天线方向图
long double imageStarttime = this->TaskSetting->getSARImageStartTime();
// 卫星
double* h_antpx, * d_antpx;
double* h_antpy, * d_antpy;
double* h_antpz, * d_antpz;
double* h_antvx, * d_antvx;
double* h_antvy, * d_antvy;
double* h_antvz, * d_antvz;
double* h_antdirectx, * d_antdirectx;
double* h_antdirecty, * d_antdirecty;
double* h_antdirectz, * d_antdirectz;
double* h_antXaxisX, * d_antXaxisX;
double* h_antXaxisY, * d_antXaxisY;
double* h_antXaxisZ, * d_antXaxisZ;
double* h_antYaxisX, * d_antYaxisX;
double* h_antYaxisY, * d_antYaxisY;
double* h_antYaxisZ, * d_antYaxisZ;
double* h_antZaxisX, * d_antZaxisX;
double* h_antZaxisY, * d_antZaxisY;
double* h_antZaxisZ, * d_antZaxisZ;
{
h_antpx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antpy = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antpz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antvx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antvy = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antvz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antdirectx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antdirecty = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antdirectz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antXaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antXaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antXaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antYaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antYaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antYaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antZaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antZaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
h_antZaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
this->EchoSimulationData->getAntPos();
std::shared_ptr<SatelliteOribtNode[]> sateOirbtNodes = this->getSatelliteOribtNodes(prf_time, dt, antflag, imageStarttime);
/** 天线方向图***************************************************/
std::shared_ptr<AbstractRadiationPattern> TransformPattern = this->TaskSetting->getTransformRadiationPattern(); // 发射天线方向图
d_antpx = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antpy = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antpz = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antvx = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antvy = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antvz = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antdirectx = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antdirecty = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antdirectz = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antXaxisX = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antXaxisY = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antXaxisZ = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antYaxisX = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antYaxisY = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antYaxisZ = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antZaxisX = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antZaxisY = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
d_antZaxisZ = (double*)mallocCUDADevice(sizeof(double) * PRFCount);
this->EchoSimulationData->getAntPos();
std::shared_ptr<SatelliteOribtNode[]> sateOirbtNodes = this->getSatelliteOribtNodes(prf_time, dt, antflag, imageStarttime);
for (long tempprfid = 0; tempprfid < PRFCount; tempprfid++) {
long prfid = tempprfid;
h_antpx[tempprfid] = sateOirbtNodes[prfid].Px;
h_antpy[tempprfid] = sateOirbtNodes[prfid].Py;
h_antpz[tempprfid] = sateOirbtNodes[prfid].Pz;
h_antvx[tempprfid] = sateOirbtNodes[prfid].Vx;
h_antvy[tempprfid] = sateOirbtNodes[prfid].Vy;
h_antvz[tempprfid] = sateOirbtNodes[prfid].Vz; //6
h_antdirectx[tempprfid] = sateOirbtNodes[prfid].AntDirecX;
h_antdirecty[tempprfid] = sateOirbtNodes[prfid].AntDirecY;
h_antdirectz[tempprfid] = sateOirbtNodes[prfid].AntDirecZ; // 9 天线指向
h_antXaxisX[tempprfid] = sateOirbtNodes[prfid].AntXaxisX;
h_antXaxisY[tempprfid] = sateOirbtNodes[prfid].AntXaxisY;
h_antXaxisZ[tempprfid] = sateOirbtNodes[prfid].AntXaxisZ;//12 天线坐标系
h_antYaxisX[tempprfid] = sateOirbtNodes[prfid].AntYaxisX;
h_antYaxisY[tempprfid] = sateOirbtNodes[prfid].AntYaxisY;
h_antYaxisZ[tempprfid] = sateOirbtNodes[prfid].AntYaxisZ;//15
h_antZaxisX[tempprfid] = sateOirbtNodes[prfid].AntZaxisX;
h_antZaxisY[tempprfid] = sateOirbtNodes[prfid].AntZaxisY;
h_antZaxisZ[tempprfid] = sateOirbtNodes[prfid].AntZaxisZ;//18
}
DeviceToDevice(h_antpx, d_antpx, sizeof(double) * PRFCount);
DeviceToDevice(h_antpy, d_antpy, sizeof(double) * PRFCount);
DeviceToDevice(h_antpz, d_antpz, sizeof(double) * PRFCount);
DeviceToDevice(h_antvx, d_antvx, sizeof(double) * PRFCount);
DeviceToDevice(h_antvy, d_antvy, sizeof(double) * PRFCount);
DeviceToDevice(h_antvz, d_antvz, sizeof(double) * PRFCount);
DeviceToDevice(h_antdirectx, d_antdirectx, sizeof(double) * PRFCount);
DeviceToDevice(h_antdirecty, d_antdirecty, sizeof(double) * PRFCount);
DeviceToDevice(h_antdirectz, d_antdirectz, sizeof(double) * PRFCount);
DeviceToDevice(h_antXaxisX, d_antXaxisX, sizeof(double) * PRFCount);
DeviceToDevice(h_antXaxisY, d_antXaxisY, sizeof(double) * PRFCount);
DeviceToDevice(h_antXaxisZ, d_antXaxisZ, sizeof(double) * PRFCount);
DeviceToDevice(h_antYaxisX, d_antYaxisX, sizeof(double) * PRFCount);
DeviceToDevice(h_antYaxisY, d_antYaxisY, sizeof(double) * PRFCount);
DeviceToDevice(h_antYaxisZ, d_antYaxisZ, sizeof(double) * PRFCount);
DeviceToDevice(h_antZaxisX, d_antZaxisX, sizeof(double) * PRFCount);
DeviceToDevice(h_antZaxisY, d_antZaxisY, sizeof(double) * PRFCount);
DeviceToDevice(h_antZaxisZ, d_antZaxisZ, sizeof(double) * PRFCount);
}
/** 天线方向图***************************************************/
std::shared_ptr<AbstractRadiationPattern> TransformPattern = this->TaskSetting->getTransformRadiationPattern(); // 发射天线方向图
std::shared_ptr<AbstractRadiationPattern> ReceivePattern = this->TaskSetting->getReceiveRadiationPattern(); // 接收天线方向图
std::shared_ptr<AbstractRadiationPattern> ReceivePattern = this->TaskSetting->getReceiveRadiationPattern(); // 接收天线方向图
POLARTYPEENUM polartype = this->TaskSetting->getPolarType();
PatternImageDesc TantPatternDesc = {};
double* h_TantPattern = nullptr;
double* d_TantPattern = nullptr;
double maxTransAntPatternValue = 0;
{
// 处理发射天线方向图
// 处理发射天线方向图
double Tminphi = TransformPattern->getMinPhi();
double Tmaxphi = TransformPattern->getMaxPhi();
double Tmintheta = TransformPattern->getMinTheta();
@ -499,9 +590,13 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
}
testOutAntPatternTrans("TransPattern.bin", h_TantPattern, TstartTheta, Tdtheta, TstartPhi, Tdphi, Tthetanum, Tphinum);
maxTransAntPatternValue = powf(10.0, h_TantPattern[0] / 10);
for (long i = 0; i < Tthetanum; i++) {
for (long j = 0; j < Tphinum; j++) {
h_TantPattern[i * Tphinum + j] = powf(10.0, h_TantPattern[i * Tphinum + j] / 10);
h_TantPattern[i * Tphinum + j] = powf(10.0, h_TantPattern[i * Tphinum + j] / 10); // 转换为线性值
if (maxTransAntPatternValue < h_TantPattern[i * Tphinum + j]) {
maxTransAntPatternValue = h_TantPattern[i * Tphinum + j];
}
}
}
HostToDevice(h_TantPattern, d_TantPattern, sizeof(double) * Tthetanum * Tphinum);
@ -516,8 +611,9 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
PatternImageDesc RantPatternDesc = {};
double* h_RantPattern = nullptr;
double* d_RantPattern = nullptr;
double maxReceiveAntPatternValue = 0;
{
// 处理接收天线方向图
// 处理接收天线方向图
double Rminphi = ReceivePattern->getMinPhi();
double Rmaxphi = ReceivePattern->getMaxPhi();
double Rmintheta = ReceivePattern->getMinTheta();
@ -534,18 +630,23 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
h_RantPattern = (double*)mallocCUDAHost(sizeof(double) * Rthetanum * Rphinum);
d_RantPattern = (double*)mallocCUDADevice(sizeof(double) * Rthetanum * Rphinum);
for (long i = 0; i < Rthetanum; i++) {
for (long j = 0; j < Rphinum; j++) {
//h_RantPattern[i * Rphinum + j] = ReceivePattern->getGainLearThetaPhi(RstartTheta + i * Rdtheta, RstartPhi + j * Rdphi);
h_RantPattern[i * Rphinum + j] = ReceivePattern->getGain(RstartTheta + i * Rdtheta, RstartPhi + j * Rdphi);
}
}
testOutAntPatternTrans("ReceivePattern.bin", h_RantPattern, Rmintheta, Rdtheta, RstartPhi, Rdphi, Rthetanum, Rphinum);
maxReceiveAntPatternValue = powf(10.0, h_RantPattern[0] / 10);
for (long i = 0; i < Rthetanum; i++) {
for (long j = 0; j < Rphinum; j++) {
h_RantPattern[i * Rphinum + j] = powf(10.0, h_RantPattern[i * Rphinum + j] / 10);
if (maxReceiveAntPatternValue < h_RantPattern[i * Rphinum + j]) {
maxReceiveAntPatternValue = h_RantPattern[i * Rphinum + j];
}
}
}
HostToDevice(h_RantPattern, d_RantPattern, sizeof(double) * Rthetanum * Rphinum);
@ -559,22 +660,22 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
/** 坐标区域点***************************************************/
gdalImage demxyz(this->demxyzPath);// 地面点坐标
gdalImage demlandcls(this->LandCoverPath);// 地表覆盖类型
gdalImage demsloperxyz(this->demsloperPath);// 地面坡向
/** 坐标区域点***************************************************/
gdalImage demxyz(this->demxyzPath);// 地面点坐标
gdalImage demlandcls(this->LandCoverPath);// 地表覆盖类型
gdalImage demsloperxyz(this->demsloperPath);// 地面坡向
long demRow = demxyz.height;
long demCol = demxyz.width;
//处理地表覆盖
//处理地表覆盖
QMap<long, long> clamap;
long clamapid = 0;
long startline = 0;
{
long blokline = getBlockRows(2e4, demCol, sizeof(double));
long blokline = getBlockRows(2e4, demCol, sizeof(double),demRow);
for (startline = 0; startline < demRow; startline = startline + blokline) {
Eigen::MatrixXd clsland = demlandcls.getData(startline, 0, blokline, demlandcls.width, 1);
long clsrows = clsland.rows();
@ -613,7 +714,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
h_clsSigmaParam[clamap[id]].p6 = tempp.p6;
}
// 打印日志
// 打印日志
std::cout << "sigma params:" << std::endl;
std::cout << "classid:\tp1\tp2\tp3\tp4\tp5\tp6" << std::endl;
for (long ii = 0; ii < clamapid; ii++) {
@ -628,30 +729,58 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
}
HostToDevice(h_clsSigmaParam, d_clsSigmaParam, sizeof(CUDASigmaParam) * clamapid);
qDebug() << "CUDA class Proces finished!!!";
// 处理地面坐标
long blockline = getBlockRows(TargetMemoryMB, demCol, sizeof(double));
double* h_dem_x = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_dem_y = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_dem_z = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_demsloper_x = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_demsloper_y = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_demsloper_z = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
long* h_demcls = (long*)mallocCUDAHost(sizeof(long) * blockline * demCol);
// 处理地面坐标
long blockline = getBlockRows(TargetMemoryMB, demCol, sizeof(double), demRow);
double* h_dem_x = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_dem_y = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_dem_z = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_demsloper_x = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_demsloper_y = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
double* h_demsloper_z = (double*)mallocCUDAHost(sizeof(double) * blockline * demCol);
long* h_demcls = (long*)mallocCUDAHost(sizeof(long) * blockline * demCol);
/** 处理回波***************************************************/
long echo_block_rows = getBlockRows(5000, freqnum, sizeof(float)*2);
echo_block_rows = echo_block_rows < PRFCount ? echo_block_rows : PRFCount;
float* h_echo_block_real = (float*)mallocCUDAHost(sizeof(float) * echo_block_rows * freqnum);
double* d_dem_x = (double*)mallocCUDADevice(sizeof(double) * blockline * demCol);
double* d_dem_y = (double*)mallocCUDADevice(sizeof(double) * blockline * demCol);
double* d_dem_z = (double*)mallocCUDADevice(sizeof(double) * blockline * demCol);
double* d_demsloper_x = (double*)mallocCUDADevice(sizeof(double) * blockline * demCol);
double* d_demsloper_y = (double*)mallocCUDADevice(sizeof(double) * blockline * demCol);
double* d_demsloper_z = (double*)mallocCUDADevice(sizeof(double) * blockline * demCol);
long* d_demcls = (long*) mallocCUDADevice(sizeof(long) * blockline * demCol);
/** 处理回波***************************************************/
long echo_block_rows = getBlockRows(1000, freqnum, sizeof(float)*2, PRFCount);
float* h_echo_block_real = (float*)mallocCUDAHost(sizeof(float) * echo_block_rows * freqnum);
float* h_echo_block_imag = (float*)mallocCUDAHost(sizeof(float) * echo_block_rows * freqnum);
float* d_echo_block_real = (float*)mallocCUDADevice(sizeof(float) * echo_block_rows * freqnum);
float* d_echo_block_imag = (float*)mallocCUDADevice(sizeof(float) * echo_block_rows * freqnum);
/** 主流程处理 ***************************************************/
float* d_temp_R = (float*)mallocCUDADevice(sizeof(float) * echo_block_rows * SHAREMEMORY_FLOAT_HALF); //2GB 距离
float* d_temp_amp = (float*)mallocCUDADevice(sizeof(float) * echo_block_rows * SHAREMEMORY_FLOAT_HALF);//2GB 强度
/** 主流程处理 ***************************************************/
qDebug() << "CUDA Main Proces";
for (long sprfid = 0; sprfid < PRFCount; sprfid = sprfid + echo_block_rows) {
long PRF_len = (sprfid + echo_block_rows) < PRFCount ? echo_block_rows : (PRFCount - sprfid);
qDebug() << "Start PRF: " << sprfid << "\t-\t" << sprfid + PRF_len << "\t:copy ant list host -> GPU";
std::shared_ptr< CUDA_AntSate_PtrList> antptrlist(malloc_AntSate_PtrList(PRF_len), Free_AntSate_PtrList);
COPY_AntStation_FROM_HOST_GPU(sateOirbtNodes, antptrlist, sprfid, PRF_len);
qDebug() << "Start PRF: " << sprfid << "\t-\t" << sprfid + PRF_len << "\t:copy echo data list host -> GPU";
std::shared_ptr<std::complex<double>> echo_temp = this->EchoSimulationData->getEchoArr(sprfid, PRF_len);
for (long ii = 0; ii < PRF_len; ii++) {
for (long jj = 0; jj < freqnum; jj++) {
@ -659,9 +788,12 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
h_echo_block_imag[ii * freqnum + jj]=echo_temp.get()[ii * freqnum + jj].imag();
}
}
HostToDevice(h_echo_block_real, d_echo_block_real, sizeof(float) * PRF_len* freqnum);
HostToDevice(h_echo_block_imag, d_echo_block_imag, sizeof(float) * PRF_len* freqnum);
for (long startline = 0; startline < demRow; startline = startline + blockline) {
Eigen::MatrixXd dem_x = demxyz.getData(startline, 0, blockline, demCol, 1); // 地面坐标
for (startline = 0; startline < demRow; startline = startline + blockline) {
Eigen::MatrixXd dem_x = demxyz.getData(startline, 0, blockline, demCol, 1); // 地面坐标
Eigen::MatrixXd dem_y = demxyz.getData(startline, 0, blockline, demCol, 2);
Eigen::MatrixXd dem_z = demxyz.getData(startline, 0, blockline, demCol, 3);
Eigen::MatrixXd demsloper_x = demsloperxyz.getData(startline, 0, blockline, demCol, 1);
@ -675,7 +807,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
// 更新数据格式
// 更新数据格式
for (long i = 0; i < temp_dem_row; i++) {
for (long j = 0; j < temp_dem_col; j++) {
h_dem_x[i * temp_dem_col + j] = double(dem_x(i, j));
@ -684,39 +816,70 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
h_demsloper_x[i * temp_dem_col + j] = double(demsloper_x(i, j));
h_demsloper_y[i * temp_dem_col + j] = double(demsloper_y(i, j));
h_demsloper_z[i * temp_dem_col + j] = double(demsloper_z(i, j));
h_demcls[i * temp_dem_col + j] = clamap[long(landcover(i, j))];
}
}
qDebug() << "Start PRF: " << sprfid << "\t-\t" << sprfid + PRF_len << "\t:copy target data ("<< startline<<" - "<< startline + blockline << ") host -> GPU";
HostToDevice(h_dem_x, d_dem_x , sizeof(double) * blockline * demCol);
HostToDevice(h_dem_y, d_dem_y , sizeof(double) * blockline * demCol);
HostToDevice(h_dem_z, d_dem_z , sizeof(double) * blockline * demCol);
HostToDevice(h_demsloper_x, d_demsloper_x , sizeof(double) * blockline * demCol);
HostToDevice(h_demsloper_y, d_demsloper_y , sizeof(double) * blockline * demCol);
HostToDevice(h_demsloper_z, d_demsloper_z , sizeof(double) * blockline * demCol);
HostToDevice(h_demcls, d_demcls ,sizeof(long)* blockline* demCol);
// 分块处理
// 分块处理
qDebug() << "Start PRF: " << sprfid << "\t-\t" << sprfid + PRF_len << "\t:GPU Computer target data (" << startline << "-" << startline + blockline << ")";
CUDA_RFPC_MainProcess(
d_antpx, d_antpy, d_antpz,
d_antXaxisX, d_antXaxisY, d_antXaxisZ, // 天线坐标系的X轴
d_antYaxisX, d_antYaxisY, d_antYaxisZ,// 天线坐标系的Y轴
d_antZaxisX, d_antZaxisY, d_antZaxisZ,// 天线坐标系的Z轴
d_antdirectx, d_antdirecty, d_antdirectz,// 天线的指向
PRFCount, freqnum,
antptrlist->d_antpx, antptrlist->d_antpy, antptrlist->d_antpz,
antptrlist->d_antXaxisX, antptrlist->d_antXaxisY, antptrlist->d_antXaxisZ, // 天线坐标系的X轴
antptrlist->d_antYaxisX, antptrlist->d_antYaxisY, antptrlist->d_antYaxisZ,// 天线坐标系的Y轴
antptrlist->d_antZaxisX, antptrlist->d_antZaxisY, antptrlist->d_antZaxisZ,// 天线坐标系的Z轴
antptrlist->d_antdirectx, antptrlist->d_antdirecty, antptrlist->d_antdirectz,// 天线的指向
PRF_len, freqnum,
f0,dfreq,
Pt,
refphaseRange,
// 天线方向图
// 天线方向图
d_TantPattern,
TantPatternDesc.startTheta, TantPatternDesc.startPhi, TantPatternDesc.dtheta, TantPatternDesc.dphi, TantPatternDesc.thetanum, TantPatternDesc.phinum,
d_RantPattern,
RantPatternDesc.startTheta, RantPatternDesc.startPhi, RantPatternDesc.dtheta, RantPatternDesc.dphi, RantPatternDesc.thetanum, RantPatternDesc.phinum,
NearRange, FarRange, // 近斜据
h_dem_x, h_dem_y, h_dem_z, h_demcls, temp_dem_count, // 地面坐标
h_demsloper_x, h_demsloper_y, h_demsloper_z, // 地表坡度矢量
maxTransAntPatternValue, maxReceiveAntPatternValue,
NearRange, FarRange, // 近斜据
d_dem_x, d_dem_y, d_dem_z, d_demcls, temp_dem_count, // 地面坐标
d_demsloper_x, d_demsloper_y, d_demsloper_z, // 地表坡度矢量
d_clsSigmaParam, clamapid,
h_echo_block_real, h_echo_block_imag// 输出回波
d_echo_block_real, d_echo_block_imag,// 输出回波
d_temp_R, d_temp_amp
);
PRINT("dem : %d - %d / %d , echo: %d -%d / %d \n", startline, startline+ temp_dem_row, demRow, sprfid, sprfid+ PRF_len, PRFCount);
PRINT("dem : %d ~ %d / %d , echo: %d ~ %d / %d \n", startline, startline+ temp_dem_row, demRow, sprfid, sprfid+ PRF_len, PRFCount);
}
#if (defined __PRFDEBUG__) && (defined __PRFDEBUG_PRFINF__)
float* h_temp_R = (float*)mallocCUDAHost(sizeof(float) * echo_block_rows * SHAREMEMORY_FLOAT_HALF); //2GB 距离
float* h_temp_amp = (float*)mallocCUDAHost(sizeof(float) * echo_block_rows * SHAREMEMORY_FLOAT_HALF);//2GB 强度
DeviceToHost(h_temp_R, d_temp_R, sizeof(float) * echo_block_rows * SHAREMEMORY_FLOAT_HALF);
DeviceToHost(h_temp_amp, d_temp_amp, sizeof(float) * echo_block_rows * SHAREMEMORY_FLOAT_HALF);
testOutAmpArr("temp_R.bin", h_temp_R, echo_block_rows, SHAREMEMORY_FLOAT_HALF);
testOutAmpArr("temp_Amp.bin", h_temp_amp, echo_block_rows, SHAREMEMORY_FLOAT_HALF);
FreeCUDAHost(h_temp_R);
FreeCUDAHost(h_temp_amp);
#endif
DeviceToHost(h_echo_block_real, d_echo_block_real, sizeof(float) * PRF_len * freqnum);
DeviceToHost(h_echo_block_imag, d_echo_block_imag, sizeof(float) * PRF_len * freqnum);
for (long ii = 0; ii < PRF_len; ii++) {
for (long jj = 0; jj < freqnum; jj++) {
@ -725,13 +888,15 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
}
}
this->EchoSimulationData->saveEchoArr(echo_temp, sprfid, PRF_len);
}
/** 内存释放***************************************************/
/** 内存释放***************************************************/
FreeCUDAHost(h_TantPattern);
FreeCUDAHost(h_RantPattern);
FreeCUDADevice(d_TantPattern);
@ -747,45 +912,24 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
FreeCUDAHost(h_echo_block_real);
FreeCUDAHost(h_echo_block_imag);
FreeCUDAHost(h_antpx);
FreeCUDAHost(h_antpy);
FreeCUDAHost(h_antpz);
FreeCUDAHost(h_antvx);
FreeCUDAHost(h_antvy);
FreeCUDAHost(h_antvz);
FreeCUDAHost(h_antdirectx);
FreeCUDAHost(h_antdirecty);
FreeCUDAHost(h_antdirectz);
FreeCUDAHost(h_antXaxisX);
FreeCUDAHost(h_antXaxisY);
FreeCUDAHost(h_antXaxisZ);
FreeCUDAHost(h_antYaxisX);
FreeCUDAHost(h_antYaxisY);
FreeCUDAHost(h_antYaxisZ);
FreeCUDAHost(h_antZaxisX);
FreeCUDAHost(h_antZaxisY);
FreeCUDAHost(h_antZaxisZ);
FreeCUDADevice(d_antpx);
FreeCUDADevice(d_antpy);
FreeCUDADevice(d_antpz);
FreeCUDADevice(d_antvx);
FreeCUDADevice(d_antvy);
FreeCUDADevice(d_antvz);
FreeCUDADevice(d_antdirectx);
FreeCUDADevice(d_antdirecty);
FreeCUDADevice(d_antdirectz);
FreeCUDADevice(d_antXaxisX);
FreeCUDADevice(d_antXaxisY);
FreeCUDADevice(d_antXaxisZ);
FreeCUDADevice(d_antYaxisX);
FreeCUDADevice(d_antYaxisY);
FreeCUDADevice(d_antYaxisZ);
FreeCUDADevice(d_antZaxisX);
FreeCUDADevice(d_antZaxisY);
FreeCUDADevice(d_antZaxisZ);
FreeCUDADevice(d_dem_x);
FreeCUDADevice(d_dem_y);
FreeCUDADevice(d_dem_z);
FreeCUDADevice(d_demsloper_x);
FreeCUDADevice(d_demsloper_y);
FreeCUDADevice(d_demsloper_z);
FreeCUDADevice(d_demcls);
FreeCUDADevice(d_echo_block_real);
FreeCUDADevice(d_echo_block_imag);
FreeCUDADevice(d_temp_R);
FreeCUDADevice(d_temp_amp);
return ErrorCode::SUCCESS;
}

20
SimulationSAR/RFPCProcessCls.h
View File

@ -28,6 +28,23 @@
#include "EchoDataFormat.h"
#include "SigmaDatabase.h"
/***** ¹¤¾ßº¯Êý *******************************/
CUDA_AntSate_PtrList* malloc_AntSate_PtrList(long PRFCount);
void Free_AntSate_PtrList(CUDA_AntSate_PtrList* ant);
void COPY_AntStation_FROM_HOST_GPU(
std::shared_ptr<SatelliteOribtNode[]> sateOirbtNodes,
std::shared_ptr<CUDA_AntSate_PtrList> gpupptr,
long startPID,
long PRF_len
);
class RFPCProcessCls
{
public:
@ -74,3 +91,6 @@ private:
void RFPCProcessMain(long num_thread,QString TansformPatternFilePath,QString ReceivePatternFilePath,QString simulationtaskName, QString OutEchoPath, QString GPSXmlPath,QString TaskXmlPath,
QString demTiffPath, QString sloperPath, QString LandCoverPath);

3
SimulationSAR/SARSatelliteSimulationAbstractCls.h
View File

@ -348,3 +348,6 @@ double getDopplerFreqRate(double& lamda, double& R, Vector3D& Rs, Vector3D& Rt,

BIN
resource/toolboxIcon.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 44 KiB

BIN
resource/toolicon.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 10 KiB