增加 GPU版本的 间接定位法
parent
b8cf53cf20
commit
f426ce1f1d
|
|
@ -5,3 +5,9 @@ BASECONSTVARIABLEAPI void PrintMsgToQDebug(char* msg)
|
||||||
qDebug() << QString(msg);
|
qDebug() << QString(msg);
|
||||||
return ;
|
return ;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
BASECONSTVARIABLEAPI void PrintTipMsgToQDebug(const char* tip, const char* msg)
|
||||||
|
{
|
||||||
|
qDebug() <<QString(tip)<<"\t:\t" << QString(msg);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
|
||||||
|
|
@ -4,6 +4,7 @@
|
||||||
#include "BaseConstVariable.h"
|
#include "BaseConstVariable.h"
|
||||||
|
|
||||||
extern "C" BASECONSTVARIABLEAPI void PrintMsgToQDebug(char* msg);
|
extern "C" BASECONSTVARIABLEAPI void PrintMsgToQDebug(char* msg);
|
||||||
|
extern "C" BASECONSTVARIABLEAPI void PrintTipMsgToQDebug(const char* tip, const char* msg);
|
||||||
#endif // !PRINTMSGTOQDEBUG_H_
|
#endif // !PRINTMSGTOQDEBUG_H_
|
||||||
|
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -25,6 +25,11 @@
|
||||||
<ClInclude Include="GPUTool\GPUBaseLibAPI.h" />
|
<ClInclude Include="GPUTool\GPUBaseLibAPI.h" />
|
||||||
<CudaCompile Include="GPUTool\GPUTool.cuh" />
|
<CudaCompile Include="GPUTool\GPUTool.cuh" />
|
||||||
</ItemGroup>
|
</ItemGroup>
|
||||||
|
<ItemGroup>
|
||||||
|
<ProjectReference Include="..\BaseCommonLibrary\BaseCommonLibrary.vcxproj">
|
||||||
|
<Project>{872ecd6f-30e3-4a1b-b17c-15e87d373ff6}</Project>
|
||||||
|
</ProjectReference>
|
||||||
|
</ItemGroup>
|
||||||
<PropertyGroup Label="Globals">
|
<PropertyGroup Label="Globals">
|
||||||
<VCProjectVersion>17.0</VCProjectVersion>
|
<VCProjectVersion>17.0</VCProjectVersion>
|
||||||
<Keyword>Win32Proj</Keyword>
|
<Keyword>Win32Proj</Keyword>
|
||||||
|
|
|
||||||
|
|
@ -11,6 +11,7 @@
|
||||||
#include <chrono>
|
#include <chrono>
|
||||||
#include "BaseConstVariable.h"
|
#include "BaseConstVariable.h"
|
||||||
#include "GPUTool.cuh"
|
#include "GPUTool.cuh"
|
||||||
|
#include "PrintMsgToQDebug.h"
|
||||||
|
|
||||||
#ifdef __CUDANVCC___
|
#ifdef __CUDANVCC___
|
||||||
#define BLOCK_DIM 1024
|
#define BLOCK_DIM 1024
|
||||||
|
|
@ -486,7 +487,8 @@ void PrintLasterError(const char* s)
|
||||||
{
|
{
|
||||||
cudaError_t err = cudaGetLastError();
|
cudaError_t err = cudaGetLastError();
|
||||||
if (err != cudaSuccess) {
|
if (err != cudaSuccess) {
|
||||||
printf("%s: %s\n", s, cudaGetErrorString(err));
|
//printf("%s: %s\n", s, cudaGetErrorString(err));
|
||||||
|
PrintTipMsgToQDebug(s, cudaGetErrorString(err));
|
||||||
exit(2);
|
exit(2);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
|
||||||
|
|
@ -5,7 +5,7 @@
|
||||||
#include "FileOperator.h"
|
#include "FileOperator.h"
|
||||||
#include "BaseConstVariable.h"
|
#include "BaseConstVariable.h"
|
||||||
#include "GPUTool.cuh"
|
#include "GPUTool.cuh"
|
||||||
|
#include "LookTableSimulationComputer.cuh"
|
||||||
#include <QDebug>
|
#include <QDebug>
|
||||||
|
|
||||||
|
|
||||||
|
|
@ -259,6 +259,7 @@ namespace LookTableSimualtionMainProcessSpace {
|
||||||
|
|
||||||
int rowcount = 0;
|
int rowcount = 0;
|
||||||
int colcount = 0;
|
int colcount = 0;
|
||||||
|
double fact_lamda = 1 / lamda;
|
||||||
for (long rid = 0; rid < demimg.height; rid = rid + GPUMemoryline) {
|
for (long rid = 0; rid < demimg.height; rid = rid + GPUMemoryline) {
|
||||||
rowcount = GPUMemoryline;
|
rowcount = GPUMemoryline;
|
||||||
colcount = demimg.width;
|
colcount = demimg.width;
|
||||||
|
|
@ -279,7 +280,20 @@ namespace LookTableSimualtionMainProcessSpace {
|
||||||
HostToDevice(host_demZ.get(), device_demZ.get(), sizeof(double)* GPUMemoryline* demimg.width);
|
HostToDevice(host_demZ.get(), device_demZ.get(), sizeof(double)* GPUMemoryline* demimg.width);
|
||||||
|
|
||||||
|
|
||||||
|
RDProcess_dopplerGPU(
|
||||||
|
device_demX.get(), device_demY.get(), device_demZ.get(),
|
||||||
|
device_Rid.get(), device_Cid.get(),
|
||||||
|
rowcount, colcount,
|
||||||
|
Xp0, Yp0, Zp0, Xv0, Yv0, Zv0,
|
||||||
|
Xp1, Yp1, Zp1, Xv1, Yv1, Zv1,
|
||||||
|
Xp2, Yp2, Zp2, Xv2, Yv2, Zv2,
|
||||||
|
Xp3, Yp3, Zp3, Xv3, Yv3, Zv3,
|
||||||
|
Xp4, Yp4, Zp4, Xv4, Yv4, Zv4,
|
||||||
|
Xp5, Yp5, Zp5, Xv5, Yv5, Zv5,
|
||||||
|
dopplerRefrenceTime, r0, r1, r2, r3, r4,
|
||||||
|
starttime, nearRange, farRange,
|
||||||
|
PRF, Fs,
|
||||||
|
fact_lamda);
|
||||||
|
|
||||||
|
|
||||||
// GPU -> ÄÚ´æ
|
// GPU -> ÄÚ´æ
|
||||||
|
|
|
||||||
|
|
@ -20,9 +20,9 @@ extern __device__ __host__ double getNumberDopplerCenterRate(double R, double r
|
||||||
return dopplerCenterRate;
|
return dopplerCenterRate;
|
||||||
}
|
}
|
||||||
|
|
||||||
__device__ __host__ double getDopplerCenterRate(double Rx, double Ry, double Rz, double Vx, double Vy, double Vz, double lamda)
|
__device__ __host__ double getDopplerCenterRate(double Rx, double Ry, double Rz, double Vx, double Vy, double Vz, double fact_lamda)
|
||||||
{
|
{
|
||||||
return __device__ __host__ double();
|
return -2*(Rx*Vx+Ry*Vy+Rz*Vz)* fact_lamda;
|
||||||
}
|
}
|
||||||
|
|
||||||
__device__ __host__ double getPolyfitNumber(double x, double a0, double a1, double a2, double a3, double a4, double a5)
|
__device__ __host__ double getPolyfitNumber(double x, double a0, double a1, double a2, double a3, double a4, double a5)
|
||||||
|
|
@ -31,6 +31,134 @@ __device__ __host__ double getPolyfitNumber(double x, double a0, double a1, doub
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
__global__ void Kernel_RDProcess_doppler(
|
||||||
|
double* demX, double* demY, double* demZ,
|
||||||
|
float* outRidx, float* outCidx,
|
||||||
|
long pixelcount,
|
||||||
|
double Xp0, double Yp0, double Zp0, double Xv0, double Yv0, double Zv0,
|
||||||
|
double Xp1, double Yp1, double Zp1, double Xv1, double Yv1, double Zv1,
|
||||||
|
double Xp2, double Yp2, double Zp2, double Xv2, double Yv2, double Zv2,
|
||||||
|
double Xp3, double Yp3, double Zp3, double Xv3, double Yv3, double Zv3,
|
||||||
|
double Xp4, double Yp4, double Zp4, double Xv4, double Yv4, double Zv4,
|
||||||
|
double Xp5, double Yp5, double Zp5, double Xv5, double Yv5, double Zv5,
|
||||||
|
double reftime, double r0, double r1, double r2, double r3, double r4,
|
||||||
|
double starttime, double nearRange, double farRange,
|
||||||
|
double PRF, double Fs,
|
||||||
|
double fact_lamda
|
||||||
|
) {
|
||||||
|
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||||
|
if (idx < pixelcount) {
|
||||||
|
double demx = demX[idx];
|
||||||
|
double demy = demY[idx];
|
||||||
|
double demz = demZ[idx];
|
||||||
|
float Rd_r = -1;
|
||||||
|
float Rd_c = -1;
|
||||||
|
double dt = 1 / PRF / 1000;
|
||||||
|
double Spx = 0, Spy = 0, Spz = 0, Svx = 0, Svy = 0, Svz = 0;
|
||||||
|
double Rx = 0, Ry = 0, Rz = 0,R = 0;
|
||||||
|
|
||||||
|
double dp1=0, dpn1=0, dp2=0, dpn2=0;
|
||||||
|
|
||||||
|
double ti = 0;
|
||||||
|
double inct = 0;
|
||||||
|
for (long i = 0; i < 10000; i++) { // ×î´óľü´úˇśÎ§
|
||||||
|
Spx = getPolyfitNumber(ti+dt, Xp0, Xp1, Xp2, Xp3, Xp4, Xp5);
|
||||||
|
Spy = getPolyfitNumber(ti+dt, Yp0, Yp1, Yp2, Yp3, Yp4, Yp5);
|
||||||
|
Spz = getPolyfitNumber(ti+dt, Zp0, Zp1, Zp2, Zp3, Zp4, Zp5);
|
||||||
|
Svx = getPolyfitNumber(ti+dt, Xv0, Xv1, Xv2, Xv3, Xv4, Xv5);
|
||||||
|
Svy = getPolyfitNumber(ti+dt, Yv0, Yv1, Yv2, Yv3, Yv4, Yv5);
|
||||||
|
Svz = getPolyfitNumber(ti+dt, Zv0, Zv1, Zv2, Zv3, Zv4, Zv5);
|
||||||
|
|
||||||
|
Rx = Spx - demx;
|
||||||
|
Ry = Spy - demy;
|
||||||
|
Rz = Spz - demz;
|
||||||
|
R = sqrt(Rx * Rx + Ry * Ry + Rz * Rz);
|
||||||
|
Rx = Rx / R;
|
||||||
|
Ry = Ry / R;
|
||||||
|
Rz = Rz / R;
|
||||||
|
|
||||||
|
dp2 = getDopplerCenterRate(Rx, Ry, Rz, Svx, Svy, Svz, fact_lamda);
|
||||||
|
dpn2 = getNumberDopplerCenterRate(R, r0, r1, r2, r3, r4, reftime);
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
// ti
|
||||||
|
Spx = getPolyfitNumber(ti, Xp0, Xp1, Xp2, Xp3, Xp4, Xp5);
|
||||||
|
Spy = getPolyfitNumber(ti, Yp0, Yp1, Yp2, Yp3, Yp4, Yp5);
|
||||||
|
Spz = getPolyfitNumber(ti, Zp0, Zp1, Zp2, Zp3, Zp4, Zp5);
|
||||||
|
Svx = getPolyfitNumber(ti, Xv0, Xv1, Xv2, Xv3, Xv4, Xv5);
|
||||||
|
Svy = getPolyfitNumber(ti, Yv0, Yv1, Yv2, Yv3, Yv4, Yv5);
|
||||||
|
Svz = getPolyfitNumber(ti, Zv0, Zv1, Zv2, Zv3, Zv4, Zv5);
|
||||||
|
|
||||||
|
Rx = Spx - demx;
|
||||||
|
Ry = Spy - demy;
|
||||||
|
Rz = Spz - demz;
|
||||||
|
R = sqrt(Rx * Rx + Ry * Ry + Rz * Rz);
|
||||||
|
Rx = Rx / R;
|
||||||
|
Ry = Ry / R;
|
||||||
|
Rz = Rz / R;
|
||||||
|
|
||||||
|
dp1=getDopplerCenterRate(Rx, Ry, Rz, Svx, Svy, Svz, fact_lamda);
|
||||||
|
dpn1 = getNumberDopplerCenterRate(R, r0, r1, r2, r3, r4, reftime);
|
||||||
|
|
||||||
|
// iter
|
||||||
|
inct = dt * (dp2 - dpn1) / (dp1 - dp2);
|
||||||
|
|
||||||
|
if (abs(inct) <= dt ||(abs(dp1 - dp2) < 1e-9)) {
|
||||||
|
Rd_r = (ti - starttime) * PRF;
|
||||||
|
Rd_c = 2 * (R - nearRange) / LIGHTSPEED * Fs;
|
||||||
|
outRidx[idx] = Rd_r;
|
||||||
|
outCidx[idx] = Rd_c;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
ti = ti - inct;
|
||||||
|
}
|
||||||
|
outRidx[idx] = -1;
|
||||||
|
outCidx[idx] = -1;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
void RDProcess_dopplerGPU(
|
||||||
|
double* demX, double* demY, double* demZ,
|
||||||
|
float* outRidx, float* outCidx,
|
||||||
|
long rowcount, long colcount,
|
||||||
|
double Xp0, double Yp0, double Zp0, double Xv0, double Yv0, double Zv0,
|
||||||
|
double Xp1, double Yp1, double Zp1, double Xv1, double Yv1, double Zv1,
|
||||||
|
double Xp2, double Yp2, double Zp2, double Xv2, double Yv2, double Zv2,
|
||||||
|
double Xp3, double Yp3, double Zp3, double Xv3, double Yv3, double Zv3,
|
||||||
|
double Xp4, double Yp4, double Zp4, double Xv4, double Yv4, double Zv4,
|
||||||
|
double Xp5, double Yp5, double Zp5, double Xv5, double Yv5, double Zv5,
|
||||||
|
double reftime, double r0, double r1, double r2, double r3, double r4,
|
||||||
|
double starttime, double nearRange, double farRange,
|
||||||
|
double PRF, double Fs,
|
||||||
|
double fact_lamda)
|
||||||
|
{
|
||||||
|
long pixelcount = rowcount * colcount;
|
||||||
|
int numBlocks = (pixelcount + BLOCK_SIZE - 1) / BLOCK_SIZE;
|
||||||
|
Kernel_RDProcess_doppler << <numBlocks, BLOCK_SIZE >> > (
|
||||||
|
demX, demY, demZ,
|
||||||
|
outRidx, outCidx,
|
||||||
|
pixelcount,
|
||||||
|
Xp0, Yp0, Zp0, Xv0, Yv0, Zv0,
|
||||||
|
Xp1, Yp1, Zp1, Xv1, Yv1, Zv1,
|
||||||
|
Xp2, Yp2, Zp2, Xv2, Yv2, Zv2,
|
||||||
|
Xp3, Yp3, Zp3, Xv3, Yv3, Zv3,
|
||||||
|
Xp4, Yp4, Zp4, Xv4, Yv4, Zv4,
|
||||||
|
Xp5, Yp5, Zp5, Xv5, Yv5, Zv5,
|
||||||
|
reftime, r0, r1, r2, r3, r4,
|
||||||
|
starttime,nearRange, farRange,
|
||||||
|
PRF, Fs,
|
||||||
|
fact_lamda
|
||||||
|
);
|
||||||
|
PrintLasterError("RD with doppler function");
|
||||||
|
cudaDeviceSynchronize();
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -10,7 +10,7 @@
|
||||||
|
|
||||||
extern __device__ __host__ double getNumberDopplerCenterRate(double R, double r0, double r1, double r2, double r3, double r4,double reftime);
|
extern __device__ __host__ double getNumberDopplerCenterRate(double R, double r0, double r1, double r2, double r3, double r4,double reftime);
|
||||||
|
|
||||||
extern __device__ __host__ double getDopplerCenterRate(double Rx,double Ry,double Rz,double Vx,double Vy,double Vz,double lamda);
|
extern __device__ __host__ double getDopplerCenterRate(double Rx,double Ry,double Rz,double Vx,double Vy,double Vz,double fact_lamda); //fact_lamda lamda 的倒数
|
||||||
|
|
||||||
extern __device__ __host__ double getPolyfitNumber(double x, double a0, double a1, double a2, double a3, double a4,double a5);
|
extern __device__ __host__ double getPolyfitNumber(double x, double a0, double a1, double a2, double a3, double a4,double a5);
|
||||||
|
|
||||||
|
|
@ -26,8 +26,25 @@ extern __device__ __host__ double getPolyfitNumber(double x, double a0, double a
|
||||||
|
|
||||||
|
|
||||||
//
|
//
|
||||||
extern "C" void RDProcess_doppler(
|
extern "C" void RDProcess_dopplerGPU(
|
||||||
|
double* demX, double* demY, double* demZ, // 处理入射坐标
|
||||||
|
float* outRidx,float* outCidx, // 输出 行列数
|
||||||
|
long rowcount,long colcount,
|
||||||
|
double Xp0 = 0, double Yp0 = 0,double Zp0 = 0,double Xv0 = 0,double Yv0 = 0,double Zv0 = 0, // 轨道参数
|
||||||
|
double Xp1 = 0, double Yp1 = 0,double Zp1 = 0,double Xv1 = 0,double Yv1 = 0,double Zv1 = 0,
|
||||||
|
double Xp2 = 0, double Yp2 = 0,double Zp2 = 0,double Xv2 = 0,double Yv2 = 0,double Zv2 = 0,
|
||||||
|
double Xp3 = 0, double Yp3 = 0,double Zp3 = 0,double Xv3 = 0,double Yv3 = 0,double Zv3 = 0,
|
||||||
|
double Xp4 = 0, double Yp4 = 0,double Zp4 = 0,double Xv4 = 0,double Yv4 = 0,double Zv4 = 0,
|
||||||
|
double Xp5 = 0, double Yp5 = 0,double Zp5 = 0,double Xv5 = 0,double Yv5 = 0,double Zv5 = 0,
|
||||||
|
double reftime=0, // 轨道参数
|
||||||
|
double r0 = 0, double r1 = 0, double r2 = 0, double r3 = 0, double r4 = 0,
|
||||||
|
|
||||||
|
double starttime,// 成像参数
|
||||||
|
double nearRange,
|
||||||
|
double farRange,
|
||||||
|
double PRF,
|
||||||
|
double Fs,
|
||||||
|
double fact_lamda // lamda 倒数
|
||||||
|
|
||||||
);
|
);
|
||||||
|
|
||||||
|
|
|
||||||
Loading…
Reference in New Issue