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

解决冲突

pull/3/head
陈增辉 2025-01-20 15:53:46 +08:00
parent 0741fab7ed 1b8340d39a
commit ea8ed343f7
9 changed files with 820 additions and 484 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
BaseTool/BaseConstVariable.h
View File

@ -26,6 +26,22 @@
#include <complex> #include <complex>
#include <time.h> #include <time.h>
/** 打印时间 ***************************************************************/
inline char* get_cur_time() {
static char s[20];
time_t t;
struct tm* ltime;
time(&t);
ltime = localtime(&t);
strftime(s, 20, "%Y-%m-%d %H:%M:%S", ltime);
return s;
}
#define PRINT(fmt, ...) printf("%s " fmt, get_cur_time(), ##__VA_ARGS__)
#define MATPLOTDRAWIMAGE #define MATPLOTDRAWIMAGE
@ -175,5 +191,12 @@ inline void PrintTime() {
printf("Current timestamp in seconds: %ld\n", (long)current_time); printf("Current timestamp in seconds: %ld\n", (long)current_time);
} }
/** 计算分块 ******************************************************************/
inline long getBlockRows(long sizeMB, long cols,long sizeMeta) {
return (round(Memory1MB * 1.0 / sizeMeta * sizeMB) + cols - 1) / cols;
}
#endif #endif

58
BaseTool/ImageOperatorBase.cpp
View File

@ -757,6 +757,64 @@ Eigen::MatrixXi gdalImage::getDatai(int start_row, int start_col, int rows_count
return datamatrix; return datamatrix;
} }
ErrorCode gdalImage::getData(double* data, int start_row, int start_col, int rows_count, int cols_count, int band_ids)
{
ErrorCode state =ErrorCode::SUCCESS;
omp_lock_t lock;
omp_init_lock(&lock);
omp_set_lock(&lock);
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
GDALDataset* rasterDataset = (GDALDataset*)(GDALOpen(this->img_path.toUtf8().constData(), GA_ReadOnly)); // 锟斤拷只斤拷式锟斤拷取斤拷影锟斤拷
GDALDataType gdal_datatype = rasterDataset->GetRasterBand(1)->GetRasterDataType();
GDALRasterBand* demBand = rasterDataset->GetRasterBand(band_ids);
rows_count = start_row + rows_count <= this->height ? rows_count : this->height - start_row;
cols_count = start_col + cols_count <= this->width ? cols_count : this->width - start_col;
if (gdal_datatype == GDT_Float64) {
demBand->RasterIO(GF_Read, start_col, start_row, cols_count, rows_count, data, cols_count,rows_count, gdal_datatype, 0, 0);
}
else {
state = ErrorCode::FAIL;
}
GDALClose((GDALDatasetH)rasterDataset);
omp_unset_lock(&lock); // 锟酵放伙拷斤拷
omp_destroy_lock(&lock); // 劫伙拷斤拷
return state;
}
ErrorCode gdalImage::getData(long* data, int start_row, int start_col, int rows_count, int cols_count, int band_ids)
{
ErrorCode state = ErrorCode::SUCCESS;
omp_lock_t lock;
omp_init_lock(&lock);
omp_set_lock(&lock);
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
GDALDataset* rasterDataset = (GDALDataset*)(GDALOpen(this->img_path.toUtf8().constData(), GA_ReadOnly)); // 锟斤拷只斤拷式锟斤拷取斤拷影锟斤拷
GDALDataType gdal_datatype = rasterDataset->GetRasterBand(1)->GetRasterDataType();
GDALRasterBand* demBand = rasterDataset->GetRasterBand(band_ids);
rows_count = start_row + rows_count <= this->height ? rows_count : this->height - start_row;
cols_count = start_col + cols_count <= this->width ? cols_count : this->width - start_col;
if (gdal_datatype == GDT_Int32) {
demBand->RasterIO(GF_Read, start_col, start_row, cols_count, rows_count, data, cols_count,rows_count, gdal_datatype, 0, 0);
}
else {
state = ErrorCode::FAIL;
}
GDALClose((GDALDatasetH)rasterDataset);
omp_unset_lock(&lock); // 锟酵放伙拷斤拷
omp_destroy_lock(&lock); // 劫伙拷斤拷
return state;
}
Eigen::MatrixXd gdalImage::getGeoTranslation() Eigen::MatrixXd gdalImage::getGeoTranslation()
{ {
return this->gt; return this->gt;

3
BaseTool/ImageOperatorBase.h
View File

@ -164,7 +164,8 @@ public: // 方法
virtual void setData(Eigen::MatrixXd,int data_band_ids=1); virtual void setData(Eigen::MatrixXd,int data_band_ids=1);
virtual Eigen::MatrixXd getData(int start_row, int start_col, int rows_count, int cols_count, int band_ids); virtual Eigen::MatrixXd getData(int start_row, int start_col, int rows_count, int cols_count, int band_ids);
virtual Eigen::MatrixXi getDatai(int start_row, int start_col, int rows_count, int cols_count, int band_ids); virtual Eigen::MatrixXi getDatai(int start_row, int start_col, int rows_count, int cols_count, int band_ids);
virtual ErrorCode getData(double* data, int start_row, int start_col, int rows_count, int cols_count, int band_ids);
virtual ErrorCode getData(long* data, int start_row, int start_col, int rows_count, int cols_count, int band_ids);
virtual Eigen::MatrixXd getGeoTranslation(); virtual Eigen::MatrixXd getGeoTranslation();
virtual GDALDataType getDataType(); virtual GDALDataType getDataType();
virtual void saveImage(Eigen::MatrixXd, int start_row, int start_col, int band_ids); virtual void saveImage(Eigen::MatrixXd, int start_row, int start_col, int band_ids);

378
GPUTool/GPURFPC.cu
View File

@ -547,83 +547,333 @@ __global__ void CUDAKernel_PRF_GeneratorEcho(float* Rarr, float* ampArr,
// 计算每块
__global__ void CUDA_Kernel_Computer_R_amp(
/** 对外封装接口 *******************************************************************************************************/ double* antX, double* antY, double* antZ,
double* antXaxisX, double* antXaxisY, double* antXaxisZ,
extern "C" void CUDA_RFPC_MainBlock( double* antYaxisX, double* antYaxisY, double* antYaxisZ,
double* antX, double* antY, double* antZ, // 天线的坐标 double* antZaxisX, double* antZaxisY, double* antZaxisZ,
double* antXaxisX, double* antXaxisY, double* antXaxisZ, // 天线坐标系的X轴 double* antDirectX, double* antDirectY, double* antDirectZ,
double* antYaxisX, double* antYaxisY, double* antYaxisZ,// 天线坐标系的Y轴 long sPid, long PRFCount,
double* antZaxisX, double* antZaxisY, double* antZaxisZ,// 天线坐标系的Z轴 double* targetX, double* targetY, double* targetZ, long* demCls, long TargetNumber,
double* antDirectX, double* antDirectY, double* antDirectZ,// 天线的指向 double* demSlopeX, double* demSlopeY, double* demSlopeZ,
long startpid,long PRFCount, // 脉冲数 long sPosId,long pixelcount,
float f0, float dfreq, long freqnum, // 频率数 CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
double* targetX, double* targetY, double* targetZ, long TargetPixelNumber, // 地面坐标 double Pt,
long* demCls, // 地表类别
double* demSlopeX, double* demSlopeY, double* demSlopeZ, // 地表坡度矢量
double Pt,// 发射能量
double refPhaseRange, double refPhaseRange,
double* TransAntpattern, double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图 double* TransAntpattern,
double* ReceiveAntpattern, double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图 double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints,
double NearR, double FarR, // 距离范围 double* ReceiveAntpattern,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图 double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,
float* out_echoReal, float* out_echoImag,// 输出回波 double NearR, double FarR,
float* temp_R, float* temp_amp long BlockPRFCount,
//, double* temp_phi, double* temp_real, double* temp_imag// 临时变量 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 blocknum = 0; double RstR = sqrt(RstX * RstX + RstY * RstY + RstZ * RstZ); // 矢量距离
long pixelcount=TargetPixelNumber; if (RstR<NearR || RstR>FarR) {
int numBlocks = 0; d_temp_R[idx] = 0;
for(long pid=0;pid<PRFCount;pid++){ d_temp_amps[idx] = 0;
numBlocks = (TargetPixelNumber + BLOCK_SIZE - 1) / BLOCK_SIZE; // 根据 pixelcount 计算网格大小 }
CUDAKernel_RFPC_Computer_R_Gain<<<numBlocks , BLOCK_SIZE >>>( else {
antX[startpid+pid], antY[startpid+pid], antZ[startpid+pid], double slopeX = demSlopeX[aposId];
targetX, targetY, targetZ, TargetPixelNumber, double slopeY = demSlopeY[aposId];
demCls, double slopeZ = demSlopeZ[aposId];
demSlopeX, demSlopeY, demSlopeZ,
antXaxisX[startpid+pid], antXaxisY[startpid+pid], antXaxisZ[startpid+pid], double slopR = sqrtf(slopeX * slopeX + slopeY * slopeY + slopeZ * slopeZ); //
antYaxisX[startpid+pid], antYaxisY[startpid+pid], antYaxisZ[startpid+pid], double dotAB = RstX * slopeX + RstY * slopeY + RstZ * slopeZ;
antZaxisX[startpid+pid], antZaxisY[startpid+pid], antZaxisZ[startpid+pid], double localangle = acosf(dotAB / (RstR * slopR)); // 局地入射角
antDirectX[startpid + pid], antDirectY[startpid + pid], antDirectZ[startpid + pid],
Pt,// 增益后发射能量 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(
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 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)); // 反射强度
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;
}
}
}
}
__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_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] ;
long tid = threadIdx.x;
long bid = blockIdx.x;
long idx= bid * blockDim.x + tid;
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];
}
__syncthreads(); // 确定所有待处理数据都已经进入程序中
long prfId = idx / FreqPoints; // 脉冲
long fId = idx % FreqPoints;// 频率
if (fId < maxfreqnum&& prfId< temp_PRF_Count) {
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);
}
d_temp_echo_real[idx] += temp_real;
d_temp_echo_imag[idx] += 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* demSlopeX, double* demSlopeY, double* demSlopeZ,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
float* out_echoReal, float* out_echoImag)
{
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;
float* h_tX = (float*)mallocCUDAHost(sizeof(float) * BlockTarlist);
float* h_tY = (float*)mallocCUDAHost(sizeof(float) * BlockTarlist);
float* h_tZ = (float*)mallocCUDAHost(sizeof(float) * BlockTarlist);
float* h_sloperX = (float*)mallocCUDAHost(sizeof(float) * BlockTarlist);
float* h_sloperY = (float*)mallocCUDAHost(sizeof(float) * BlockTarlist);
float* h_sloperZ = (float*)mallocCUDAHost(sizeof(float) * 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 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);
}
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];
}
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);
}
// 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, refPhaseRange,
TransAntpattern, TransAntpattern,
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints, Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
ReceiveAntpattern, ReceiveAntpattern,
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints, Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
NearR, FarR, NearR, FarR,
sigma0Paramslist, sigmaparamslistlen, BlockPRFCount,
//factorj, freqnum, BlockTarlist, // 模块
temp_R, // 输出距离 d_temp_R, d_temp_amp// 计算输出
temp_amp
//out_echoReal, out_echoImag, pid // 输出振幅
); );
//cudaDeviceSynchronize();
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDA_RFPC_MainBlock [CUDAKernel_RFPC_Computer_R_Gain] CUDA Error [pid:%d] : %s\n", startpid + pid, cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
} }
#endif // __CUDADEBUG__ // step 3 计算回波
blocknum = (pixelcount + GPU_SHARE_STEP - 1) / GPU_SHARE_STEP; {
numBlocks = (blocknum + BLOCK_SIZE - 1) / BLOCK_SIZE; // 网格数量 cudaBlocknum = (BlockPRFCount * freqpoints + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDAKernel_PRF_GeneratorEcho << <numBlocks, BLOCK_SIZE >> > CUDA_Kernel_Computer_echo << <cudaBlocknum, BLOCK_SIZE >> > (
(temp_R, temp_amp, pixelcount, d_temp_R, d_temp_amp, BlockTarlist,
f0,dfreq,freqnum, f0, dfreq, freqpoints, FreqNum,
out_echoReal, out_echoImag, pid); d_temp_echo_real, d_temp_echo_imag,
#ifdef __CUDADEBUG__ BlockPRFCount
err = cudaGetLastError(); );
if (err != cudaSuccess) {
printf("CUDA_RFPC_MainBlock [CUDAKernel_PRF_GeneratorEcho] CUDA Error [pid:%d] : %s\n", startpid+pid,cudaGetErrorString(err));
// Possibly: exit(-1) if program cannot continue....
} }
//cudaDeviceSynchronize();
#endif // __CUDADEBUG__ PRINT("PRF %d , TargetID: %d / %d", spid+ BlockPRFCount, sTi, sTi+ BlockTarlist);
} }
cudaDeviceSynchronize();
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);
} }

54
GPUTool/GPURFPC.cuh
View File

@ -99,29 +99,61 @@ 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_MainBlock(
double* antX, double* antY, double* antZ, // 天线的坐标 extern "C" void CUDA_RFPC_MainProcess(
// 天线
double* antX, double* antY, double* antZ, // 天线坐标
double* antXaxisX, double* antXaxisY, double* antXaxisZ, // 天线坐标系的X轴 double* antXaxisX, double* antXaxisY, double* antXaxisZ, // 天线坐标系的X轴
double* antYaxisX, double* antYaxisY, double* antYaxisZ,// 天线坐标系的Y轴 double* antYaxisX, double* antYaxisY, double* antYaxisZ,// 天线坐标系的Y轴
double* antZaxisX, double* antZaxisY, double* antZaxisZ,// 天线坐标系的Z轴 double* antZaxisX, double* antZaxisY, double* antZaxisZ,// 天线坐标系的Z轴
double* antDirectX, double* antDirectY, double* antDirectZ,// 天线的指向 double* antDirectX, double* antDirectY, double* antDirectZ,// 天线的指向
long startpid, long PRFCount, // 脉冲数 long PRFCount, long FreqNum, // 脉冲数量，频率数量
float f0, float dfreq, long freqnum,// 频率数 float f0, float dfreq,// 起始频率，终止频率
double* targetX, double* targetY, double* targetZ, long TargetPixelNumber, // 地面坐标
long* demCls, // 地表类别
double* demSlopeX, double* demSlopeY, double* demSlopeZ, // 地表坡度矢量
double Pt,// 发射能量 double Pt,// 发射能量
double refPhaseRange, double refPhaseRange,
// 天线方向图
double* TransAntpattern, double Transtarttheta, double Transstartphi, double Transdtheta, double Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图 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* ReceiveAntpattern, double Receivestarttheta, double Receivestartphi, double Receivedtheta, double Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
double NearR, double FarR, // 距离范围 double NearR, double FarR, // 距离范围
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
float* out_echoReal, float* out_echoImag,// 输出回波
float* temp_R, float* temp_amp // 地面
//,double* temp_phi ,double* temp_real, double* tmep_imag// 临时变量 double* targetX, double* targetY, double* targetZ, long* demCls, long TargetPixelNumber, // 地面坐标、地表覆盖类型，像素数
double* demSlopeX, double* demSlopeY, double* demSlopeZ,// 地表坡度矢量
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图像
float* out_echoReal, float* out_echoImag// 输出回波
); );
#endif #endif

63
GPUTool/GPUTool.cu
View File

@ -16,6 +16,45 @@
#define BLOCK_DIM 1024 #define BLOCK_DIM 1024
#define REDUCE_SCALE 4 #define REDUCE_SCALE 4
// ´òÓ¡GPU²ÎÊý
void printDeviceInfo(int deviceId) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, deviceId);
std::cout << "Device " << deviceId << ": " << deviceProp.name << std::endl;
std::cout << " Compute Capability: " << deviceProp.major << "." << deviceProp.minor << std::endl;
std::cout << " Total Global Memory: " << deviceProp.totalGlobalMem / (1024 * 1024) << " MB" << std::endl;
std::cout << " Shared Memory per Block: " << deviceProp.sharedMemPerBlock << " Bytes" << std::endl;
std::cout << " Registers per Block: " << deviceProp.regsPerBlock << std::endl;
std::cout << " Warp Size: " << deviceProp.warpSize << std::endl;
std::cout << " Max Threads per Block: " << deviceProp.maxThreadsPerBlock << std::endl;
std::cout << " Max Threads Dim: (" << deviceProp.maxThreadsDim[0] << ", "
<< deviceProp.maxThreadsDim[1] << ", " << deviceProp.maxThreadsDim[2] << ")" << std::endl;
std::cout << " Max Grid Size: (" << deviceProp.maxGridSize[0] << ", "
<< deviceProp.maxGridSize[1] << ", " << deviceProp.maxGridSize[2] << ")" << std::endl;
std::cout << " Multiprocessor Count: " << deviceProp.multiProcessorCount << std::endl;
std::cout << " Clock Rate: " << deviceProp.clockRate / 1000 << " MHz" << std::endl;
std::cout << " Memory Clock Rate: " << deviceProp.memoryClockRate / 1000 << " MHz" << std::endl;
std::cout << " Memory Bus Width: " << deviceProp.memoryBusWidth << " bits" << std::endl;
std::cout << " L2 Cache Size: " << deviceProp.l2CacheSize / 1024 << " KB" << std::endl;
std::cout << " Max Texture Dimensions: (" << deviceProp.maxTexture1D << ", "
<< deviceProp.maxTexture2D[0] << "x" << deviceProp.maxTexture2D[1] << ", "
<< deviceProp.maxTexture3D[0] << "x" << deviceProp.maxTexture3D[1] << "x" << deviceProp.maxTexture3D[2] << ")" << std::endl;
std::cout << " Unified Addressing: " << (deviceProp.unifiedAddressing ? "Yes" : "No") << std::endl;
std::cout << " Concurrent Kernels: " << (deviceProp.concurrentKernels ? "Yes" : "No") << std::endl;
std::cout << " ECC Enabled: " << (deviceProp.ECCEnabled ? "Yes" : "No") << std::endl;
std::cout << " PCI Bus ID: " << deviceProp.pciBusID << std::endl;
std::cout << " PCI Device ID: " << deviceProp.pciDeviceID << std::endl;
std::cout << " PCI Domain ID: " << deviceProp.pciDomainID << std::endl;
std::cout << std::endl;
}
// ¶¨Òå²ÎÊý // ¶¨Òå²ÎÊý
__device__ cuComplex cuCexpf(cuComplex d) __device__ cuComplex cuCexpf(cuComplex d)
{ {
@ -45,7 +84,15 @@ __device__ float GPU_CosAngle_VectorA_VectorB(CUDAVector A, CUDAVector B) {
__global__ void CUDAKernel_MemsetBlock(cuComplex* data, cuComplex init0, long len) { extern __global__ void CUDAKernel_MemsetBlock(cuComplex* data, cuComplex init0, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
data[idx] = init0;
}
}
extern __global__ void CUDAKernel_MemsetBlock(float* data, float init0, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x; long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) { if (idx < len) {
data[idx] = init0; data[idx] = init0;
@ -83,15 +130,6 @@ __global__ void CUDACkernel_SUM_reduce_dynamicshared(float* d_x, float* d_y, lo
__global__ void CUDA_DistanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* R, long len) { __global__ void CUDA_DistanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* R, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x; long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) { if (idx < len) {
@ -439,6 +477,11 @@ extern "C" void CUDADCos(double* y, double* X, int n)
} }
long NextBlockPad(long num, long blocksize)
{
return ((num + blocksize - 1) / blocksize) * blocksize;
}
#endif #endif

35
GPUTool/GPUTool.cuh
View File

@ -17,8 +17,11 @@
#define BLOCK_SIZE 256 #define BLOCK_SIZE 256
#define SHAREMEMORY_BYTE 49152
#define SHAREMEMORY_FLOAT_HALF 6144
// 打印GPU参数
void printDeviceInfo(int deviceId);
enum LAMPGPUDATETYPE { enum LAMPGPUDATETYPE {
@ -57,14 +60,6 @@ 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_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); extern __global__ void CUDA_B_DistanceA(float* Ax, float* Ay, float* Az, float Bx, float By, float Bz, float* R, long len);
@ -74,19 +69,9 @@ extern __global__ void CUDA_cosAngle_VA_AB(float* Ax, float* Ay, float* Az, floa
extern __global__ void CUDA_GridPoint_Linear_Interp1(float* v, float* q, float* qv, long xlen, long qlen); extern __global__ void CUDA_GridPoint_Linear_Interp1(float* v, float* q, float* qv, long xlen, long qlen);
extern __global__ void CUDA_D_sin(double* y, double* X, int n); extern __global__ void CUDA_D_sin(double* y, double* X, int n);
extern __global__ void CUDA_D_cos(double* y, double* X, int n); extern __global__ void CUDA_D_cos(double* y, double* X, int n);
extern __global__ void CUDAKernel_MemsetBlock(cuComplex* data, cuComplex init0, long len);
extern __global__ void CUDAKernel_MemsetBlock(float* data, float init0, long len);
@ -115,5 +100,13 @@ extern "C" void CUDAcosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, f
extern "C" void CUDAGridPointLinearInterp1(float* v, float* q, float* qv,long xlen, long qlen); extern "C" void CUDAGridPointLinearInterp1(float* v, float* q, float* qv,long xlen, long qlen);
extern "C" void CUDADSin(double* y, double* X, int n); extern "C" void CUDADSin(double* y, double* X, int n);
extern "C" void CUDADCos(double* y, double* X, int n); extern "C" void CUDADCos(double* y, double* X, int n);
// 估算分块整数
extern "C" long NextBlockPad(long num,long blocksize);
#endif #endif
#endif #endif

8
RasterProcessTool.vcxproj.filters
View File

@ -249,6 +249,7 @@
<ClInclude Include="BaseToolbox\WGS84_J2000.h"> <ClInclude Include="BaseToolbox\WGS84_J2000.h">
<Filter>BaseToolbox</Filter> <Filter>BaseToolbox</Filter>
</ClInclude> </ClInclude>
<ClInclude Include="GPUTool\GPUGarbage.cuh" />
</ItemGroup> </ItemGroup>
<ItemGroup> <ItemGroup>
<QtMoc Include="QMergeRasterProcessDialog.h"> <QtMoc Include="QMergeRasterProcessDialog.h">
@ -263,9 +264,6 @@
<QtMoc Include="QSimulationRFPCGUI.h"> <QtMoc Include="QSimulationRFPCGUI.h">
<Filter>Header Files</Filter> <Filter>Header Files</Filter>
</QtMoc> </QtMoc>
<QtMoc Include="BaseTool\QToolProcessBarDialog.h">
<Filter>Header Files</Filter>
</QtMoc>
<QtMoc Include="SimulationSAR\QToolAbstract.h"> <QtMoc Include="SimulationSAR\QToolAbstract.h">
<Filter>Header Files</Filter> <Filter>Header Files</Filter>
</QtMoc> </QtMoc>
@ -293,6 +291,9 @@
<QtMoc Include="BaseToolbox\QOrthSlrRaster.h"> <QtMoc Include="BaseToolbox\QOrthSlrRaster.h">
<Filter>BaseToolbox</Filter> <Filter>BaseToolbox</Filter>
</QtMoc> </QtMoc>
<QtMoc Include="BaseTool\QToolProcessBarDialog.h">
<Filter>BaseTool</Filter>
</QtMoc>
</ItemGroup> </ItemGroup>
<ItemGroup> <ItemGroup>
<QtUic Include="QMergeRasterProcessDialog.ui"> <QtUic Include="QMergeRasterProcessDialog.ui">
@ -342,6 +343,7 @@
<CudaCompile Include="GPUTool\GPUTool.cu"> <CudaCompile Include="GPUTool\GPUTool.cu">
<Filter>GPUTool</Filter> <Filter>GPUTool</Filter>
</CudaCompile> </CudaCompile>
<CudaCompile Include="GPUTool\GPUGarbage.cu" />
</ItemGroup> </ItemGroup>
<ItemGroup> <ItemGroup>
<None Include="cpp.hint" /> <None Include="cpp.hint" />

618
SimulationSAR/RFPCProcessCls.cpp
View File

@ -25,7 +25,7 @@
#include <cuda_runtime.h> #include <cuda_runtime.h>
#include <cublas_v2.h> #include <cublas_v2.h>
#endif // __CUDANVCC___ #endif // __CUDANVCC___
#include <Imageshow/ImageShowDialogClass.h> //#include <Imageshow/ImageShowDialogClass.h>
@ -329,130 +329,148 @@ void RFPCProcessMain(long num_thread,
} }
ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
{ /** 内存分配***************************************************/
double widthSpace = LIGHTSPEED / 2 / this->TaskSetting->getFs(); long TargetMemoryMB = 500;
double prf_time = 0;
double dt = 1 / this->TaskSetting->getPRF();// 获取每次脉冲的时间间隔
bool antflag = true; // 计算天线方向图 /** 参数区域***************************************************/
Landpoint LandP{ 0,0,0 };
Point3 GERpoint{ 0,0,0 };
double R = 0;
double dem_row = 0, dem_col = 0, dem_alt = 0;
QVector<double> freqlist = this->TaskSetting->getFreqList(); QVector<double> freqlist = this->TaskSetting->getFreqList();
long freqnum = freqlist.count(); long freqnum = freqlist.count();
std::shared_ptr<double> freqPtr(new double[freqnum], delArrPtr);
for (long ii = 0; ii < freqlist.count(); ii++) {
freqPtr.get()[ii] = freqlist[ii];
}
testOutAmpArr("freqlist.bin", (double*)(freqPtr.get()), freqnum, 1);
float f0 = float(freqlist[0] / 1e9); float f0 = float(freqlist[0] / 1e9);
float dfreq = float((freqlist[1] - freqlist[0]) / 1e9); float dfreq = float((freqlist[1] - freqlist[0]) / 1e9);
long PRFCount = this->EchoSimulationData->getPluseCount();
long double imageStarttime = 0;
imageStarttime = this->TaskSetting->getSARImageStartTime();
//std::vector<SatelliteOribtNode> sateOirbtNodes(this->PluseCount);
std::shared_ptr<SatelliteOribtNode[]> sateOirbtNodes = this->getSatelliteOribtNodes(prf_time, dt, antflag, imageStarttime); // 获取天线坐标
long echoIdx = 0;
double NearRange = this->EchoSimulationData->getNearRange(); // 近斜距 double NearRange = this->EchoSimulationData->getNearRange(); // 近斜距
double FarRange = this->EchoSimulationData->getFarRange(); double FarRange = this->EchoSimulationData->getFarRange();
double TimgNearRange = 2 * NearRange / LIGHTSPEED;
double TimgFarRange = 2 * FarRange / LIGHTSPEED;
double dx = (FarRange - NearRange) / (PlusePoint - 1);
double Fs = this->TaskSetting->getFs(); // 距离向采样率
double Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 发射电压 1v double Pt = this->TaskSetting->getPt() * this->TaskSetting->getGri();// 发射电压 1v
//double GainAntLen = -3;// -3dB 为天线半径
long pluseCount = this->PluseCount;
double lamda = this->TaskSetting->getCenterLamda(); // 波长 double lamda = this->TaskSetting->getCenterLamda(); // 波长
double refphaseRange = this->TaskSetting->getRefphaseRange(); // 参考相位斜距 double refphaseRange = this->TaskSetting->getRefphaseRange(); // 参考相位斜距
// 天线方向图
std::shared_ptr<AbstractRadiationPattern> TransformPattern = this->TaskSetting->getTransformRadiationPattern(); // 发射天线方向图
std::shared_ptr<AbstractRadiationPattern> ReceivePattern = this->TaskSetting->getReceiveRadiationPattern(); // 接收天线方向图
long PlusePoint = this->EchoSimulationData->getPlusePoints();
POLARTYPEENUM polartype = this->TaskSetting->getPolarType();
gdalImage echoMaskImg(this->OutEchoMaskPath);
long echoblockline = Memory1GB / 8 / 2 / PlusePoint * 2;
#ifndef __CUDANVCC___ double prf_time = 0;
QMessageBox::information(this, u8"程序提示", u8"请确定安装了CUDA库"); double dt = 1 / this->TaskSetting->getPRF();// 获取每次脉冲的时间间隔
#else 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;
double* antpx,*antpy,*antpz,
*antvx,*antvy,*antvz,
*antdirectx,*antdirecty,*antdirectz
,*antXaxisX,*antXaxisY,*antXaxisZ,
*antYaxisX,*antYaxisY,*antYaxisZ,
*antZaxisX,*antZaxisY,*antZaxisZ;
{ {
antpx = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antpx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antpy = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antpy = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antpz = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antpz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antvx = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antvx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antvy = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antvy = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antvz = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antvz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antdirectx = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antdirectx = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antdirecty = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antdirecty = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antdirectz = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antdirectz = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antXaxisX = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antXaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antXaxisY = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antXaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antXaxisZ = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antXaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antYaxisX = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antYaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antYaxisY = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antYaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antYaxisZ = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antYaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antZaxisX = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antZaxisX = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antZaxisY = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antZaxisY = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
antZaxisZ = (double*)mallocCUDAHost(sizeof(double) * pluseCount); h_antZaxisZ = (double*)mallocCUDAHost(sizeof(double) * PRFCount);
for (long tempprfid = 0; tempprfid < pluseCount; tempprfid++) { 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; long prfid = tempprfid;
antpx[tempprfid] = sateOirbtNodes[prfid].Px; h_antpx[tempprfid] = sateOirbtNodes[prfid].Px;
antpy[tempprfid] = sateOirbtNodes[prfid].Py; h_antpy[tempprfid] = sateOirbtNodes[prfid].Py;
antpz[tempprfid] = sateOirbtNodes[prfid].Pz; h_antpz[tempprfid] = sateOirbtNodes[prfid].Pz;
antvx[tempprfid] = sateOirbtNodes[prfid].Vx; h_antvx[tempprfid] = sateOirbtNodes[prfid].Vx;
antvy[tempprfid] = sateOirbtNodes[prfid].Vy; h_antvy[tempprfid] = sateOirbtNodes[prfid].Vy;
antvz[tempprfid] = sateOirbtNodes[prfid].Vz; //6 h_antvz[tempprfid] = sateOirbtNodes[prfid].Vz; //6
antdirectx[tempprfid] = sateOirbtNodes[prfid].AntDirecX; h_antdirectx[tempprfid] = sateOirbtNodes[prfid].AntDirecX;
antdirecty[tempprfid] = sateOirbtNodes[prfid].AntDirecY; h_antdirecty[tempprfid] = sateOirbtNodes[prfid].AntDirecY;
antdirectz[tempprfid] = sateOirbtNodes[prfid].AntDirecZ; // 9 天线指向 h_antdirectz[tempprfid] = sateOirbtNodes[prfid].AntDirecZ; // 9 天线指向
antXaxisX[tempprfid] = sateOirbtNodes[prfid].AntXaxisX; h_antXaxisX[tempprfid] = sateOirbtNodes[prfid].AntXaxisX;
antXaxisY[tempprfid] = sateOirbtNodes[prfid].AntXaxisY; h_antXaxisY[tempprfid] = sateOirbtNodes[prfid].AntXaxisY;
antXaxisZ[tempprfid] = sateOirbtNodes[prfid].AntXaxisZ;//12 天线坐标系 h_antXaxisZ[tempprfid] = sateOirbtNodes[prfid].AntXaxisZ;//12 天线坐标系
antYaxisX[tempprfid] = sateOirbtNodes[prfid].AntYaxisX; h_antYaxisX[tempprfid] = sateOirbtNodes[prfid].AntYaxisX;
antYaxisY[tempprfid] = sateOirbtNodes[prfid].AntYaxisY; h_antYaxisY[tempprfid] = sateOirbtNodes[prfid].AntYaxisY;
antYaxisZ[tempprfid] = sateOirbtNodes[prfid].AntYaxisZ;//15 h_antYaxisZ[tempprfid] = sateOirbtNodes[prfid].AntYaxisZ;//15
antZaxisX[tempprfid] = sateOirbtNodes[prfid].AntZaxisX; h_antZaxisX[tempprfid] = sateOirbtNodes[prfid].AntZaxisX;
antZaxisY[tempprfid] = sateOirbtNodes[prfid].AntZaxisY; h_antZaxisY[tempprfid] = sateOirbtNodes[prfid].AntZaxisY;
antZaxisZ[tempprfid] = sateOirbtNodes[prfid].AntZaxisZ;//18 h_antZaxisZ[tempprfid] = sateOirbtNodes[prfid].AntZaxisZ;//18
}
} }
// RFPC CUDA版本 DeviceToDevice(h_antpx, d_antpx, sizeof(double) * PRFCount);
if (pluseCount * 4 * 18 > Memory1MB * 100) { DeviceToDevice(h_antpy, d_antpy, sizeof(double) * PRFCount);
long max = Memory1MB * 100 / 4 / 20 / PluseCount; DeviceToDevice(h_antpz, d_antpz, sizeof(double) * PRFCount);
QMessageBox::warning(nullptr, u8"仿真场景太大了", u8"当前频点数下，脉冲数量最多为：" + QString::number(max)); 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);
} }
gdalImage demxyz(this->demxyzPath);// 地面点坐标
gdalImage demlandcls(this->LandCoverPath);// 地表覆盖类型
gdalImage demsloperxyz(this->demsloperPath);// 地面坡向
// 参数与分块计算
long demRow = demxyz.height;
long demCol = demxyz.width;
long blokline = 100;
// 每块 250MB*16 = 4GB /** 天线方向图***************************************************/
std::shared_ptr<AbstractRadiationPattern> TransformPattern = this->TaskSetting->getTransformRadiationPattern(); // 发射天线方向图
blokline = Memory1MB / 8 / demCol * 500; std::shared_ptr<AbstractRadiationPattern> ReceivePattern = this->TaskSetting->getReceiveRadiationPattern(); // 接收天线方向图
blokline = blokline < 1 ? 1 : blokline;
bool bloklineflag = false; POLARTYPEENUM polartype = this->TaskSetting->getPolarType();
PatternImageDesc TantPatternDesc = {}; PatternImageDesc TantPatternDesc = {};
double* h_TantPattern=nullptr; double* h_TantPattern = nullptr;
double* d_TantPattern=nullptr; double* d_TantPattern = nullptr;
{ {
// 处理发射天线方向图 // 处理发射天线方向图
@ -498,7 +516,6 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
PatternImageDesc RantPatternDesc = {}; PatternImageDesc RantPatternDesc = {};
double* h_RantPattern = nullptr; double* h_RantPattern = nullptr;
double* d_RantPattern = nullptr; double* d_RantPattern = nullptr;
{ {
// 处理接收天线方向图 // 处理接收天线方向图
double Rminphi = ReceivePattern->getMinPhi(); double Rminphi = ReceivePattern->getMinPhi();
@ -539,12 +556,25 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
RantPatternDesc.phinum = Rphinum; RantPatternDesc.phinum = Rphinum;
RantPatternDesc.thetanum = Rthetanum; RantPatternDesc.thetanum = Rthetanum;
} }
/** 坐标区域点***************************************************/
gdalImage demxyz(this->demxyzPath);// 地面点坐标
gdalImage demlandcls(this->LandCoverPath);// 地表覆盖类型
gdalImage demsloperxyz(this->demsloperPath);// 地面坡向
long demRow = demxyz.height;
long demCol = demxyz.width;
//处理地表覆盖 //处理地表覆盖
QMap<long, long> clamap; QMap<long, long> clamap;
long clamapid = 0; long clamapid = 0;
long startline = 0; long startline = 0;
{ {
long blokline = getBlockRows(2e4, demCol, sizeof(double));
for (startline = 0; startline < demRow; startline = startline + blokline) { for (startline = 0; startline < demRow; startline = startline + blokline) {
Eigen::MatrixXd clsland = demlandcls.getData(startline, 0, blokline, demlandcls.width, 1); Eigen::MatrixXd clsland = demlandcls.getData(startline, 0, blokline, demlandcls.width, 1);
long clsrows = clsland.rows(); long clsrows = clsland.rows();
@ -599,269 +629,173 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
HostToDevice(h_clsSigmaParam, d_clsSigmaParam, sizeof(CUDASigmaParam) * clamapid); HostToDevice(h_clsSigmaParam, d_clsSigmaParam, sizeof(CUDASigmaParam) * clamapid);
long blockwidth = demxyz.width; // 处理地面坐标
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 echo_block_rows = getBlockRows(5000, freqnum, sizeof(float));
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);
/** 主流程处理 ***************************************************/
for (long sprfid = 0; sprfid < PRFCount; sprfid = sprfid + echo_block_rows) {
long PRF_len = (sprfid + echo_block_rows) < PRFCount ? echo_block_rows : (PRFCount - sprfid);
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++) {
h_echo_block_real[ii * freqnum + jj]=echo_temp.get()[ii * freqnum + jj].real();
h_echo_block_imag[ii * freqnum + jj]=echo_temp.get()[ii * freqnum + jj].imag();
}
}
for (long 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);
Eigen::MatrixXd demsloper_y = demsloperxyz.getData(startline, 0, blockline, demCol, 2);
Eigen::MatrixXd demsloper_z = demsloperxyz.getData(startline, 0, blockline, demCol, 3);
Eigen::MatrixXd landcover = demlandcls.getData(startline, 0, blockline, demCol, 1);
long temp_dem_row = dem_x.rows();
long temp_dem_col = dem_x.cols();
long temp_dem_count = dem_x.count();
// 更新数据格式
for (long i = 0; i < temp_dem_row; i++) {
for (long j = 0; j < temp_dem_col; j++) {
#ifdef __PRFDEBUG__ #ifdef __PRFDEBUG__
blokline = 1; h_dem_x[i * temp_dem_col + j] = -2028380.6250000; double(dem_x(i, j));
blockwidth = 1; h_dem_y[i * temp_dem_col + j] = 4139373.250000; double(dem_y(i, j));
#endif h_dem_z[i * temp_dem_col + j] = 4393382.500000; double(dem_z(i, j));
// 地面 XYZ h_demsloper_x[i * temp_dem_col + j] = 4393382.500000; double(demsloper_x(i, j));
Eigen::MatrixXd dem_x = demxyz.getData(0, 0, blokline, blockwidth, 1); // 地面坐标 h_demsloper_y[i * temp_dem_col + j] = 446.923950; double(demsloper_y(i, j));
long tempDemRows = dem_x.rows(); h_demsloper_z[i * temp_dem_col + j] = -219.002213; double(demsloper_z(i, j));
long tempDemCols = dem_x.cols(); h_demcls[i * temp_dem_col + j] = clamap[80];// clamap[long(landcover(i, j))];
Eigen::MatrixXd dem_y = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd dem_z = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd demsloper_x = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd demsloper_y = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
Eigen::MatrixXd demsloper_z = Eigen::MatrixXd::Zero(tempDemRows, tempDemCols);
double* h_dem_x = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* h_dem_y = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* h_dem_z = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* h_demsloper_x = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* h_demsloper_y = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* h_demsloper_z = (double*)mallocCUDAHost(sizeof(double) * blokline * tempDemCols);
double* d_dem_x = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols); // 7
double* d_dem_y = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols);
double* d_dem_z = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols);
double* d_demsloper_x = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols);
double* d_demsloper_y = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols);
double* d_demsloper_z = (double*)mallocCUDADevice(sizeof(double) * blokline * tempDemCols);
// 提前声明参数变量
float* h_R = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* d_R = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
float* h_amp = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
float* d_amp = (float*)mallocCUDADevice(sizeof(float) * blokline * tempDemCols);
float* h_PRFEcho_real = (float*)mallocCUDAHost(sizeof(float) * echoblockline * PlusePoint);
float* h_PRFEcho_imag = (float*)mallocCUDAHost(sizeof(float) * echoblockline * PlusePoint);
float* d_PRFEcho_real = (float*)mallocCUDADevice(sizeof(float) * echoblockline * PlusePoint);
float* d_PRFEcho_imag = (float*)mallocCUDADevice(sizeof(float) * echoblockline * PlusePoint);
double* h_factorj = (double*)mallocCUDAHost(sizeof(double) * freqlist.size());
double* h_freqlist = (double*)mallocCUDAHost(sizeof(double) * freqlist.size());
for (long ii = 0; ii < freqlist.size(); ii++) {
h_factorj[ii] = -4 * PI * freqlist[ii] / LIGHTSPEED;
h_freqlist[ii] = freqlist[ii];
}
testOutAmpArr("freqlist.bin", h_freqlist, freqlist.size(), 1);
testOutAmpArr("factorj.bin", h_factorj, freqlist.size(), 1);
// 地表覆盖类型
Eigen::MatrixXd landcover = Eigen::MatrixXd::Zero(blokline, tempDemCols);// 地面覆盖类型
long* h_demcls = (long*)mallocCUDAHost(sizeof(long) * blokline * tempDemCols);
long* d_demcls = (long*)mallocCUDADevice(sizeof(long) * blokline * tempDemCols);
for (startline = 0; startline < demRow; startline = startline + blokline) {
long newblokline = blokline;
if ((startline + blokline) >= demRow) {
newblokline = demRow - startline;
bloklineflag = true;
}
dem_x = demxyz.getData(startline, 0, newblokline, blockwidth, 1); // 地面坐标
dem_y = demxyz.getData(startline, 0, newblokline, blockwidth, 2);
dem_z = demxyz.getData(startline, 0, newblokline, blockwidth, 3);
demsloper_x = demsloperxyz.getData(startline, 0, newblokline, blockwidth, 1);
demsloper_y = demsloperxyz.getData(startline, 0, newblokline, blockwidth, 2);
demsloper_z = demsloperxyz.getData(startline, 0, newblokline, blockwidth, 3);
landcover = demlandcls.getData(startline, 0, newblokline, blockwidth, 1);
long calpluseFreqBufferLen = Memory1GB / 8 / 2 / PlusePoint * 2;
if (calpluseFreqBufferLen < 1000) {
qDebug() << "frequency point has morn than 50000";
QMessageBox::warning(nullptr, u8"frequency point has morn than 50000", u8"frequency point has morn than 50000");
}
if (bloklineflag) {
FreeCUDAHost(h_dem_x); FreeCUDADevice(d_dem_x);
FreeCUDAHost(h_dem_y); FreeCUDADevice(d_dem_y);
FreeCUDAHost(h_dem_z); FreeCUDADevice(d_dem_z);
FreeCUDAHost(h_demsloper_x); FreeCUDADevice(d_demsloper_x);
FreeCUDAHost(h_demsloper_y); FreeCUDADevice(d_demsloper_y);
FreeCUDAHost(h_demsloper_z); FreeCUDADevice(d_demsloper_z); //6
FreeCUDAHost(h_demcls); FreeCUDADevice(d_demcls);
FreeCUDAHost(h_R); FreeCUDADevice(d_R);
FreeCUDAHost(h_amp); FreeCUDADevice(d_amp);
//FreeCUDAHost(h_phi); FreeCUDADevice(d_phi);
//FreeCUDAHost(h_real); FreeCUDADevice(d_real);
//FreeCUDAHost(h_imag); FreeCUDADevice(d_imag);
h_dem_x = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_dem_y = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_dem_z = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_demsloper_x = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_demsloper_y = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_demsloper_z = (double*)mallocCUDAHost(sizeof(double) * newblokline * tempDemCols);
h_demcls = (long*)mallocCUDAHost(sizeof(long) * newblokline * tempDemCols);
d_dem_x = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);
d_dem_y = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);
d_dem_z = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);
d_demsloper_x = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);
d_demsloper_y = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);
d_demsloper_z = (double*)mallocCUDADevice(sizeof(double) * newblokline * tempDemCols);//6
d_demcls = (long*)mallocCUDADevice(sizeof(long) * newblokline * tempDemCols);
// 临时变量
h_R = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
d_R = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
h_amp = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
d_amp = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
}
//# pragma omp parallel for
for (long i = 0; i < newblokline; i++) {
for (long j = 0; j < blockwidth; j++) {
#ifdef __PRFDEBUG__
h_dem_x[i * blockwidth + j] = -2028380.6250000; double(dem_x(i, j));
h_dem_y[i * blockwidth + j] = 4139373.250000; double(dem_y(i, j));
h_dem_z[i * blockwidth + j] = 4393382.500000; double(dem_z(i, j));
h_demsloper_x[i * blockwidth + j] = 4393382.500000; double(demsloper_x(i, j));
h_demsloper_y[i * blockwidth + j] = 446.923950; double(demsloper_y(i, j));
h_demsloper_z[i * blockwidth + j] = -219.002213; double(demsloper_z(i, j));
h_demcls[i * blockwidth + j] = clamap[80];// clamap[long(landcover(i, j))];
#else #else
h_dem_x[i * blockwidth + j] = double(dem_x(i, j)); h_dem_x[i * temp_dem_col + j] = double(dem_x(i, j));
h_dem_y[i * blockwidth + j] = double(dem_y(i, j)); h_dem_y[i * temp_dem_col + j] = double(dem_y(i, j));
h_dem_z[i * blockwidth + j] = double(dem_z(i, j)); h_dem_z[i * temp_dem_col + j] = double(dem_z(i, j));
h_demsloper_x[i * blockwidth + j] = double(demsloper_x(i, j)); h_demsloper_x[i * temp_dem_col + j] = double(demsloper_x(i, j));
h_demsloper_y[i * blockwidth + j] = double(demsloper_y(i, j)); h_demsloper_y[i * temp_dem_col + j] = double(demsloper_y(i, j));
h_demsloper_z[i * blockwidth + j] = double(demsloper_z(i, j)); h_demsloper_z[i * temp_dem_col + j] = double(demsloper_z(i, j));
h_demcls[i * blockwidth + j] = clamap[long(landcover(i, j))]; h_demcls[i * temp_dem_col + j] = clamap[long(landcover(i, j))];
#endif #endif
} }
} }
HostToDevice((void*)h_dem_x, (void*)d_dem_x, sizeof(double) * newblokline * tempDemCols); // 复制 机器 -> GPU
HostToDevice((void*)h_dem_y, (void*)d_dem_y, sizeof(double) * newblokline * tempDemCols);
HostToDevice((void*)h_dem_z, (void*)d_dem_z, sizeof(double) * newblokline * tempDemCols);
HostToDevice((void*)h_demsloper_x, (void*)d_demsloper_x, sizeof(double) * newblokline * tempDemCols);
HostToDevice((void*)h_demsloper_y, (void*)d_demsloper_y, sizeof(double) * newblokline * tempDemCols);
HostToDevice((void*)h_demsloper_z, (void*)d_demsloper_z, sizeof(double) * newblokline * tempDemCols);
HostToDevice((void*)h_demcls, (void*)d_demcls, sizeof(long) * newblokline * tempDemCols);
#ifdef __PRFDEBUG__ && __PRFDEBUG_PRFINF__ // 分块处理
printf("tatgetPs=[%f,%f,%f]\n", h_dem_x[0], h_dem_y[0], h_dem_z[0]); CUDA_RFPC_MainProcess(
std::shared_ptr<double> h_temp_R(new double[PluseCount], delArrPtr); d_antpx, d_antpy, d_antpz,
#endif // __PRFDEBUG__ d_antXaxisX, d_antXaxisY, d_antXaxisZ, // 天线坐标系的X轴
d_antYaxisX, d_antYaxisY, d_antYaxisZ,// 天线坐标系的Y轴
long pixelcount = newblokline * tempDemCols; d_antZaxisX, d_antZaxisY, d_antZaxisZ,// 天线坐标系的Z轴
long startprfid = 0; d_antdirectx, d_antdirecty, d_antdirectz,// 天线的指向
for (startprfid = 0; startprfid < pluseCount; startprfid = startprfid + echoblockline) { PRFCount, freqnum,
std::cout << "[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] dem:\t" << startline << "\t-\t" << startline + newblokline << "\t:\t pluse :\t" << startprfid << " / " << pluseCount << std::endl; f0,dfreq,
long templine = startprfid + echoblockline < PluseCount ? echoblockline : PluseCount - startprfid; Pt,
std::shared_ptr<std::complex<double>> echotemp = this->EchoSimulationData->getEchoArr(startprfid, templine);
for (long tempprfid = 0; tempprfid < templine; tempprfid++) {
for (long freqid = 0; freqid < PlusePoint; freqid++) {
h_PRFEcho_real[tempprfid * PlusePoint + freqid] = 0;// echotemp.get()[tempprfid * PlusePoint + freqid].real();
h_PRFEcho_imag[tempprfid * PlusePoint + freqid] = 0;// echotemp.get()[tempprfid * PlusePoint + freqid].imag();
}
}
HostToDevice(h_PRFEcho_real, d_PRFEcho_real, sizeof(float) * echoblockline * PlusePoint);
HostToDevice(h_PRFEcho_imag, d_PRFEcho_imag, sizeof(float) * echoblockline * PlusePoint);
CUDA_RFPC_MainBlock(
antpx, antpy, antpz, // 天线的坐标
antXaxisX, antXaxisY, antXaxisZ, // 天线坐标系的X轴
antYaxisX, antYaxisY, antYaxisZ,// 天线坐标系的Y轴
antZaxisX, antZaxisY, antZaxisZ,// 天线坐标系的Z轴
antdirectx, antdirecty, antdirectz,// 天线的指向
startprfid,templine, // 脉冲数
//h_freqlist, h_factorj, PlusePoint,// 频率数
f0, dfreq, PlusePoint,// 频率数
d_dem_x, d_dem_y, d_dem_z, pixelcount, // 地面坐标
d_demcls,
d_demsloper_x, d_demsloper_y, d_demsloper_z, // 地表坡度矢量
Pt,// 增益后发射能量
refphaseRange, refphaseRange,
// 天线方向图
d_TantPattern, d_TantPattern,
TantPatternDesc.startTheta, TantPatternDesc.startPhi, TantPatternDesc.dtheta,TantPatternDesc.dphi,TantPatternDesc.thetanum,TantPatternDesc.phinum, TantPatternDesc.startTheta, TantPatternDesc.startPhi, TantPatternDesc.dtheta, TantPatternDesc.dphi, TantPatternDesc.thetanum, TantPatternDesc.phinum,
d_RantPattern, d_RantPattern,
RantPatternDesc.startTheta, RantPatternDesc.startPhi, RantPatternDesc.dtheta, RantPatternDesc.dphi, RantPatternDesc.thetanum, RantPatternDesc.phinum, RantPatternDesc.startTheta, RantPatternDesc.startPhi, RantPatternDesc.dtheta, RantPatternDesc.dphi, RantPatternDesc.thetanum, RantPatternDesc.phinum,
NearRange, FarRange,
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, // 地表坡度矢量
d_clsSigmaParam, clamapid, d_clsSigmaParam, clamapid,
d_PRFEcho_real, d_PRFEcho_imag,// 输出回波
d_R, d_amp h_echo_block_real, h_echo_block_imag// 输出回波
//, d_phi, d_real, d_imag// 临时变量
); );
DeviceToHost(h_PRFEcho_real, d_PRFEcho_real, sizeof(float) * echoblockline * PlusePoint); PRINT("dem : %d - %d / %d , echo: %d -%d / %d", startline, startline+ temp_dem_row, demRow, sprfid, sprfid+ PRF_len, PRFCount);
DeviceToHost(h_PRFEcho_imag, d_PRFEcho_imag, sizeof(float) * echoblockline * PlusePoint);
for (long tempprfid = 0; tempprfid < templine; tempprfid++) {
for (long freqid = 0; freqid < PlusePoint; freqid++) {
echotemp.get()[tempprfid * PlusePoint + freqid].real(
echotemp.get()[tempprfid * PlusePoint + freqid].real() + h_PRFEcho_real[tempprfid * PlusePoint + freqid]);
echotemp.get()[tempprfid * PlusePoint + freqid].imag(
echotemp.get()[tempprfid * PlusePoint + freqid].imag() + h_PRFEcho_imag[tempprfid * PlusePoint + freqid]);
}
}
this->EchoSimulationData->saveEchoArr(echotemp, startprfid, templine);
} }
#ifdef __PRFDEBUG__ && __PRFDEBUG_PRFINF__
break; for (long ii = 0; ii < PRF_len; ii++) {
#endif // __PRFDEBUG__ for (long jj = 0; jj < freqnum; jj++) {
echo_temp.get()[ii * freqnum + jj].real(h_echo_block_real[ii * freqnum + jj]);
echo_temp.get()[ii * freqnum + jj].imag(h_echo_block_imag[ii * freqnum + jj]);
}
}
this->EchoSimulationData->saveEchoArr(echo_temp, sprfid, PRF_len);
} }
std::cout << std::endl;
// 地面数据释放
FreeCUDAHost(h_dem_x); FreeCUDADevice(d_dem_x);
FreeCUDAHost(h_dem_y); FreeCUDADevice(d_dem_y);
FreeCUDAHost(h_dem_z); FreeCUDADevice(d_dem_z);
FreeCUDAHost(h_demsloper_x); FreeCUDADevice(d_demsloper_x);
FreeCUDAHost(h_demsloper_y); FreeCUDADevice(d_demsloper_y);
FreeCUDAHost(h_demsloper_z); FreeCUDADevice(d_demsloper_z); //6
// 临时变量释放
FreeCUDAHost(h_R); FreeCUDADevice(d_R);
FreeCUDAHost(h_amp); FreeCUDADevice(d_amp);
FreeCUDAHost(h_demcls); FreeCUDADevice(d_demcls);
FreeCUDAHost(h_factorj); //FreeCUDADevice(d_factorj);
FreeCUDAHost(h_freqlist); //FreeCUDADevice(d_freqlist);
FreeCUDAHost(h_PRFEcho_real); FreeCUDADevice(d_PRFEcho_real);
FreeCUDAHost(h_PRFEcho_imag); FreeCUDADevice(d_PRFEcho_imag);
//FreeCUDAHost(h_phi); FreeCUDADevice(d_phi);
//FreeCUDAHost(h_real); FreeCUDADevice(d_real);
//FreeCUDAHost(h_imag); FreeCUDADevice(d_imag);
FreeCUDAHost(antpx); // 回收局部数据
FreeCUDAHost(antpy);
FreeCUDAHost(antpz);
FreeCUDAHost(antvx);
FreeCUDAHost(antvy);
FreeCUDAHost(antvz);
FreeCUDAHost(antdirectx);
FreeCUDAHost(antdirecty);
FreeCUDAHost(antdirectz);
FreeCUDAHost(antXaxisX);
FreeCUDAHost(antXaxisY);
FreeCUDAHost(antXaxisZ);
FreeCUDAHost(antYaxisX);
FreeCUDAHost(antYaxisY);
FreeCUDAHost(antYaxisZ);
FreeCUDAHost(antZaxisX);
FreeCUDAHost(antZaxisY);
FreeCUDAHost(antZaxisZ);
#endif
this->EchoSimulationData->saveToXml(); /** 内存释放***************************************************/
FreeCUDAHost(h_TantPattern);
FreeCUDAHost(h_RantPattern);
FreeCUDADevice(d_TantPattern);
FreeCUDADevice(d_RantPattern);
FreeCUDAHost(h_dem_x);
FreeCUDAHost(h_dem_y);
FreeCUDAHost(h_dem_z);
FreeCUDAHost(h_demsloper_x);
FreeCUDAHost(h_demsloper_y);
FreeCUDAHost(h_demsloper_z);
FreeCUDAHost(h_demsloper_z);
FreeCUDAHost(h_demcls);
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);
return ErrorCode::SUCCESS; return ErrorCode::SUCCESS;
} }