调整仿真代码
parent
6737300246
commit
1b8340d39a
|
|
@ -26,6 +26,22 @@
|
|||
#include <complex>
|
||||
#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
|
||||
|
||||
|
||||
|
|
@ -175,5 +191,12 @@ inline void PrintTime() {
|
|||
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
|
||||
|
|
@ -757,6 +757,64 @@ Eigen::MatrixXi gdalImage::getDatai(int start_row, int start_col, int rows_count
|
|||
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()
|
||||
{
|
||||
return this->gt;
|
||||
|
|
|
|||
|
|
@ -164,7 +164,8 @@ public: // 方法
|
|||
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::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 GDALDataType getDataType();
|
||||
virtual void saveImage(Eigen::MatrixXd, int start_row, int start_col, int band_ids);
|
||||
|
|
|
|||
|
|
@ -530,10 +530,10 @@ __global__ void CUDAKernel_PRF_GeneratorEcho(float* Rarr, float* ampArr,
|
|||
for (long j = 0; j < GPU_SHARE_STEP; j++) {
|
||||
dataid = j * BLOCK_SIZE + tid; //
|
||||
temp_phi = s_R[dataid] * factorjTemp;
|
||||
temp_amp = 1; s_Amp[dataid];
|
||||
temp_amp = s_Amp[dataid];
|
||||
|
||||
temp_real += 1; //temp_amp* cosf(temp_phi);
|
||||
temp_imag += 1;// temp_amp* sinf(temp_phi);
|
||||
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); // 更新虚部
|
||||
|
|
@ -544,84 +544,417 @@ __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,
|
||||
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,
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
|
||||
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,
|
||||
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];
|
||||
|
||||
double RstR = sqrt(RstX * RstX + RstY * RstY + RstZ * RstZ); // 矢量距离
|
||||
if (RstR<NearR || RstR>FarR) {
|
||||
d_temp_R[idx] = 0;
|
||||
d_temp_amps[idx] = 0;
|
||||
}
|
||||
else {
|
||||
double slopeX = demSlopeX[aposId];
|
||||
double slopeY = demSlopeY[aposId];
|
||||
double slopeZ = demSlopeZ[aposId];
|
||||
|
||||
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[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,
|
||||
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 , TargetID: %d / %d", spid+ BlockPRFCount, sTi, sTi+ BlockTarlist);
|
||||
|
||||
|
||||
}
|
||||
|
||||
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);
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
/** 对外封装接口 *******************************************************************************************************/
|
||||
|
||||
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 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* out_echoReal, float* out_echoImag,// 输出回波
|
||||
float* temp_R, float* temp_amp
|
||||
//, double* temp_phi, double* temp_real, double* temp_imag// 临时变量
|
||||
) {
|
||||
|
||||
long blocknum = 0;
|
||||
long pixelcount=TargetPixelNumber;
|
||||
int numBlocks = 0;
|
||||
for(long pid=0;pid<PRFCount;pid++){
|
||||
numBlocks = (TargetPixelNumber + BLOCK_SIZE - 1) / BLOCK_SIZE; // 根据 pixelcount 计算网格大小
|
||||
CUDAKernel_RFPC_Computer_R_Gain<<<numBlocks , BLOCK_SIZE >>>(
|
||||
antX[startpid+pid], antY[startpid+pid], antZ[startpid+pid],
|
||||
targetX, targetY, targetZ, TargetPixelNumber,
|
||||
demCls,
|
||||
demSlopeX, demSlopeY, demSlopeZ,
|
||||
antXaxisX[startpid+pid], antXaxisY[startpid+pid], antXaxisZ[startpid+pid],
|
||||
antYaxisX[startpid+pid], antYaxisY[startpid+pid], antYaxisZ[startpid+pid],
|
||||
antZaxisX[startpid+pid], antZaxisY[startpid+pid], antZaxisZ[startpid+pid],
|
||||
antDirectX[startpid + pid], antDirectY[startpid + pid], antDirectZ[startpid + pid],
|
||||
Pt,// 增益后发射能量
|
||||
refPhaseRange,
|
||||
TransAntpattern,
|
||||
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
|
||||
ReceiveAntpattern,
|
||||
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
|
||||
NearR, FarR,
|
||||
sigma0Paramslist, sigmaparamslistlen,
|
||||
//factorj, freqnum,
|
||||
temp_R, // 输出距离
|
||||
temp_amp
|
||||
//out_echoReal, out_echoImag, pid // 输出振幅
|
||||
);
|
||||
#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__
|
||||
blocknum = (pixelcount + GPU_SHARE_STEP - 1) / GPU_SHARE_STEP;
|
||||
numBlocks = (blocknum + BLOCK_SIZE - 1) / BLOCK_SIZE; // 网格数量
|
||||
CUDAKernel_PRF_GeneratorEcho << <numBlocks, BLOCK_SIZE >> >
|
||||
(temp_R, temp_amp, pixelcount,
|
||||
f0,dfreq,freqnum,
|
||||
out_echoReal, out_echoImag, pid);
|
||||
#ifdef __CUDADEBUG__
|
||||
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....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
}
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
}
|
||||
|
||||
//
|
||||
//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 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* out_echoReal, float* out_echoImag,// 输出回波
|
||||
// float* temp_R, float* temp_amp
|
||||
// //, double* temp_phi, double* temp_real, double* temp_imag// 临时变量
|
||||
//) {
|
||||
//
|
||||
// long blocknum = 0;
|
||||
// long pixelcount=TargetPixelNumber;
|
||||
// int numBlocks = 0;
|
||||
// for(long pid=0;pid<PRFCount;pid++){
|
||||
// numBlocks = (TargetPixelNumber + BLOCK_SIZE - 1) / BLOCK_SIZE; // 根据 pixelcount 计算网格大小
|
||||
// CUDAKernel_RFPC_Computer_R_Gain<<<numBlocks , BLOCK_SIZE >>>(
|
||||
// antX[startpid+pid], antY[startpid+pid], antZ[startpid+pid],
|
||||
// targetX, targetY, targetZ, TargetPixelNumber,
|
||||
// demCls,
|
||||
// demSlopeX, demSlopeY, demSlopeZ,
|
||||
// antXaxisX[startpid+pid], antXaxisY[startpid+pid], antXaxisZ[startpid+pid],
|
||||
// antYaxisX[startpid+pid], antYaxisY[startpid+pid], antYaxisZ[startpid+pid],
|
||||
// antZaxisX[startpid+pid], antZaxisY[startpid+pid], antZaxisZ[startpid+pid],
|
||||
// antDirectX[startpid + pid], antDirectY[startpid + pid], antDirectZ[startpid + pid],
|
||||
// Pt,// 增益后发射能量
|
||||
// refPhaseRange,
|
||||
// TransAntpattern,
|
||||
// Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
|
||||
// ReceiveAntpattern,
|
||||
// Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
|
||||
// NearR, FarR,
|
||||
// sigma0Paramslist, sigmaparamslistlen,
|
||||
// //factorj, freqnum,
|
||||
// temp_R, // 输出距离
|
||||
// temp_amp
|
||||
// //out_echoReal, out_echoImag, pid // 输出振幅
|
||||
// );
|
||||
//#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__
|
||||
// blocknum = (pixelcount + GPU_SHARE_STEP - 1) / GPU_SHARE_STEP;
|
||||
// numBlocks = (blocknum + BLOCK_SIZE - 1) / BLOCK_SIZE; // 网格数量
|
||||
// CUDAKernel_PRF_GeneratorEcho << <numBlocks, BLOCK_SIZE >> >
|
||||
// (temp_R, temp_amp, pixelcount,
|
||||
// f0,dfreq,freqnum,
|
||||
// out_echoReal, out_echoImag, pid);
|
||||
//#ifdef __CUDADEBUG__
|
||||
// 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....
|
||||
// }
|
||||
//#endif // __CUDADEBUG__
|
||||
// }
|
||||
// cudaDeviceSynchronize();
|
||||
//
|
||||
//}
|
||||
//
|
||||
|
||||
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -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, // 天线的坐标
|
||||
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, // 地表坡度矢量
|
||||
|
||||
extern "C" void CUDA_RFPC_MainProcess(
|
||||
// 天线
|
||||
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 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* 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* 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
|
||||
|
||||
|
||||
|
|
|
|||
|
|
@ -16,6 +16,45 @@
|
|||
#define BLOCK_DIM 1024
|
||||
#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)
|
||||
{
|
||||
|
|
@ -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;
|
||||
if (idx < len) {
|
||||
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) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
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
|
||||
|
|
|
|||
|
|
@ -17,8 +17,11 @@
|
|||
|
||||
#define BLOCK_SIZE 256
|
||||
|
||||
#define SHAREMEMORY_BYTE 49152
|
||||
#define SHAREMEMORY_FLOAT_HALF 6144
|
||||
|
||||
|
||||
// 打印GPU参数
|
||||
void printDeviceInfo(int deviceId);
|
||||
|
||||
|
||||
enum LAMPGPUDATETYPE {
|
||||
|
|
@ -55,15 +58,7 @@ 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);
|
||||
|
|
@ -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_D_sin(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 CUDADSin(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
|
||||
|
|
|
|||
|
|
@ -249,6 +249,7 @@
|
|||
<ClInclude Include="BaseToolbox\WGS84_J2000.h">
|
||||
<Filter>BaseToolbox</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="GPUTool\GPUGarbage.cuh" />
|
||||
</ItemGroup>
|
||||
<ItemGroup>
|
||||
<QtMoc Include="QMergeRasterProcessDialog.h">
|
||||
|
|
@ -263,9 +264,6 @@
|
|||
<QtMoc Include="QSimulationRFPCGUI.h">
|
||||
<Filter>Header Files</Filter>
|
||||
</QtMoc>
|
||||
<QtMoc Include="BaseTool\QToolProcessBarDialog.h">
|
||||
<Filter>Header Files</Filter>
|
||||
</QtMoc>
|
||||
<QtMoc Include="SimulationSAR\QToolAbstract.h">
|
||||
<Filter>Header Files</Filter>
|
||||
</QtMoc>
|
||||
|
|
@ -293,6 +291,9 @@
|
|||
<QtMoc Include="BaseToolbox\QOrthSlrRaster.h">
|
||||
<Filter>BaseToolbox</Filter>
|
||||
</QtMoc>
|
||||
<QtMoc Include="BaseTool\QToolProcessBarDialog.h">
|
||||
<Filter>BaseTool</Filter>
|
||||
</QtMoc>
|
||||
</ItemGroup>
|
||||
<ItemGroup>
|
||||
<QtUic Include="QMergeRasterProcessDialog.ui">
|
||||
|
|
@ -342,6 +343,7 @@
|
|||
<CudaCompile Include="GPUTool\GPUTool.cu">
|
||||
<Filter>GPUTool</Filter>
|
||||
</CudaCompile>
|
||||
<CudaCompile Include="GPUTool\GPUGarbage.cu" />
|
||||
</ItemGroup>
|
||||
<ItemGroup>
|
||||
<None Include="cpp.hint" />
|
||||
|
|
|
|||
|
|
@ -25,7 +25,7 @@
|
|||
#include <cuda_runtime.h>
|
||||
#include <cublas_v2.h>
|
||||
#endif // __CUDANVCC___
|
||||
#include <Imageshow/ImageShowDialogClass.h>
|
||||
//#include <Imageshow/ImageShowDialogClass.h>
|
||||
|
||||
|
||||
|
||||
|
|
@ -329,131 +329,149 @@ void RFPCProcessMain(long num_thread,
|
|||
}
|
||||
|
||||
|
||||
ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
|
||||
{
|
||||
double widthSpace = LIGHTSPEED / 2 / this->TaskSetting->getFs();
|
||||
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;
|
||||
ErrorCode RFPCProcessCls::RFPCMainProcess_GPU() {
|
||||
/** 内存分配***************************************************/
|
||||
long TargetMemoryMB = 500;
|
||||
|
||||
|
||||
/** 参数区域***************************************************/
|
||||
QVector<double> freqlist = this->TaskSetting->getFreqList();
|
||||
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 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 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 GainAntLen = -3;// -3dB 为天线半径
|
||||
long pluseCount = this->PluseCount;
|
||||
|
||||
double lamda = this->TaskSetting->getCenterLamda(); // 波长
|
||||
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___
|
||||
QMessageBox::information(this, u8"程序提示", u8"请确定安装了CUDA库");
|
||||
#else
|
||||
|
||||
|
||||
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);
|
||||
antpy = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antpz = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antvx = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antvy = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antvz = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antdirectx = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antdirecty = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antdirectz = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antXaxisX = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antXaxisY = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antXaxisZ = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antYaxisX = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antYaxisY = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antYaxisZ = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antZaxisX = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antZaxisY = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
antZaxisZ = (double*)mallocCUDAHost(sizeof(double) * pluseCount);
|
||||
|
||||
|
||||
for (long tempprfid = 0; tempprfid < pluseCount; tempprfid++) {
|
||||
long prfid = tempprfid;
|
||||
antpx[tempprfid] = sateOirbtNodes[prfid].Px;
|
||||
antpy[tempprfid] = sateOirbtNodes[prfid].Py;
|
||||
antpz[tempprfid] = sateOirbtNodes[prfid].Pz;
|
||||
antvx[tempprfid] = sateOirbtNodes[prfid].Vx;
|
||||
antvy[tempprfid] = sateOirbtNodes[prfid].Vy;
|
||||
antvz[tempprfid] = sateOirbtNodes[prfid].Vz; //6
|
||||
antdirectx[tempprfid] = sateOirbtNodes[prfid].AntDirecX;
|
||||
antdirecty[tempprfid] = sateOirbtNodes[prfid].AntDirecY;
|
||||
antdirectz[tempprfid] = sateOirbtNodes[prfid].AntDirecZ; // 9 天线指向
|
||||
antXaxisX[tempprfid] = sateOirbtNodes[prfid].AntXaxisX;
|
||||
antXaxisY[tempprfid] = sateOirbtNodes[prfid].AntXaxisY;
|
||||
antXaxisZ[tempprfid] = sateOirbtNodes[prfid].AntXaxisZ;//12 天线坐标系
|
||||
antYaxisX[tempprfid] = sateOirbtNodes[prfid].AntYaxisX;
|
||||
antYaxisY[tempprfid] = sateOirbtNodes[prfid].AntYaxisY;
|
||||
antYaxisZ[tempprfid] = sateOirbtNodes[prfid].AntYaxisZ;//15
|
||||
antZaxisX[tempprfid] = sateOirbtNodes[prfid].AntZaxisX;
|
||||
antZaxisY[tempprfid] = sateOirbtNodes[prfid].AntZaxisY;
|
||||
antZaxisZ[tempprfid] = sateOirbtNodes[prfid].AntZaxisZ;//18
|
||||
}
|
||||
}
|
||||
|
||||
// RFPC CUDA版本
|
||||
if (pluseCount * 4 * 18 > Memory1MB * 100) {
|
||||
long max = Memory1MB * 100 / 4 / 20 / PluseCount;
|
||||
QMessageBox::warning(nullptr, u8"仿真场景太大了", u8"当前频点数下，脉冲数量最多为：" + QString::number(max));
|
||||
}
|
||||
|
||||
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
|
||||
|
||||
blokline = Memory1MB / 8 / demCol * 500;
|
||||
blokline = blokline < 1 ? 1 : blokline;
|
||||
bool bloklineflag = false;
|
||||
PatternImageDesc TantPatternDesc = {};
|
||||
double* h_TantPattern=nullptr;
|
||||
double* d_TantPattern=nullptr;
|
||||
double prf_time = 0;
|
||||
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);
|
||||
|
||||
|
||||
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(); // 接收天线方向图
|
||||
|
||||
POLARTYPEENUM polartype = this->TaskSetting->getPolarType();
|
||||
PatternImageDesc TantPatternDesc = {};
|
||||
double* h_TantPattern = nullptr;
|
||||
double* d_TantPattern = nullptr;
|
||||
|
||||
{
|
||||
// 处理发射天线方向图
|
||||
double Tminphi = TransformPattern->getMinPhi();
|
||||
|
|
@ -494,11 +512,10 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
|
|||
TantPatternDesc.phinum = Tphinum;
|
||||
TantPatternDesc.thetanum = Tthetanum;
|
||||
}
|
||||
|
||||
|
||||
PatternImageDesc RantPatternDesc = {};
|
||||
double* h_RantPattern = nullptr;
|
||||
double* d_RantPattern = nullptr;
|
||||
|
||||
{
|
||||
// 处理接收天线方向图
|
||||
double Rminphi = ReceivePattern->getMinPhi();
|
||||
|
|
@ -539,12 +556,25 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
|
|||
RantPatternDesc.phinum = Rphinum;
|
||||
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;
|
||||
long clamapid = 0;
|
||||
long startline = 0;
|
||||
|
||||
|
||||
{
|
||||
long blokline = getBlockRows(2e4, demCol, sizeof(double));
|
||||
for (startline = 0; startline < demRow; startline = startline + blokline) {
|
||||
Eigen::MatrixXd clsland = demlandcls.getData(startline, 0, blokline, demlandcls.width, 1);
|
||||
long clsrows = clsland.rows();
|
||||
|
|
@ -599,269 +629,173 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU()
|
|||
|
||||
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__
|
||||
blokline = 1;
|
||||
blockwidth = 1;
|
||||
#endif
|
||||
// 地面 XYZ
|
||||
Eigen::MatrixXd dem_x = demxyz.getData(0, 0, blokline, blockwidth, 1); // 地面坐标
|
||||
long tempDemRows = dem_x.rows();
|
||||
long tempDemCols = dem_x.cols();
|
||||
|
||||
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))];
|
||||
h_dem_x[i * temp_dem_col + j] = -2028380.6250000; double(dem_x(i, j));
|
||||
h_dem_y[i * temp_dem_col + j] = 4139373.250000; double(dem_y(i, j));
|
||||
h_dem_z[i * temp_dem_col + j] = 4393382.500000; double(dem_z(i, j));
|
||||
h_demsloper_x[i * temp_dem_col + j] = 4393382.500000; double(demsloper_x(i, j));
|
||||
h_demsloper_y[i * temp_dem_col + j] = 446.923950; double(demsloper_y(i, j));
|
||||
h_demsloper_z[i * temp_dem_col + j] = -219.002213; double(demsloper_z(i, j));
|
||||
h_demcls[i * temp_dem_col + j] = clamap[80];// clamap[long(landcover(i, j))];
|
||||
#else
|
||||
h_dem_x[i * blockwidth + j] = double(dem_x(i, j));
|
||||
h_dem_y[i * blockwidth + j] = double(dem_y(i, j));
|
||||
h_dem_z[i * blockwidth + j] = double(dem_z(i, j));
|
||||
h_demsloper_x[i * blockwidth + j] = double(demsloper_x(i, j));
|
||||
h_demsloper_y[i * blockwidth + j] = double(demsloper_y(i, j));
|
||||
h_demsloper_z[i * blockwidth + j] = double(demsloper_z(i, j));
|
||||
h_demcls[i * blockwidth + j] = clamap[long(landcover(i, j))];
|
||||
h_dem_x[i * temp_dem_col + j] = double(dem_x(i, j));
|
||||
h_dem_y[i * temp_dem_col + j] = double(dem_y(i, j));
|
||||
h_dem_z[i * temp_dem_col + j] = double(dem_z(i, j));
|
||||
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))];
|
||||
#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]);
|
||||
std::shared_ptr<double> h_temp_R(new double[PluseCount], delArrPtr);
|
||||
#endif // __PRFDEBUG__
|
||||
|
||||
long pixelcount = newblokline * tempDemCols;
|
||||
long startprfid = 0;
|
||||
for (startprfid = 0; startprfid < pluseCount; startprfid = startprfid + echoblockline) {
|
||||
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;
|
||||
long templine = startprfid + echoblockline < PluseCount ? echoblockline : PluseCount - startprfid;
|
||||
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,// 增益后发射能量
|
||||
// 分块处理
|
||||
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,
|
||||
f0,dfreq,
|
||||
Pt,
|
||||
refphaseRange,
|
||||
// 天线方向图
|
||||
d_TantPattern,
|
||||
TantPatternDesc.startTheta, TantPatternDesc.startPhi, TantPatternDesc.dtheta,TantPatternDesc.dphi,TantPatternDesc.thetanum,TantPatternDesc.phinum,
|
||||
d_RantPattern,
|
||||
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,
|
||||
|
||||
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_PRFEcho_real, d_PRFEcho_imag,// 输出回波
|
||||
d_R, d_amp
|
||||
//, d_phi, d_real, d_imag// 临时变量
|
||||
|
||||
h_echo_block_real, h_echo_block_imag// 输出回波
|
||||
);
|
||||
|
||||
DeviceToHost(h_PRFEcho_real, d_PRFEcho_real, sizeof(float) * echoblockline * PlusePoint);
|
||||
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);
|
||||
PRINT("dem : %d - %d / %d , echo: %d -%d / %d", startline, startline+ temp_dem_row, demRow, sprfid, sprfid+ PRF_len, PRFCount);
|
||||
}
|
||||
|
||||
#ifdef __PRFDEBUG__ && __PRFDEBUG_PRFINF__
|
||||
break;
|
||||
#endif // __PRFDEBUG__
|
||||
|
||||
for (long ii = 0; ii < PRF_len; ii++) {
|
||||
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;
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue