将计算振幅提前
parent
e3a6585adb
commit
19dbddd6b5
|
|
@ -19,7 +19,7 @@
|
|||
|
||||
|
||||
__device__ float GPU_getSigma0dB(CUDASigmaParam param, float theta) {//线性值
|
||||
float sigma= param.p1 + param.p2 * exp(-param.p3 * theta) + param.p4 * cos(param.p5 * theta + param.p6);
|
||||
float sigma = param.p1 + param.p2 * exp(-param.p3 * theta) + param.p4 * cos(param.p5 * theta + param.p6);
|
||||
return sigma;
|
||||
}
|
||||
|
||||
|
|
@ -104,7 +104,7 @@ __device__ CUDAVectorEllipsoidal GPU_SatelliteAntDirectNormal(
|
|||
if (isnan(PhiAnt)) {
|
||||
printf("V=[%f,%f,%f];norm=%f;thetaAnt=%f;phiAnt=%f;\n", Xant, Yant, Zant, Norm, ThetaAnt, PhiAnt);
|
||||
}
|
||||
|
||||
|
||||
|
||||
result.theta = ThetaAnt;
|
||||
result.phi = PhiAnt;
|
||||
|
|
@ -131,7 +131,7 @@ __device__ float GPU_BillerInterpAntPattern(float* antpattern,
|
|||
long nextTheta = lasttheta + 1;
|
||||
long lastphi = floorf(pphiid);
|
||||
long nextPhi = lastphi + 1;
|
||||
|
||||
|
||||
|
||||
if (lasttheta < 0 || nextTheta < 0 || lastphi < 0 || nextPhi < 0 ||
|
||||
lasttheta >= thetapoints || nextTheta >= thetapoints || lastphi >= phipoints || nextPhi >= phipoints)
|
||||
|
|
@ -175,7 +175,7 @@ __device__ cuComplex GPU_calculationEcho(float sigma0, float TransAnt, float Re
|
|||
float phi = (-4 * LAMP_CUDA_PI / lamda) * R;
|
||||
cuComplex echophi = make_cuComplex(0, phi);
|
||||
cuComplex echophiexp = cuCexpf(echophi);
|
||||
cuComplex echo=make_cuComplex(echophiexp.x * amp, echophiexp.y * amp);
|
||||
cuComplex echo = make_cuComplex(echophiexp.x * amp, echophiexp.y * amp);
|
||||
return echo;
|
||||
}
|
||||
|
||||
|
|
@ -189,9 +189,9 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
|
|||
, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
float Xst = -1 * RstX[idx]; // 卫星 --> 地面
|
||||
float Yst = -1 * RstY[idx];
|
||||
float Zst = -1 * RstZ[idx];
|
||||
float Xst = -1 * RstX[idx]; // 卫星 --> 地面
|
||||
float Yst = -1 * RstY[idx];
|
||||
float Zst = -1 * RstZ[idx];
|
||||
float AntXaxisX = antXaxisX;
|
||||
float AntXaxisY = antXaxisY;
|
||||
float AntXaxisZ = antXaxisZ;
|
||||
|
|
@ -203,7 +203,7 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
|
|||
float AntZaxisZ = antZaxisZ;
|
||||
|
||||
// 归一化
|
||||
float RstNorm = sqrtf(Xst * Xst + Yst * Yst + Zst * Zst);
|
||||
float RstNorm = sqrtf(Xst * Xst + Yst * Yst + Zst * Zst);
|
||||
float AntXaxisNorm = sqrtf(AntXaxisX * AntXaxisX + AntXaxisY * AntXaxisY + AntXaxisZ * AntXaxisZ);
|
||||
float AntYaxisNorm = sqrtf(AntYaxisX * AntYaxisX + AntYaxisY * AntYaxisY + AntYaxisZ * AntYaxisZ);
|
||||
float AntZaxisNorm = sqrtf(AntZaxisX * AntZaxisX + AntZaxisY * AntZaxisY + AntZaxisZ * AntZaxisZ);
|
||||
|
|
@ -222,11 +222,11 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
|
|||
float Zy = AntZaxisY / AntZaxisNorm;
|
||||
float Zz = AntZaxisZ / AntZaxisNorm;
|
||||
|
||||
float Xant = (Rx * Yy * Zz - Rx * Yz * Zy - Ry * Yx * Zz + Ry * Yz * Zx + Rz * Yx * Zy - Rz * Yy * Zx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
|
||||
float Yant = -(Rx * Xy * Zz - Rx * Xz * Zy - Ry * Xx * Zz + Ry * Xz * Zx + Rz * Xx * Zy - Rz * Xy * Zx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
|
||||
float Zant = (Rx * Xy * Yz - Rx * Xz * Yy - Ry * Xx * Yz + Ry * Xz * Yx + Rz * Xx * Yy - Rz * Xy * Yx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
|
||||
|
||||
|
||||
float Xant = (Rx * Yy * Zz - Rx * Yz * Zy - Ry * Yx * Zz + Ry * Yz * Zx + Rz * Yx * Zy - Rz * Yy * Zx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
|
||||
float Yant = -(Rx * Xy * Zz - Rx * Xz * Zy - Ry * Xx * Zz + Ry * Xz * Zx + Rz * Xx * Zy - Rz * Xy * Zx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
|
||||
float Zant = (Rx * Xy * Yz - Rx * Xz * Yy - Ry * Xx * Yz + Ry * Xz * Yx + Rz * Xx * Yy - Rz * Xy * Yx) / (Xx * Yy * Zz - Xx * Yz * Zy - Xy * Yx * Zz + Xy * Yz * Zx + Xz * Yx * Zy - Xz * Yy * Zx);
|
||||
|
||||
|
||||
// 计算theta 与 phi
|
||||
float Norm = sqrtf(Xant * Xant + Yant * Yant + Zant * Zant); // 计算 pho
|
||||
float ThetaAnt = acosf(Zant / Norm); // theta 与 Z轴的夹角
|
||||
|
|
@ -249,7 +249,7 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
|
|||
PhiAnt = PI + PhiAnt;
|
||||
}
|
||||
else {
|
||||
PhiAnt = -PI+PhiAnt ;
|
||||
PhiAnt = -PI + PhiAnt;
|
||||
}
|
||||
}
|
||||
else { // Xant>0 X 正轴
|
||||
|
|
@ -257,7 +257,7 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
|
|||
}
|
||||
|
||||
if (isnan(PhiAnt)) {
|
||||
printf("V=[%f,%f,%f];norm=%f;thetaAnt=%f;phiAnt=%f;\n", Xant, Yant, Zant,Norm, ThetaAnt, PhiAnt);
|
||||
printf("V=[%f,%f,%f];norm=%f;thetaAnt=%f;phiAnt=%f;\n", Xant, Yant, Zant, Norm, ThetaAnt, PhiAnt);
|
||||
}
|
||||
|
||||
//if (abs(ThetaAnt - 0) < PRECISIONTOLERANCE) {
|
||||
|
|
@ -266,8 +266,8 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
|
|||
//else {}
|
||||
|
||||
|
||||
thetaAnt[idx] = ThetaAnt*r2d;
|
||||
phiAnt[idx] = PhiAnt*r2d;
|
||||
thetaAnt[idx] = ThetaAnt * r2d;
|
||||
phiAnt[idx] = PhiAnt * r2d;
|
||||
//printf("Rst=[%f,%f,%f];AntXaxis = [%f, %f, %f];AntYaxis=[%f,%f,%f];AntZaxis=[%f,%f,%f];phiAnt=%f;thetaAnt=%f;\n", Xst, Yst, Zst
|
||||
// , AntXaxisX, AntXaxisY, AntXaxisZ
|
||||
// , AntYaxisX, AntYaxisY, AntYaxisZ
|
||||
|
|
@ -354,7 +354,7 @@ __global__ void CUDA_calculationEcho(float* sigma0, float* TransAnt, float* Reci
|
|||
// amp = 0;
|
||||
//}
|
||||
|
||||
cuComplex echo = make_cuComplex(echophiexp.x , echophiexp.y);
|
||||
cuComplex echo = make_cuComplex(echophiexp.x, echophiexp.y);
|
||||
echoArr[idx] = echo;
|
||||
FreqID[idx] = timeID;
|
||||
}
|
||||
|
|
@ -364,14 +364,14 @@ __global__ void CUDA_calculationEcho(float* sigma0, float* TransAnt, float* Reci
|
|||
__global__ void CUDA_AntPatternInterpGain(float* anttheta, float* antphi, float* gain,
|
||||
float* antpattern, float starttheta, float startphi, float dtheta, float dphi, int thetapoints, int phipoints, long len) {
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
|
||||
if (idx < len) {
|
||||
|
||||
float temptheta = anttheta[idx];
|
||||
float tempphi = antphi[idx];
|
||||
float antPatternGain = GPU_BillerInterpAntPattern(antpattern,
|
||||
float antPatternGain = GPU_BillerInterpAntPattern(antpattern,
|
||||
starttheta, startphi, dtheta, dphi, thetapoints, phipoints,
|
||||
temptheta, tempphi) ;
|
||||
temptheta, tempphi);
|
||||
gain[idx] = antPatternGain;
|
||||
}
|
||||
}
|
||||
|
|
@ -385,16 +385,16 @@ __global__ void CUDA_InterpSigma(
|
|||
long clsid = demcls[idx];
|
||||
float localangle = localanglearr[idx];
|
||||
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
|
||||
if (localangle < 0 || localangle >= LAMP_CUDA_PI/2) {
|
||||
if (localangle < 0 || localangle >= LAMP_CUDA_PI / 2) {
|
||||
sigmaAmp[idx] = 0;
|
||||
}
|
||||
else {}
|
||||
|
||||
if (abs(tempsigma.p1)< PRECISIONTOLERANCE&&
|
||||
if (abs(tempsigma.p1) < PRECISIONTOLERANCE &&
|
||||
abs(tempsigma.p2) < PRECISIONTOLERANCE &&
|
||||
abs(tempsigma.p3) < PRECISIONTOLERANCE &&
|
||||
abs(tempsigma.p4) < PRECISIONTOLERANCE&&
|
||||
abs(tempsigma.p5) < PRECISIONTOLERANCE&&
|
||||
abs(tempsigma.p4) < PRECISIONTOLERANCE &&
|
||||
abs(tempsigma.p5) < PRECISIONTOLERANCE &&
|
||||
abs(tempsigma.p6) < PRECISIONTOLERANCE
|
||||
) {
|
||||
sigmaAmp[idx] = 0;
|
||||
|
|
@ -409,8 +409,8 @@ __global__ void CUDA_InterpSigma(
|
|||
}
|
||||
|
||||
|
||||
__global__ void CUDA_CalculationEchoAmp(float* sigma0, float* TransAnt, float* ReciveAnt, float* R,
|
||||
float Pt,
|
||||
__global__ void CUDA_CalculationEchoAmp(float* sigma0, float* TransAnt, float* ReciveAnt, float* R,
|
||||
float Pt,
|
||||
float* ampArr, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
|
|
@ -434,7 +434,7 @@ __global__ void CUDA_CalculationEchoPhase(float* R, float lamda, float* phaseArr
|
|||
|
||||
|
||||
__global__ void CUDA_CombinationEchoAmpAndPhase(float* R,
|
||||
float* echoAmp,float* echoPhase,
|
||||
float* echoAmp, float* echoPhase,
|
||||
float nearRange, float Fs, long plusepoints,
|
||||
cuComplex* echo, long* FreqID, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
|
@ -451,7 +451,7 @@ __global__ void CUDA_CombinationEchoAmpAndPhase(float* R,
|
|||
timeID = 0;
|
||||
amp = 0;
|
||||
}
|
||||
cuComplex echotemp = make_cuComplex(echophiexp.x*amp, echophiexp.y*amp);
|
||||
cuComplex echotemp = make_cuComplex(echophiexp.x * amp, echophiexp.y * amp);
|
||||
echo[idx] = echotemp;
|
||||
FreqID[idx] = timeID;
|
||||
}
|
||||
|
|
@ -459,21 +459,22 @@ __global__ void CUDA_CombinationEchoAmpAndPhase(float* R,
|
|||
|
||||
|
||||
__global__ void CUDAKernel_RFPC_Caluation_R_Gain(
|
||||
float antX, float antY, float antZ,
|
||||
float* targetX, float* targetY, float* targetZ, long len,
|
||||
float* demSlopeX, float* demSlopeY, float* demSlopeZ,
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ,
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,
|
||||
float antDirectX, float antDirectY, float antDirectZ,
|
||||
float* TransAntpattern,
|
||||
float Transtarttheta, float Transstartphi, float Transdtheta, float Transdphi, int Transthetapoints, int Transphipoints,
|
||||
float* ReceiveAntpattern,
|
||||
float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,
|
||||
float NearR, float FarR,
|
||||
float* outR,
|
||||
float* outLocalAngle,
|
||||
float* AmpGain) {
|
||||
float antX, float antY, float antZ, // 天线的坐标
|
||||
float* targetX, float* targetY, float* targetZ, long len, // 地面坐标
|
||||
long* demCls,
|
||||
float* demSlopeX, float* demSlopeY, float* demSlopeZ, // 地表坡度矢量
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ, // 天线坐标系的X轴
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,// 天线坐标系的Y轴
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,// 天线坐标系的Z轴
|
||||
float antDirectX, float antDirectY, float antDirectZ,// 天线的指向
|
||||
float Pt,// 发射能量
|
||||
float* TransAntpattern, float Transtarttheta, float Transstartphi, float Transdtheta, float Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图
|
||||
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
|
||||
float NearR, float FarR, // 距离范围
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
|
||||
float* outR, // 输出距离
|
||||
float* outAmp // 输出增益
|
||||
) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
float RstX = antX - targetX[idx]; // 计算坐标矢量
|
||||
|
|
@ -482,18 +483,18 @@ __global__ void CUDAKernel_RFPC_Caluation_R_Gain(
|
|||
float slopeX = demSlopeX[idx];
|
||||
float slopeY = demSlopeY[idx];
|
||||
float slopeZ = demSlopeZ[idx];
|
||||
float RstR= sqrtf(RstX* RstX + RstY* RstY + RstZ* RstZ); // 矢量距离
|
||||
float RstR = sqrtf(RstX * RstX + RstY * RstY + RstZ * RstZ); // 矢量距离
|
||||
|
||||
if (RstR<NearR || RstR>FarR) {
|
||||
outLocalAngle[idx] = 0;
|
||||
|
||||
outR[idx] = 0;
|
||||
AmpGain[idx] = 0;
|
||||
outAmp[idx] = 0;
|
||||
}
|
||||
else {
|
||||
// 求解坡度
|
||||
float slopR = sqrtf(slopeX * slopeX + slopeY * slopeY + slopeZ * slopeZ); //
|
||||
float dotAB = RstX * slopeX + RstY * slopeY + RstZ * slopeZ;
|
||||
outLocalAngle[idx] = acosf(dotAB / (RstR * slopR)); // 局地入射角
|
||||
float localangle = acosf(dotAB / (RstR * slopR)); // 局地入射角
|
||||
float ampGain = 0;
|
||||
// 求解天线方向图指向
|
||||
CUDAVectorEllipsoidal antVector = GPU_SatelliteAntDirectNormal(
|
||||
|
|
@ -519,14 +520,39 @@ __global__ void CUDAKernel_RFPC_Caluation_R_Gain(
|
|||
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
|
||||
temptheta, tempphi);
|
||||
|
||||
// 计算
|
||||
float sigma0 = 0;
|
||||
{
|
||||
long clsid = demCls[idx];
|
||||
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 {
|
||||
float 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)); // 反射强度
|
||||
AmpGain[idx] = ampGain;
|
||||
outAmp[idx] = ampGain * Pt * sigma0;
|
||||
outR[idx] = RstR;
|
||||
}
|
||||
else {
|
||||
outR[idx] = 0;
|
||||
AmpGain[idx] = 0;
|
||||
outAmp[idx] = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -584,7 +610,7 @@ __global__ void CUDACkernel_Complex_SUM_reduce_dynamicshared(cuComplex* d_x, cu
|
|||
const int bid = blockIdx.x; // 某个block在网格grid内的标号 index
|
||||
const int n = bid * blockDim.x + tid; // n 是某个线程的标号 index
|
||||
__shared__ cuComplex s_y[128]; // 分配共享内存空间，不同的block都有共享内存变量的副本
|
||||
s_y[tid] = (n < N) ? d_x[n] : make_cuComplex(0.0,0.0); // 每个block的共享内存变量副本，都用全局内存数组d_x来赋值，最后一个多出来的用0
|
||||
s_y[tid] = (n < N) ? d_x[n] : make_cuComplex(0.0, 0.0); // 每个block的共享内存变量副本，都用全局内存数组d_x来赋值，最后一个多出来的用0
|
||||
__syncthreads(); // 线程块内部直接同步
|
||||
|
||||
for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1) // 折半
|
||||
|
|
@ -663,9 +689,9 @@ extern "C" void AntPatternInterpGain(float* anttheta, float* antphi, float* ga
|
|||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
//printf("\nCUDA_RFPC_SiglePRF blockSize:%d ,numBlock:%d\n", blockSize, numBlocks);
|
||||
|
||||
CUDA_AntPatternInterpGain << <numBlocks, blockSize >> > ( anttheta,antphi, gain,
|
||||
antpattern,
|
||||
starttheta, startphi, dtheta, dphi, thetapoints, phipoints,
|
||||
CUDA_AntPatternInterpGain << <numBlocks, blockSize >> > (anttheta, antphi, gain,
|
||||
antpattern,
|
||||
starttheta, startphi, dtheta, dphi, thetapoints, phipoints,
|
||||
len);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
|
|
@ -676,7 +702,7 @@ extern "C" void AntPatternInterpGain(float* anttheta, float* antphi, float* ga
|
|||
#endif // __CUDADEBUG__
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
|
||||
|
||||
extern "C" void calculationEcho(float* sigma0, float* TransAnt, float* ReciveAnt,
|
||||
float* localangle, float* R, float* slopeangle,
|
||||
|
|
@ -709,7 +735,7 @@ extern "C" void CUDACalculationEchoAmp(float* sigma0, float* TransAnt, float* Re
|
|||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_CalculationEchoAmp << <numBlocks, blockSize >> > (
|
||||
sigma0, TransAnt, ReciveAnt, R, Pt, ampArr, len);
|
||||
sigma0, TransAnt, ReciveAnt, R, Pt, ampArr, len);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
|
|
@ -749,7 +775,7 @@ extern "C" void CUDACombinationEchoAmpAndPhase(float* R,
|
|||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_CombinationEchoAmpAndPhase << <numBlocks, blockSize >> > (
|
||||
R,
|
||||
R,
|
||||
echoAmp, echoPhase,
|
||||
nearRange, Fs, plusepoints, echo, FreqID, len
|
||||
);
|
||||
|
|
@ -766,7 +792,7 @@ extern "C" void CUDACombinationEchoAmpAndPhase(float* R,
|
|||
|
||||
|
||||
extern "C" void CUDAInterpSigma(
|
||||
long* demcls,float* sigmaAmp, float* localanglearr,long len,
|
||||
long* demcls, float* sigmaAmp, float* localanglearr, long len,
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen) {// 地表覆盖类型-sigma插值对应函数-ulaby
|
||||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
|
|
@ -785,40 +811,45 @@ extern "C" void CUDAInterpSigma(
|
|||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
extern "C" void CUDARFPC_Caluation_R_Gain(float antX, float antY, float antZ,
|
||||
float* targetX, float* targetY, float* targetZ, long TargetPixelNumber,
|
||||
float* demSlopeX, float* demSlopeY, float* demSlopeZ,
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ,
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,
|
||||
float antDirectX, float antDirectY, float antDirectZ,
|
||||
float* TransAntpattern, float Transtarttheta, float Transstartphi, float Transdtheta, float Transdphi, int Transthetapoints, int Transphipoints,
|
||||
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,
|
||||
float NearR, float FarR,
|
||||
float* outR,
|
||||
float* outLocalAngle,
|
||||
float* AmpGain)
|
||||
extern "C" void CUDARFPC_Caluation_R_Gain(
|
||||
float antX, float antY, float antZ, // 天线的坐标
|
||||
float* targetX, float* targetY, float* targetZ, long TargetPixelNumber, // 地面坐标
|
||||
long* demCls,
|
||||
float* demSlopeX, float* demSlopeY, float* demSlopeZ, // 地表坡度矢量
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ, // 天线坐标系的X轴
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,// 天线坐标系的Y轴
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,// 天线坐标系的Z轴
|
||||
float antDirectX, float antDirectY, float antDirectZ,// 天线的指向
|
||||
float Pt,// 发射能量
|
||||
float* TransAntpattern, float Transtarttheta, float Transstartphi, float Transdtheta, float Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图
|
||||
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
|
||||
float NearR, float FarR, // 距离范围
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
|
||||
float* outR, // 输出距离
|
||||
float* outAmp // 输出增益
|
||||
)
|
||||
{
|
||||
|
||||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (TargetPixelNumber + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDAKernel_RFPC_Caluation_R_Gain << <numBlocks, blockSize >> > (
|
||||
antX, antY, antZ,
|
||||
targetX,targetY, targetZ, TargetPixelNumber,
|
||||
antX, antY, antZ,
|
||||
targetX, targetY, targetZ, TargetPixelNumber,
|
||||
demCls,
|
||||
demSlopeX, demSlopeY, demSlopeZ,
|
||||
antXaxisX, antXaxisY, antXaxisZ,
|
||||
antYaxisX, antYaxisY, antYaxisZ,
|
||||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ,
|
||||
TransAntpattern,
|
||||
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
|
||||
ReceiveAntpattern,
|
||||
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
|
||||
antXaxisX, antXaxisY, antXaxisZ,
|
||||
antYaxisX, antYaxisY, antYaxisZ,
|
||||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ,
|
||||
Pt,
|
||||
TransAntpattern,
|
||||
Transtarttheta, Transstartphi, Transdtheta, Transdphi, Transthetapoints, Transphipoints,
|
||||
ReceiveAntpattern,
|
||||
Receivestarttheta, Receivestartphi, Receivedtheta, Receivedphi, Receivethetapoints, Receivephipoints,
|
||||
NearR, FarR,
|
||||
outR,
|
||||
outLocalAngle,
|
||||
AmpGain
|
||||
outAmp
|
||||
);
|
||||
|
||||
|
||||
|
|
@ -836,12 +867,12 @@ extern "C" void CUDARFPC_Caluation_R_Gain(float antX, float antY, float antZ,
|
|||
}
|
||||
|
||||
extern "C" void CUDARFPC_Target_Freq_EchoData(float* InR,
|
||||
float* InlocalAngle,
|
||||
float* InampGain,
|
||||
long* Indemcls,
|
||||
float* InlocalAngle,
|
||||
float* InampGain,
|
||||
long* Indemcls,
|
||||
long len,
|
||||
float Pt, float freq,
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
|
||||
float Pt, float freq,
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,
|
||||
cuComplex* OutechoArr
|
||||
|
||||
)
|
||||
|
|
@ -871,13 +902,13 @@ extern "C" void CUDARFPC_Target_Freq_EchoData(float* InR,
|
|||
}
|
||||
|
||||
|
||||
extern "C" void CUDA_DemEchoSUM_NoMalloc(cuComplex* d_dem_echo,long N,
|
||||
extern "C" void CUDA_DemEchoSUM_NoMalloc(cuComplex* d_dem_echo, long N,
|
||||
cuComplex* d_echosum_temp, int grid_size) {
|
||||
|
||||
|
||||
long NUM_REPEATS = 100;
|
||||
const int smem = sizeof(float) * BLOCK_SIZE;
|
||||
|
||||
CUDACkernel_Complex_SUM_reduce_dynamicshared << <grid_size, BLOCK_SIZE, smem >> > (d_dem_echo, d_echosum_temp,N); //归约求和
|
||||
|
||||
CUDACkernel_Complex_SUM_reduce_dynamicshared << <grid_size, BLOCK_SIZE, smem >> > (d_dem_echo, d_echosum_temp, N); //归约求和
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
|
|
@ -886,7 +917,7 @@ extern "C" void CUDA_DemEchoSUM_NoMalloc(cuComplex* d_dem_echo,long N,
|
|||
}
|
||||
#endif // __CUDADEBUG__
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
|
||||
}
|
||||
|
||||
extern "C" void CUDA_SumPRF_Temp(cuComplex* d_dem_echo, long plusepoints, long grid_size, cuComplex* d_echo_PRF)
|
||||
|
|
|
|||
|
|
@ -67,18 +67,19 @@ extern "C" void CUDAInterpSigma(
|
|||
extern "C" void CUDARFPC_Caluation_R_Gain(
|
||||
float antX,float antY,float antZ, // 天线的坐标
|
||||
float* targetX,float* targetY, float* targetZ, long TargetPixelNumber, // 地面坐标
|
||||
long* demCls,
|
||||
float* demSlopeX, float* demSlopeY, float* demSlopeZ, // 地表坡度矢量
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ, // 天线坐标系的X轴
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,// 天线坐标系的Y轴
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,// 天线坐标系的Z轴
|
||||
float antDirectX, float antDirectY, float antDirectZ,// 天线的指向
|
||||
float Pt,// 发射能量
|
||||
float* TransAntpattern, float Transtarttheta, float Transstartphi, float Transdtheta, float Transdphi, int Transthetapoints, int Transphipoints, // 发射天线方向图
|
||||
float* ReceiveAntpattern, float Receivestarttheta, float Receivestartphi, float Receivedtheta, float Receivedphi, int Receivethetapoints, int Receivephipoints,//接收天线方向图
|
||||
float NearR,float FarR, // 距离范围
|
||||
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 插值图
|
||||
float* outR, // 输出距离
|
||||
float* outLocalAngle, // 输出局地坐标系
|
||||
float* AmpGain // 输出增益
|
||||
float* outAmp // 输出增益
|
||||
);
|
||||
|
||||
|
||||
|
|
|
|||
|
|
@ -568,7 +568,6 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
h_demsloper_x=(float* )mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
|
||||
h_demsloper_y=(float* )mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
|
||||
h_demsloper_z=(float* )mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
|
||||
h_demsloper_angle= (float*)mallocCUDAHost( sizeof(float) * blokline * tempDemCols);
|
||||
|
||||
d_dem_x=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols); // 7
|
||||
d_dem_y=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
|
|
@ -576,7 +575,6 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
d_demsloper_x=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
d_demsloper_y=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
d_demsloper_z=(float* )mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
d_demsloper_angle= (float*)mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
|
||||
// 提前声明参数变量
|
||||
float* h_R;// 辐射方向
|
||||
|
|
@ -591,8 +589,8 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
d_localangle= (float*)mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
|
||||
|
||||
float* h_gain = (float*)mallocCUDAHost( sizeof(float)* blokline* tempDemCols);
|
||||
float* d_gain = (float*)mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
float* h_amp = (float*)mallocCUDAHost( sizeof(float)* blokline* tempDemCols);
|
||||
float* d_amp = (float*)mallocCUDADevice( sizeof(float) * blokline * tempDemCols);
|
||||
|
||||
// 回波
|
||||
cuComplex* h_echo;
|
||||
|
|
@ -628,10 +626,8 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
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_demsloper_angle); FreeCUDADevice(d_demsloper_angle);//7
|
||||
FreeCUDAHost(h_R); FreeCUDADevice(d_R);
|
||||
FreeCUDAHost(h_localangle); FreeCUDADevice(d_localangle); //11
|
||||
FreeCUDAHost(h_gain); FreeCUDADevice(d_gain);
|
||||
FreeCUDAHost(h_amp); FreeCUDADevice(d_amp);
|
||||
FreeCUDAHost(h_echo); FreeCUDADevice(d_echo);//19
|
||||
FreeCUDAHost(h_demcls); FreeCUDADevice(d_demcls);
|
||||
|
||||
|
|
@ -643,11 +639,9 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
h_demsloper_z = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
|
||||
h_demsloper_angle = (float*)mallocCUDAHost(sizeof(float) * blokline * tempDemCols);
|
||||
h_R = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
|
||||
h_localangle = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
|
||||
h_gain = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
|
||||
h_amp = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
|
||||
h_echo = (cuComplex*)mallocCUDAHost(sizeof(cuComplex) * newblokline * tempDemCols);
|
||||
h_demcls = (long*)mallocCUDAHost(sizeof(long) * newblokline * tempDemCols);
|
||||
h_R = (float*)mallocCUDAHost(sizeof(float) * newblokline * tempDemCols);
|
||||
|
||||
d_dem_x=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
d_dem_y=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
|
|
@ -656,9 +650,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
d_demsloper_y=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
d_demsloper_z=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);//6
|
||||
d_demsloper_angle=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);//7
|
||||
d_R=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
d_localangle=(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
d_gain =(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
d_amp =(float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
d_echo=(cuComplex*)mallocCUDADevice(sizeof(cuComplex) * newblokline * tempDemCols);
|
||||
d_demcls = (long*)mallocCUDADevice(sizeof(long) * newblokline * tempDemCols);
|
||||
d_R = (float*)mallocCUDADevice(sizeof(float) * newblokline * tempDemCols);
|
||||
|
|
@ -683,7 +675,6 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
HostToDevice((void*)h_demsloper_x, (void*)d_demsloper_x, sizeof(float) * newblokline * tempDemCols);
|
||||
HostToDevice((void*)h_demsloper_y, (void*)d_demsloper_y, sizeof(float) * newblokline * tempDemCols);
|
||||
HostToDevice((void*)h_demsloper_z, (void*)d_demsloper_z, sizeof(float) * newblokline * tempDemCols);
|
||||
HostToDevice((void*)h_demsloper_angle, (void*)d_demsloper_angle, sizeof(float) * newblokline * tempDemCols);
|
||||
HostToDevice((void*)h_demcls, (void*)d_demcls, sizeof(long) * newblokline* tempDemCols);
|
||||
|
||||
|
||||
|
|
@ -752,49 +743,49 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
CUDARFPC_Caluation_R_Gain(
|
||||
antpx, antpy, antpz, // 天线的坐标
|
||||
d_dem_x, d_dem_y, d_dem_z, pixelcount, // 地面坐标
|
||||
d_demcls,
|
||||
d_demsloper_x, d_demsloper_y, d_demsloper_z, // 地表坡度矢量
|
||||
|
||||
antXaxisX, antXaxisY, antXaxisZ, // 天线坐标系的X轴
|
||||
antYaxisX, antYaxisY, antYaxisZ,// 天线坐标系的Y轴
|
||||
antZaxisX, antZaxisY, antZaxisZ,// 天线坐标系的Z轴
|
||||
antdirectx, antdirecty, antdirectz,// 天线的指向
|
||||
|
||||
Pt,
|
||||
d_TantPattern, TstartTheta, TstartPhi, Tdtheta, Tdphi, Tthetanum, Tphinum, // 发射天线方向图
|
||||
d_RantPattern, RstartTheta, RstartPhi, Rdtheta, Rdphi, Rthetanum, Rphinum,//接收天线方向图
|
||||
NearRange, FarRange,
|
||||
d_clsSigmaParam, clamapid,
|
||||
d_R, // 输出距离
|
||||
d_localangle, // 输出局地入射角
|
||||
d_gain // 输出增益
|
||||
d_amp // Êä³öÔöÒæ
|
||||
);
|
||||
// 计算某个具体回波
|
||||
|
||||
//DeviceToHost(h_echosum_temp, d_echosum_temp, sizeof(cuComplex)* templine* PlusePoint);
|
||||
for (long freqid = 0; freqid < freqlist.size(); freqid++) {
|
||||
float freqpoint = freqlist[freqid];
|
||||
CUDARFPC_Target_Freq_EchoData(d_R,
|
||||
d_localangle,
|
||||
d_gain,
|
||||
d_demcls,
|
||||
pixelcount,
|
||||
Pt, freqpoint,
|
||||
d_clsSigmaParam, clamapid,
|
||||
d_echo);
|
||||
// 数据求和
|
||||
CUDA_MemsetBlock(d_echosum_temp, initSumTemp, grid_size);
|
||||
long tempechoid = tempprfid * PlusePoint + freqid;
|
||||
CUDA_DemEchoSUM_NoMalloc(d_echo, pixelcount,
|
||||
d_echosum_temp + grid_size * freqid, grid_size
|
||||
);
|
||||
}
|
||||
////DeviceToHost(h_echosum_temp, d_echosum_temp, sizeof(cuComplex)* templine* PlusePoint);
|
||||
//for (long freqid = 0; freqid < freqlist.size(); freqid++) {
|
||||
// float freqpoint = freqlist[freqid];
|
||||
// CUDARFPC_Target_Freq_EchoData(d_R,
|
||||
// d_localangle,
|
||||
// d_demcls,
|
||||
// pixelcount,
|
||||
// Pt, freqpoint,
|
||||
// d_clsSigmaParam, clamapid,
|
||||
// d_echo);
|
||||
// // Êý¾ÝÇóºÍ
|
||||
// CUDA_MemsetBlock(d_echosum_temp, initSumTemp, grid_size);
|
||||
// long tempechoid = tempprfid * PlusePoint + freqid;
|
||||
// CUDA_DemEchoSUM_NoMalloc(d_echo, pixelcount,
|
||||
// d_echosum_temp + grid_size * freqid, grid_size
|
||||
// );
|
||||
//}
|
||||
|
||||
CUDA_SumPRF_Temp(d_echosum_temp, PlusePoint, grid_size, d_echo_PRF);
|
||||
DeviceToHost(h_echo_PRF, d_echo_PRF, sizeof(cuComplex) * PlusePoint);
|
||||
for (long freqid = 0; freqid < PlusePoint; freqid++) {
|
||||
echotemp.get()[tempprfid * PlusePoint + freqid] = echotemp.get()[tempprfid * PlusePoint + freqid]
|
||||
+ std::complex<float>(
|
||||
h_echo_PRF[freqid].x,
|
||||
h_echo_PRF[freqid].y);
|
||||
}
|
||||
//CUDA_SumPRF_Temp(d_echosum_temp, PlusePoint, grid_size, d_echo_PRF);
|
||||
//DeviceToHost(h_echo_PRF, d_echo_PRF, sizeof(cuComplex) * PlusePoint);
|
||||
//for (long freqid = 0; freqid < PlusePoint; freqid++) {
|
||||
// echotemp.get()[tempprfid * PlusePoint + freqid] =
|
||||
// echotemp.get()[tempprfid * PlusePoint + freqid]
|
||||
// + std::complex<float>(
|
||||
// h_echo_PRF[freqid].x,
|
||||
// h_echo_PRF[freqid].y);
|
||||
//}
|
||||
if (prfid % 100 == 0) {
|
||||
std::cout << "[" << QDateTime::currentDateTime().toString("yyyy-MM-dd hh:mm:ss.zzz").toStdString() << "] dem:\t" << startline << "\t-\t" << startline + newblokline << "\t:\t pluse :\t" << prfid << " / " << pluseCount << std::endl;
|
||||
}
|
||||
|
|
@ -825,8 +816,7 @@ ErrorCode RFPCProcessCls::RFPCMainProcess_GPU( )
|
|||
|
||||
// 临时变量释放
|
||||
FreeCUDAHost(h_R); FreeCUDADevice(d_R);
|
||||
FreeCUDAHost(h_localangle); FreeCUDADevice(d_localangle); //11
|
||||
FreeCUDAHost(h_gain); FreeCUDADevice(d_gain);
|
||||
FreeCUDAHost(h_amp); FreeCUDADevice(d_amp);
|
||||
FreeCUDAHost(h_demcls); FreeCUDADevice(d_demcls);
|
||||
|
||||
FreeCUDAHost(freqpoints);
|
||||
|
|
|
|||
Loading…
Reference in New Issue