修复RTPC算法的切片错误
parent
7bf440dc70
commit
798e02352f
|
|
@ -26,14 +26,15 @@
|
|||
#define MATPLOTDRAWIMAGE
|
||||
|
||||
|
||||
#define PI_180 180/3.141592653589793238462643383279
|
||||
#define T180_PI 3.141592653589793238462643383279/180
|
||||
#define r2d 180/3.141592653589793238462643383279
|
||||
#define d2r 3.141592653589793238462643383279/180
|
||||
#define LIGHTSPEED 299792458
|
||||
#define PRECISIONTOLERANCE 1e-9
|
||||
#define Radians2Degrees(Radians) Radians*PI_180
|
||||
#define Degrees2Radians(Degrees) Degrees*T180_PI
|
||||
#define EARTHWE 0.000072292115
|
||||
|
||||
#define PI 3.141592653589793238462643383279
|
||||
|
||||
|
||||
|
||||
#define earthRoute 0.000072292115
|
||||
|
|
@ -43,11 +44,10 @@
|
|||
|
||||
const std::complex<double> imagI(0, 1);
|
||||
|
||||
const double PI = 3.141592653589793238462643383279;
|
||||
|
||||
const double epsilon = 0.000000000000001;
|
||||
const double pi = 3.14159265358979323846;
|
||||
const double d2r = pi / 180;
|
||||
const double r2d = 180 / pi;
|
||||
|
||||
|
||||
const double a = 6378137.0; //椭球长半轴
|
||||
const double ae = 6378137.0; //椭球长半轴
|
||||
|
|
|
|||
572
GPUTool.cu
572
GPUTool.cu
|
|
@ -8,7 +8,7 @@
|
|||
#include <cuda_runtime.h>
|
||||
#include <cublas_v2.h>
|
||||
#include <cuComplex.h>
|
||||
|
||||
|
||||
#include "BaseConstVariable.h"
|
||||
#include "GPUTool.cuh"
|
||||
|
||||
|
|
@ -20,7 +20,7 @@
|
|||
|
||||
|
||||
// 定义参数
|
||||
__device__ cuComplex cuCexpf(cuComplex x)
|
||||
__device__ cuComplex cuCexpf(cuComplex x)
|
||||
{
|
||||
float factor = exp(x.x);
|
||||
return make_cuComplex(factor * cos(x.y), factor * sin(x.y));
|
||||
|
|
@ -28,136 +28,135 @@
|
|||
|
||||
// 定义仿真所需参数
|
||||
|
||||
__device__ float GPU_getSigma0dB(CUDASigmaParam param,float theta) {
|
||||
return param.p1 + param.p2 * exp(-param.p3 * theta) + param.p4 * cos(param.p5 * theta + param.p6);
|
||||
__device__ float GPU_getSigma0dB(CUDASigmaParam param, float theta) {
|
||||
return param.p1 + param.p2 * exp(-param.p3 * theta) + param.p4 * cos(param.p5 * theta + param.p6);
|
||||
}
|
||||
|
||||
__device__ CUDAVector GPU_VectorAB(CUDAVector A, CUDAVector B) {
|
||||
CUDAVector C;
|
||||
C.x = B.x - A.x;
|
||||
C.y = B.y - A.y;
|
||||
C.z = B.z - A.z;
|
||||
return C;
|
||||
__device__ CUDAVector GPU_VectorAB(CUDAVector A, CUDAVector B) {
|
||||
CUDAVector C;
|
||||
C.x = B.x - A.x;
|
||||
C.y = B.y - A.y;
|
||||
C.z = B.z - A.z;
|
||||
return C;
|
||||
}
|
||||
|
||||
__device__ float GPU_VectorNorm2(CUDAVector A) {
|
||||
return sqrtf(A.x * A.x + A.y * A.y + A.z * A.z);
|
||||
}
|
||||
|
||||
__device__ float GPU_dotVector(CUDAVector A, CUDAVector B) {
|
||||
return A.x * B.x + A.y * B.y + A.z * B.z;
|
||||
__device__ float GPU_VectorNorm2(CUDAVector A) {
|
||||
return sqrtf(A.x * A.x + A.y * A.y + A.z * A.z);
|
||||
}
|
||||
|
||||
__device__ float GPU_CosAngle_VectorA_VectorB(CUDAVector A, CUDAVector B) {
|
||||
return GPU_dotVector(A, B) / (GPU_VectorNorm2(A)*GPU_VectorNorm2(B));
|
||||
__device__ float GPU_dotVector(CUDAVector A, CUDAVector B) {
|
||||
return A.x * B.x + A.y * B.y + A.z * B.z;
|
||||
}
|
||||
|
||||
__device__ CUDAVectorEllipsoidal GPU_SatelliteAntDirectNormal(float RstX, float RstY, float RstZ,
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ,
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,
|
||||
float antDirectX, float antDirectY, float antDirectZ
|
||||
) {
|
||||
CUDAVectorEllipsoidal result{0,0,-1};
|
||||
float Xst = -1 * RstX; // 卫星 --> 地面
|
||||
float Yst = -1 * RstY;
|
||||
float Zst = -1 * RstZ;
|
||||
float AntXaxisX = antXaxisX;
|
||||
float AntXaxisY = antXaxisY;
|
||||
float AntXaxisZ = antXaxisZ;
|
||||
float AntYaxisX = antYaxisX;
|
||||
float AntYaxisY = antYaxisY;
|
||||
float AntYaxisZ = antYaxisZ;
|
||||
float AntZaxisX = antZaxisX;
|
||||
float AntZaxisY = antZaxisY;
|
||||
float AntZaxisZ = antZaxisZ;
|
||||
// 天线指向在天线坐标系下的值
|
||||
float Xant = (Xst * (AntYaxisY * AntZaxisZ - AntYaxisZ * AntZaxisY) + Xst * (AntXaxisZ * AntZaxisY - AntXaxisY * AntZaxisZ) + Xst * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Yant = (Yst * (AntYaxisZ * AntZaxisX - AntYaxisX * AntZaxisZ) + Yst * (AntXaxisX * AntZaxisZ - AntXaxisZ * AntZaxisX) + Yst * (AntYaxisX * AntXaxisZ - AntXaxisX * AntYaxisZ)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Zant = (Zst * (AntYaxisX * AntZaxisY - AntYaxisY * AntZaxisX) + Zst * (AntXaxisY * AntZaxisX - AntXaxisX * AntZaxisY) + Zst * (AntXaxisX * AntYaxisY - AntYaxisX * AntXaxisY)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
// 计算theta 与 phi
|
||||
float Norm = sqrtf(Xant * Xant + Yant * Yant + Zant * Zant); // 计算 pho
|
||||
float ThetaAnt = acosf(Zant / Norm); // theta 与 Z轴的夹角
|
||||
float YsinTheta = Yant / sinf(ThetaAnt);
|
||||
float PhiAnt = (YsinTheta / abs(YsinTheta)) * acosf(Xant / (Norm * sinf(ThetaAnt)));
|
||||
result.theta = ThetaAnt;
|
||||
result.phi = PhiAnt;
|
||||
result.pho = Norm;
|
||||
return result;
|
||||
__device__ float GPU_CosAngle_VectorA_VectorB(CUDAVector A, CUDAVector B) {
|
||||
return GPU_dotVector(A, B) / (GPU_VectorNorm2(A) * GPU_VectorNorm2(B));
|
||||
}
|
||||
|
||||
/**
|
||||
__device__ CUDAVectorEllipsoidal GPU_SatelliteAntDirectNormal(float RstX, float RstY, float RstZ,
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ,
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,
|
||||
float antZaxisX, float antZaxisY, float antZaxisZ,
|
||||
float antDirectX, float antDirectY, float antDirectZ
|
||||
) {
|
||||
CUDAVectorEllipsoidal result{ 0,0,-1 };
|
||||
float Xst = -1 * RstX; // 卫星 --> 地面
|
||||
float Yst = -1 * RstY;
|
||||
float Zst = -1 * RstZ;
|
||||
float AntXaxisX = antXaxisX;
|
||||
float AntXaxisY = antXaxisY;
|
||||
float AntXaxisZ = antXaxisZ;
|
||||
float AntYaxisX = antYaxisX;
|
||||
float AntYaxisY = antYaxisY;
|
||||
float AntYaxisZ = antYaxisZ;
|
||||
float AntZaxisX = antZaxisX;
|
||||
float AntZaxisY = antZaxisY;
|
||||
float AntZaxisZ = antZaxisZ;
|
||||
// 天线指向在天线坐标系下的值
|
||||
float Xant = (Xst * (AntYaxisY * AntZaxisZ - AntYaxisZ * AntZaxisY) + Xst * (AntXaxisZ * AntZaxisY - AntXaxisY * AntZaxisZ) + Xst * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Yant = (Yst * (AntYaxisZ * AntZaxisX - AntYaxisX * AntZaxisZ) + Yst * (AntXaxisX * AntZaxisZ - AntXaxisZ * AntZaxisX) + Yst * (AntYaxisX * AntXaxisZ - AntXaxisX * AntYaxisZ)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Zant = (Zst * (AntYaxisX * AntZaxisY - AntYaxisY * AntZaxisX) + Zst * (AntXaxisY * AntZaxisX - AntXaxisX * AntZaxisY) + Zst * (AntXaxisX * AntYaxisY - AntYaxisX * AntXaxisY)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
// 计算theta 与 phi
|
||||
float Norm = sqrtf(Xant * Xant + Yant * Yant + Zant * Zant); // 计算 pho
|
||||
float ThetaAnt = acosf(Zant / Norm); // theta 与 Z轴的夹角
|
||||
float YsinTheta = Yant / sinf(ThetaAnt);
|
||||
float PhiAnt = (YsinTheta / abs(YsinTheta)) * acosf(Xant / (Norm * sinf(ThetaAnt)));
|
||||
result.theta = ThetaAnt;
|
||||
result.phi = PhiAnt;
|
||||
result.pho = Norm;
|
||||
return result;
|
||||
}
|
||||
|
||||
/**
|
||||
天线方向图插值方法，以双线性插值算法为基础，由theta与phi组合得到的矩阵图为基础数据，通过插值计算的方法获取目标点的数据。
|
||||
其中行是theta、列是phi
|
||||
*/
|
||||
__device__ float GPU_BillerInterpAntPattern(float* antpattern,
|
||||
float starttheta, float startphi, float dtheta, float dphi,
|
||||
long thetapoints, long phipoints,
|
||||
float searththeta, float searchphi) {
|
||||
float stheta = searththeta;
|
||||
float sphi = searchphi;
|
||||
float pthetaid = (stheta - starttheta) / dtheta;//
|
||||
float pphiid = (sphi - startphi) / dphi;
|
||||
__device__ float GPU_BillerInterpAntPattern(float* antpattern,
|
||||
float starttheta, float startphi, float dtheta, float dphi,
|
||||
long thetapoints, long phipoints,
|
||||
float searththeta, float searchphi) {
|
||||
float stheta = searththeta;
|
||||
float sphi = searchphi;
|
||||
float pthetaid = (stheta - starttheta) / dtheta;//
|
||||
float pphiid = (sphi - startphi) / dphi;
|
||||
|
||||
long lasttheta = floorf(pthetaid);
|
||||
long nextTheta = lasttheta + 1;
|
||||
long lastphi = floorf(pphiid);
|
||||
long nextPhi = lastphi + 1;
|
||||
long lasttheta = floorf(pthetaid);
|
||||
long nextTheta = lasttheta + 1;
|
||||
long lastphi = floorf(pphiid);
|
||||
long nextPhi = lastphi + 1;
|
||||
|
||||
if (lasttheta < 0 || nextTheta < 0 || lastphi < 0 || nextPhi < 0 ||
|
||||
lasttheta >= thetapoints || nextTheta >= thetapoints || lastphi >= phipoints || nextPhi >= phipoints)
|
||||
{
|
||||
return 0;
|
||||
}
|
||||
else {
|
||||
float x = stheta;
|
||||
float y = sphi;
|
||||
if (lasttheta < 0 || nextTheta < 0 || lastphi < 0 || nextPhi < 0 ||
|
||||
lasttheta >= thetapoints || nextTheta >= thetapoints || lastphi >= phipoints || nextPhi >= phipoints)
|
||||
{
|
||||
return 0;
|
||||
}
|
||||
else {
|
||||
float x = stheta;
|
||||
float y = sphi;
|
||||
|
||||
float x1 = lasttheta * dtheta + starttheta;
|
||||
float x2 = nextTheta * dtheta + starttheta;
|
||||
float y1 = lastphi * dphi + startphi;
|
||||
float y2 = nextPhi * dphi + startphi;
|
||||
float x1 = lasttheta * dtheta + starttheta;
|
||||
float x2 = nextTheta * dtheta + starttheta;
|
||||
float y1 = lastphi * dphi + startphi;
|
||||
float y2 = nextPhi * dphi + startphi;
|
||||
|
||||
float z11 = antpattern[lasttheta * phipoints + lastphi];
|
||||
float z12 = antpattern[lasttheta * phipoints + nextPhi];
|
||||
float z21 = antpattern[nextTheta * phipoints + lastphi];
|
||||
float z22 = antpattern[nextTheta * phipoints + nextPhi];
|
||||
float z11 = antpattern[lasttheta * phipoints + lastphi];
|
||||
float z12 = antpattern[lasttheta * phipoints + nextPhi];
|
||||
float z21 = antpattern[nextTheta * phipoints + lastphi];
|
||||
float z22 = antpattern[nextTheta * phipoints + nextPhi];
|
||||
|
||||
|
||||
z11 = powf(10, z11 / 10);
|
||||
z12 = powf(10, z12 / 10);
|
||||
z21 = powf(10, z21 / 10);
|
||||
z22 = powf(10, z22 / 10);
|
||||
z11 = powf(10, z11 / 10);
|
||||
z12 = powf(10, z12 / 10);
|
||||
z21 = powf(10, z21 / 10);
|
||||
z22 = powf(10, z22 / 10);
|
||||
|
||||
float GainValue = (z11 * (x2 - x) * (y2 - y)
|
||||
+ z21 * (x - x1) * (y2 - y)
|
||||
+ z12 * (x2 - x) * (y - y1)
|
||||
+ z22 * (x - x1) * (y - y1));
|
||||
GainValue = GainValue / ((x2 - x1) * (y2 - y1));
|
||||
return GainValue;
|
||||
}
|
||||
}
|
||||
float GainValue = (z11 * (x2 - x) * (y2 - y)
|
||||
+ z21 * (x - x1) * (y2 - y)
|
||||
+ z12 * (x2 - x) * (y - y1)
|
||||
+ z22 * (x - x1) * (y - y1));
|
||||
GainValue = GainValue / ((x2 - x1) * (y2 - y1));
|
||||
return GainValue;
|
||||
}
|
||||
}
|
||||
|
||||
__device__ cuComplex GPU_calculationEcho(float sigma0, float TransAnt, float ReciveAnt,
|
||||
float localangle, float R, float slopeangle, float Pt, float lamda) {
|
||||
float r = R;
|
||||
float amp = Pt * TransAnt * ReciveAnt;
|
||||
amp = amp * sigma0;
|
||||
amp = amp / (powf(4 * LAMP_CUDA_PI, 2) * powf(r, 4)); // 反射强度
|
||||
float phi = (-4 * LAMP_CUDA_PI / lamda) * r;
|
||||
cuComplex echophi = make_cuComplex(0, phi);
|
||||
cuComplex echophiexp = cuCexpf(echophi);
|
||||
cuComplex echo;
|
||||
echo.x = echophiexp.x * amp;
|
||||
echo.y = echophiexp.y * amp;
|
||||
return echo;
|
||||
}
|
||||
|
||||
__device__ cuComplex GPU_calculationEcho(float sigma0, float TransAnt, float ReciveAnt,
|
||||
float localangle, float R, float slopeangle,float Pt, float lamda ) {
|
||||
float r = R;
|
||||
float amp = Pt * TransAnt * ReciveAnt;
|
||||
amp = amp * sigma0;
|
||||
amp = amp / (powf(4 * LAMP_CUDA_PI, 2) * powf(r, 4)); // 反射强度
|
||||
float phi = (-4 * LAMP_CUDA_PI / lamda) * r;
|
||||
cuComplex echophi = make_cuComplex(0, phi);
|
||||
cuComplex echophiexp = cuCexpf(echophi);
|
||||
cuComplex echo;
|
||||
echo.x = echophiexp.x * amp;
|
||||
echo.y = echophiexp.y * amp;
|
||||
return echo;
|
||||
}
|
||||
|
||||
__global__ void CUDA_DistanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz,float *R, long len) {
|
||||
__global__ void CUDA_DistanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* R, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
R[idx] = sqrtf(powf(Ax[idx]-Bx[idx], 2) + powf(Ay[idx] - By[idx], 2) + powf(Az[idx] - Bz[idx], 2));
|
||||
R[idx] = sqrtf(powf(Ax[idx] - Bx[idx], 2) + powf(Ay[idx] - By[idx], 2) + powf(Az[idx] - Bz[idx], 2));
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -176,15 +175,15 @@ __global__ void CUDA_make_VectorA_B(float sX, float sY, float sZ, float* tX, flo
|
|||
RstZ[idx] = sZ - tZ[idx];
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void CUDA_Norm_Vector(float* Vx, float* Vy, float* Vz,float *R, long len) {
|
||||
|
||||
__global__ void CUDA_Norm_Vector(float* Vx, float* Vy, float* Vz, float* R, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
R[idx] = sqrtf(powf(Vx[idx],2)+powf(Vy[idx],2)+powf(Vz[idx], 2));
|
||||
R[idx] = sqrtf(powf(Vx[idx], 2) + powf(Vy[idx], 2) + powf(Vz[idx], 2));
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void CUDA_cosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* anglecos,long len) {
|
||||
__global__ void CUDA_cosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* anglecos, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
float tAx = Ax[idx];
|
||||
|
|
@ -193,88 +192,53 @@ __global__ void CUDA_cosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx,
|
|||
float tBx = Bx[idx];
|
||||
float tBy = By[idx];
|
||||
float tBz = Bz[idx];
|
||||
float AR = sqrtf(powf(tAx,2) + powf(tAy,2) + powf(tAz,2));
|
||||
float BR = sqrtf(powf(tBx,2) + powf(tBy,2) + powf(tBz,2));
|
||||
float AR = sqrtf(powf(tAx, 2) + powf(tAy, 2) + powf(tAz, 2));
|
||||
float BR = sqrtf(powf(tBx, 2) + powf(tBy, 2) + powf(tBz, 2));
|
||||
float dotAB = tAx * tBx + tAy * tBy + tAz * tBz;
|
||||
float result =acosf( dotAB / (AR * BR));
|
||||
float result = acosf(dotAB / (AR * BR));
|
||||
anglecos[idx] = result;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void CUDA_SatelliteAntDirectNormal(float* RstX,float* RstY,float* RstZ,
|
||||
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* thetaAnt,float* phiAnt
|
||||
__global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* RstZ,
|
||||
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* thetaAnt, float* phiAnt
|
||||
, 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 AntXaxisX=antXaxisX ;
|
||||
float AntXaxisY=antXaxisY ;
|
||||
float AntXaxisZ=antXaxisZ ;
|
||||
float AntYaxisX=antYaxisX ;
|
||||
float AntYaxisY=antYaxisY ;
|
||||
float AntYaxisZ=antYaxisZ ;
|
||||
float AntZaxisX=antZaxisX ;
|
||||
float AntZaxisY=antZaxisY ;
|
||||
float AntZaxisZ=antZaxisZ ;
|
||||
float Xst = -1 * RstX[idx]; // 卫星 --> 地面
|
||||
float Yst = -1 * RstY[idx];
|
||||
float Zst = -1 * RstZ[idx];
|
||||
float AntXaxisX = antXaxisX;
|
||||
float AntXaxisY = antXaxisY;
|
||||
float AntXaxisZ = antXaxisZ;
|
||||
float AntYaxisX = antYaxisX;
|
||||
float AntYaxisY = antYaxisY;
|
||||
float AntYaxisZ = antYaxisZ;
|
||||
float AntZaxisX = antZaxisX;
|
||||
float AntZaxisY = antZaxisY;
|
||||
float AntZaxisZ = antZaxisZ;
|
||||
// 天线指向在天线坐标系下的值
|
||||
float Xant = (Xst * (AntYaxisY * AntZaxisZ - AntYaxisZ * AntZaxisY) + Xst * ( AntXaxisZ * AntZaxisY - AntXaxisY * AntZaxisZ) + Xst * ( AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY)) / ( AntXaxisX * ( AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * ( AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * ( AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Yant = (Yst * (AntYaxisZ * AntZaxisX - AntYaxisX * AntZaxisZ) + Yst * ( AntXaxisX * AntZaxisZ - AntXaxisZ * AntZaxisX) + Yst * ( AntYaxisX * AntXaxisZ - AntXaxisX * AntYaxisZ)) / ( AntXaxisX * ( AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * ( AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * ( AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Zant = (Zst * (AntYaxisX * AntZaxisY - AntYaxisY * AntZaxisX) + Zst * ( AntXaxisY * AntZaxisX - AntXaxisX * AntZaxisY) + Zst * ( AntXaxisX * AntYaxisY - AntYaxisX * AntXaxisY)) / ( AntXaxisX * ( AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * ( AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * ( AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Xant = (Xst * (AntYaxisY * AntZaxisZ - AntYaxisZ * AntZaxisY) + Xst * (AntXaxisZ * AntZaxisY - AntXaxisY * AntZaxisZ) + Xst * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Yant = (Yst * (AntYaxisZ * AntZaxisX - AntYaxisX * AntZaxisZ) + Yst * (AntXaxisX * AntZaxisZ - AntXaxisZ * AntZaxisX) + Yst * (AntYaxisX * AntXaxisZ - AntXaxisX * AntYaxisZ)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
float Zant = (Zst * (AntYaxisX * AntZaxisY - AntYaxisY * AntZaxisX) + Zst * (AntXaxisY * AntZaxisX - AntXaxisX * AntZaxisY) + Zst * (AntXaxisX * AntYaxisY - AntYaxisX * AntXaxisY)) / (AntXaxisX * (AntYaxisY * AntZaxisZ - AntZaxisY * AntYaxisZ) - AntYaxisX * (AntXaxisY * AntZaxisZ - AntXaxisZ * AntZaxisY) + AntZaxisX * (AntXaxisY * AntYaxisZ - AntXaxisZ * AntYaxisY));
|
||||
// 计算theta 与 phi
|
||||
float Norm = sqrtf(Xant * Xant + Yant * Yant + Zant * Zant); // 计算 pho
|
||||
float ThetaAnt = acosf(Zant / Norm); // theta 与 Z轴的夹角
|
||||
float Norm = sqrtf(Xant * Xant + Yant * Yant + Zant * Zant); // 计算 pho
|
||||
float ThetaAnt = acosf(Zant / Norm); // theta 与 Z轴的夹角
|
||||
float YsinTheta = Yant / sinf(ThetaAnt);
|
||||
float PhiAnt = (YsinTheta/abs(YsinTheta)) * acosf( Xant / (Norm * sinf(ThetaAnt)));
|
||||
float PhiAnt = (YsinTheta / abs(YsinTheta)) * acosf(Xant / (Norm * sinf(ThetaAnt)));
|
||||
thetaAnt[idx] = ThetaAnt;
|
||||
phiAnt[idx] = PhiAnt;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
__global__ void CUDA_calculationEcho(float* sigma0, float* TransAnt, float* ReciveAnt,
|
||||
float* localangle, float* R,float* slopeangle,
|
||||
float nearRange, float Fs,float Pt,float lamda,long FreqIDmax,
|
||||
cuComplex* echoArr , long* FreqID,
|
||||
long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
float r = R[idx];
|
||||
float amp = Pt * TransAnt[idx] * ReciveAnt[idx];
|
||||
amp= amp * sigma0[idx];
|
||||
amp = amp / (powf(4* LAMP_CUDA_PI,2)*powf(r,4)); // 反射强度
|
||||
|
||||
// 处理相位
|
||||
float phi = (-4 * LAMP_CUDA_PI / lamda) * r;
|
||||
cuComplex echophi = make_cuComplex(0, phi) ;
|
||||
cuComplex echophiexp = cuCexpf(echophi);
|
||||
|
||||
float timeR = 2 * (r - nearRange) / LIGHTSPEED * Fs;
|
||||
long timeID = floorf(timeR);
|
||||
if (timeID < 0 || timeID >= FreqIDmax) {
|
||||
timeID = 0;
|
||||
amp = 0;
|
||||
}
|
||||
|
||||
cuComplex echo;
|
||||
echo.x = echophiexp.x * amp;
|
||||
echo.y = echophiexp.y * amp;
|
||||
|
||||
|
||||
echoArr[idx] = echo;
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void CUDA_BillerInterpAntPattern(float* antpattern,
|
||||
float starttheta, float startphi, float dtheta, float dphi,
|
||||
long thetapoints, long phipoints,
|
||||
float* searththeta, float* searchphi,float* searchantpattern,
|
||||
float* searththeta, float* searchphi, float* searchantpattern,
|
||||
long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
|
|
@ -324,31 +288,36 @@ __global__ void CUDA_BillerInterpAntPattern(float* antpattern,
|
|||
}
|
||||
|
||||
|
||||
__global__ void CUDA_Test_HelloWorld(float a, long len) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
printf("\nidx:\t %d %d \n", idx, len);
|
||||
}
|
||||
|
||||
__global__ void CUDA_RTPC_Kernel(
|
||||
|
||||
__global__ void CUDA_RTPC(
|
||||
float antPx, float antPy, float antPz,// 天线坐标
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ, // 天线坐标系
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ, //
|
||||
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* demx, float* demy, float* demz, long* demcls, // 地面坐标
|
||||
float* demslopex, float* demslopey, float* demslopez, float* demslopeangle,// 地面坡度
|
||||
float* Tantpattern, float Tstarttheta, float Tstartphi, float Tdtheta, float Tdphi, long Tthetapoints, long Tphipoints,// 天线方向图相关
|
||||
float* Rantpattern, float Rstarttheta, float Rstartphi, float Rdtheta, float Rdphi, long Rthetapoints, long Rphipoints,// 天线方向图相关
|
||||
float antDirectX, float antDirectY, float antDirectZ,
|
||||
float* demx, float* demy, float* demz, long* demcls,
|
||||
float* demslopex, float* demslopey, float* demslopez, float* demslopeangle,
|
||||
float* Tantpattern, float Tstarttheta, float Tstartphi, float Tdtheta, float Tdphi, long Tthetapoints, long Tphipoints,
|
||||
float* Rantpattern, float Rstarttheta, float Rstartphi, float Rdtheta, float Rdphi, long Rthetapoints, long Rphipoints,
|
||||
float lamda, float fs, float nearrange, float Pt, long Freqnumbers, // 参数
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 地表覆盖类型-sigma插值对应函数-ulaby
|
||||
cuComplex* outecho, long* d_echoAmpFID,
|
||||
long len
|
||||
) {
|
||||
long idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx < len) {
|
||||
cuComplex* outecho, int* d_echoAmpFID,
|
||||
int linecount,int plusepoint) {
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
//printf("\nidx:\t %d %d %d\n", idx, linecount, plusepoint);
|
||||
if (idx < linecount* plusepoint) {
|
||||
long clsid = demcls[idx];
|
||||
CUDAVector Rs{ antPx,antPy,antPz };
|
||||
CUDAVector Rt{ demx[idx],demy[idx],demz[idx] };
|
||||
CUDAVector Rst{ Rs.x - Rt.x,Rs.y - Rt.y,Rs.z - Rt.z };
|
||||
CUDAVector Vslope{ demslopex[idx],demslopey[idx],demslopez[idx] };
|
||||
float R = GPU_VectorNorm2(Rst); // 斜距
|
||||
|
||||
float slopeangle = demslopeangle[idx];
|
||||
CUDAVectorEllipsoidal Rtanttheta = GPU_SatelliteAntDirectNormal( // 地面目标在天线的位置
|
||||
Rst.x, Rst.y, Rst.z,
|
||||
antXaxisX, antXaxisY, antXaxisZ,
|
||||
|
|
@ -356,21 +325,23 @@ __global__ void CUDA_RTPC_Kernel(
|
|||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ);
|
||||
|
||||
float localangle=GPU_CosAngle_VectorA_VectorB(Rst, Vslope); // 距地入射角
|
||||
float sigma = GPU_getSigma0dB(sigma0Paramslist[clsid], localangle);
|
||||
float localangle = GPU_CosAngle_VectorA_VectorB(Rst, Vslope); // 距地入射角
|
||||
float sigma = GPU_getSigma0dB(sigma0Paramslist[clsid], localangle * r2d);
|
||||
sigma = powf(10.0, sigma / 10.0);// 后向散射系数
|
||||
|
||||
//printf("\ntheta: %f\t,%f ,%f ,%f ,%f ,%f ,%f \n", localangle * r2d, sigma0Paramslist[clsid].p1, sigma0Paramslist[clsid].p2, sigma0Paramslist[clsid].p3,
|
||||
// sigma0Paramslist[clsid].p4, sigma0Paramslist[clsid].p5, sigma0Paramslist[clsid].p6);
|
||||
// 发射方向图
|
||||
float transPattern = GPU_BillerInterpAntPattern(Tantpattern,
|
||||
Tstarttheta, Tstartphi, Tdtheta, Tdphi,Tthetapoints, Tphipoints,
|
||||
Rtanttheta.theta, Rtanttheta.phi);
|
||||
|
||||
Tstarttheta, Tstartphi, Tdtheta, Tdphi, Tthetapoints, Tphipoints,
|
||||
Rtanttheta.theta, Rtanttheta.phi) * r2d;
|
||||
|
||||
// 接收方向图
|
||||
float receivePattern = GPU_BillerInterpAntPattern(Rantpattern,
|
||||
Rstarttheta, Rstartphi, Rdtheta, Rdphi, Rthetapoints, Rphipoints,
|
||||
Rtanttheta.theta, Rtanttheta.phi);
|
||||
Rtanttheta.theta, Rtanttheta.phi) * r2d;
|
||||
// 计算振幅、相位
|
||||
float amp = Pt * transPattern * receivePattern * sigma / (powf(4 * LAMP_CUDA_PI, 2) * powf(R, 4));
|
||||
float amp = Pt * transPattern * receivePattern * sigma * (1 / cos(slopeangle) * sin(localangle));
|
||||
amp = amp / (powf(4 * LAMP_CUDA_PI, 2) * powf(R, 4));
|
||||
float phi = (-4 * LAMP_CUDA_PI / lamda) * R;
|
||||
|
||||
// 构建回波
|
||||
|
|
@ -381,7 +352,7 @@ __global__ void CUDA_RTPC_Kernel(
|
|||
if (timeID < 0 || timeID >= Freqnumbers) {
|
||||
timeID = 0;
|
||||
amp = 0;
|
||||
|
||||
|
||||
}
|
||||
else {}
|
||||
|
||||
|
|
@ -391,6 +362,8 @@ __global__ void CUDA_RTPC_Kernel(
|
|||
outecho[idx] = echo;
|
||||
d_echoAmpFID[idx] = timeID;
|
||||
}
|
||||
|
||||
|
||||
}
|
||||
|
||||
|
||||
|
|
@ -402,48 +375,111 @@ void checkCudaError(cudaError_t err, const char* msg) {
|
|||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
// 主机参数内存声明
|
||||
extern "C" void* mallocCUDAHost( long memsize) {
|
||||
extern "C" void* mallocCUDAHost(long memsize) {
|
||||
void* ptr;
|
||||
cudaMallocHost(&ptr, memsize);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("mallocCUDAHost CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
return ptr;
|
||||
}
|
||||
|
||||
// 主机参数内存释放
|
||||
extern "C" void FreeCUDAHost(void* ptr) {
|
||||
cudaFreeHost(ptr);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("FreeCUDAHost CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
|
||||
}
|
||||
|
||||
// GPU参数内存声明
|
||||
extern "C" void* mallocCUDADevice( long memsize) {
|
||||
extern "C" void* mallocCUDADevice(long memsize) {
|
||||
void* ptr;
|
||||
cudaMalloc(&ptr, memsize);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("mallocCUDADevice CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
|
||||
return ptr;
|
||||
}
|
||||
|
||||
// GPU参数内存释放
|
||||
extern "C" void FreeCUDADevice(void* ptr) {
|
||||
cudaFree(ptr);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("FreeCUDADevice CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
|
||||
}
|
||||
|
||||
// GPU 内存数据转移
|
||||
extern "C" void HostToDevice(void* hostptr, void* deviceptr, long memsize) {
|
||||
cudaMemcpy(deviceptr, hostptr, memsize, cudaMemcpyHostToDevice);
|
||||
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("HostToDevice CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
|
||||
|
||||
}
|
||||
|
||||
extern "C" void DeviceToHost(void* hostptr, void* deviceptr, long memsize) {
|
||||
cudaMemcpy(hostptr, deviceptr, memsize, cudaMemcpyDeviceToHost);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("DeviceToHost CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
|
||||
}
|
||||
|
||||
|
||||
|
||||
extern "C" void distanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* R,long len) {
|
||||
extern "C" void CUDATestHelloWorld(float a,long len) {
|
||||
// 设置 CUDA 核函数的网格和块的尺寸
|
||||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_DistanceAB << <blockSize, numBlocks >> > ( Ax, Ay, Az, Bx, By, Bz, R, len);
|
||||
CUDA_Test_HelloWorld << <blockSize, numBlocks >> > (a, len);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("FreeCUDADevice CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
extern "C" void distanceAB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* R, long len) {
|
||||
// 设置 CUDA 核函数的网格和块的尺寸
|
||||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_DistanceAB << <blockSize, numBlocks >> > (Ax, Ay, Az, Bx, By, Bz, R, len);
|
||||
}
|
||||
|
||||
extern "C" void BdistanceAs(float* Ax, float* Ay, float* Az, float Bx, float By, float Bz, float* R, long len) {
|
||||
|
|
@ -460,7 +496,7 @@ extern "C" void make_VectorA_B(float sX, float sY, float sZ, float* tX, float* t
|
|||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_make_VectorA_B << <blockSize, numBlocks >> > (sX, sY, sZ,tX, tY, tZ, RstX,RstY, RstZ, len);
|
||||
CUDA_make_VectorA_B << <blockSize, numBlocks >> > (sX, sY, sZ, tX, tY, tZ, RstX, RstY, RstZ, len);
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
|
|
@ -469,7 +505,7 @@ extern "C" void Norm_Vector(float* Vx, float* Vy, float* Vz, float* R, long len)
|
|||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_Norm_Vector << <blockSize, numBlocks >> > (Vx,Vy,Vz,R, len);
|
||||
CUDA_Norm_Vector << <blockSize, numBlocks >> > (Vx, Vy, Vz, R, len);
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
|
|
@ -480,7 +516,7 @@ extern "C" void cosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, float
|
|||
CUDA_cosAngle_VA_AB << <blockSize, numBlocks >> > (Ax, Ay, Az, Bx, By, Bz, anglecos, len);
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
|
||||
extern "C" void SatelliteAntDirectNormal(float* RstX, float* RstY, float* RstZ,
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ,
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ,
|
||||
|
|
@ -492,79 +528,91 @@ extern "C" void SatelliteAntDirectNormal(float* RstX, float* RstY, float* RstZ,
|
|||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_SatelliteAntDirectNormal << <blockSize, numBlocks >> > ( RstX, RstY, RstZ,
|
||||
antXaxisX, antXaxisY, antXaxisZ,
|
||||
antYaxisX, antYaxisY, antYaxisZ,
|
||||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ,
|
||||
thetaAnt, phiAnt
|
||||
, len);
|
||||
CUDA_SatelliteAntDirectNormal << <blockSize, numBlocks >> > (RstX, RstY, RstZ,
|
||||
antXaxisX, antXaxisY, antXaxisZ,
|
||||
antYaxisX, antYaxisY, antYaxisZ,
|
||||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ,
|
||||
thetaAnt, phiAnt
|
||||
, len);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
extern "C" void calculationEcho(float* sigma0,float* TransAnt,float* ReciveAnt,
|
||||
float* localangle,float* R, float* slopeangle,
|
||||
float nearRange,float Fs, float pt, float lamda, long FreqIDmax,
|
||||
cuComplex* echoAmp,long* FreqID,
|
||||
long len)
|
||||
{
|
||||
extern "C" void CUDARTPCPRF(float antPx, long len) {
|
||||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_calculationEcho << <blockSize, numBlocks >> > ( sigma0, TransAnt,ReciveAnt,
|
||||
localangle, R, slopeangle,
|
||||
nearRange, Fs, pt, lamda, FreqIDmax,
|
||||
echoAmp, FreqID,
|
||||
len);
|
||||
printf("\nCUDA_RTPC_SiglePRF blockSize:%d ,numBlock:%d\n", blockSize, numBlocks);
|
||||
CUDA_Test_HelloWorld << <blockSize, numBlocks >> > (antPx, len);
|
||||
|
||||
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA_RTPC_SiglePRF CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
extern "C" void CUDA_RTPC_SiglePRF(
|
||||
float antPx, float antPy, float antPZ,// 天线坐标
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ, // 天线坐标系
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ, //
|
||||
float antPx, float antPy, float antPZ,
|
||||
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* demx, float* demy, float* demz, long* demcls, // 地面坐标
|
||||
float* demslopex, float* demslopey, float* demslopez, float* demslopeangle,// 地面坡度
|
||||
float* Tantpattern, float Tstarttheta, float Tstartphi, float Tdtheta, float Tdphi, long Tthetapoints, long Tphipoints,// 天线方向图相关
|
||||
float* Rantpattern, float Rstarttheta, float Rstartphi, float Rdtheta, float Rdphi, long Rthetapoints, long Rphipoints,// 天线方向图相关
|
||||
float lamda, float fs, float nearrange, float Pt, long Freqnumbers, // 参数
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 地表覆盖类型-sigma插值对应函数-ulaby
|
||||
cuComplex* outecho, long* d_echoAmpFID,
|
||||
long len
|
||||
) {
|
||||
float antDirectX, float antDirectY, float antDirectZ,
|
||||
float* demx, float* demy, float* demz, long* demcls,
|
||||
float* demslopex, float* demslopey, float* demslopez, float* demslopeangle,
|
||||
float* Tantpattern, float Tstarttheta, float Tstartphi, float Tdtheta, float Tdphi, int Tthetapoints, int Tphipoints,
|
||||
float* Rantpattern, float Rstarttheta, float Rstartphi, float Rdtheta, float Rdphi, int Rthetapoints, int Rphipoints,
|
||||
float lamda, float fs, float nearrange, float Pt, int Freqnumbers,
|
||||
CUDASigmaParam* sigma0Paramslist, int sigmaparamslistlen,
|
||||
cuComplex* outecho, int* d_echoAmpFID,
|
||||
int linecount,int colcount) {
|
||||
|
||||
int blockSize = 256; // 每个块的线程数
|
||||
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
// 调用 CUDA 核函数
|
||||
CUDA_RTPC_Kernel<<<blockSize,numBlocks>>>(
|
||||
antPx, antPy, antPZ,// 天线坐标
|
||||
antXaxisX, antXaxisY, antXaxisZ, // 天线坐标系
|
||||
antYaxisX, antYaxisY, antYaxisZ, //
|
||||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ,// 天线指向
|
||||
demx, demy, demz, demcls, // 地面坐标
|
||||
demslopex, demslopey, demslopez, demslopeangle,// 地面坡度
|
||||
Tantpattern, Tstarttheta, Tstartphi, Tdtheta, Tdphi, Tthetapoints, Tphipoints,// 天线方向图相关
|
||||
Rantpattern, Rstarttheta, Rstartphi, Rdtheta, Rdphi, Rthetapoints, Rphipoints,// 天线方向图相关
|
||||
lamda, fs, nearrange, Pt, Freqnumbers, // 参数
|
||||
sigma0Paramslist, sigmaparamslistlen,// 地表覆盖类型-sigma插值对应函数-ulaby
|
||||
outecho, d_echoAmpFID,
|
||||
len
|
||||
);
|
||||
cudaDeviceSynchronize();
|
||||
int numBlocks = (linecount* colcount + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
|
||||
//printf("\nCUDA_RTPC_SiglePRF blockSize:%d ,numBlock:%d\n",blockSize,numBlocks);
|
||||
// 调用 CUDA 核函数 CUDA_RTPC_Kernel
|
||||
|
||||
CUDA_RTPC << <numBlocks, blockSize >> > (
|
||||
antPx, antPy, antPZ,// 天线坐标
|
||||
antXaxisX, antXaxisY, antXaxisZ, // 天线坐标系
|
||||
antYaxisX, antYaxisY, antYaxisZ, //
|
||||
antZaxisX, antZaxisY, antZaxisZ,
|
||||
antDirectX, antDirectY, antDirectZ,// 天线指向
|
||||
demx, demy, demz,
|
||||
demcls, // 地面坐标
|
||||
demslopex, demslopey, demslopez, demslopeangle,// 地面坡度
|
||||
Tantpattern, Tstarttheta, Tstartphi, Tdtheta, Tdphi, Tthetapoints, Tphipoints,// 天线方向图相关
|
||||
Rantpattern, Rstarttheta, Rstartphi, Rdtheta, Rdphi, Rthetapoints, Rphipoints,// 天线方向图相关
|
||||
lamda, fs, nearrange, Pt, Freqnumbers, // 参数
|
||||
sigma0Paramslist, sigmaparamslistlen,// 地表覆盖类型-sigma插值对应函数-ulaby
|
||||
outecho, d_echoAmpFID,
|
||||
linecount, colcount
|
||||
);
|
||||
#ifdef __CUDADEBUG__
|
||||
cudaError_t err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA_RTPC_SiglePRF CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
// Possibly: exit(-1) if program cannot continue....
|
||||
}
|
||||
#endif // __CUDADEBUG__
|
||||
cudaDeviceSynchronize();
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
|
||||
|
|
|
|||
31
GPUTool.cuh
31
GPUTool.cuh
|
|
@ -7,6 +7,8 @@
|
|||
#include <cublas_v2.h>
|
||||
#include <cuComplex.h>
|
||||
|
||||
#define __CUDADEBUG__
|
||||
|
||||
// ĬÈÏÏÔ´æ·Ö²¼
|
||||
|
||||
|
||||
|
|
@ -56,25 +58,28 @@ extern "C" void make_VectorA_B(float sX, float sY, float sZ, float* tX, float* t
|
|||
extern "C" void Norm_Vector(float* Vx, float* Vy, float* Vz, float* R, long member);
|
||||
extern "C" void cosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, float* By, float* Bz, float* anglecos, long len);
|
||||
extern "C" void SatelliteAntDirectNormal(float* RstX, float* RstY, float* RstZ, 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* thetaAnt, float* phiAnt, long len);
|
||||
extern "C" void calculationEcho(float* sigma0, float* TransAnt, float* ReciveAnt,float* localangle, float* R, float* slopeangle,float nearRange, float Fs, float pt, float lamda, long FreqIDmax,cuComplex* echoAmp, long* FreqID, long len);
|
||||
extern "C" void CUDA_RTPC_SiglePRF(
|
||||
float antPx, float antPy, float antPZ,// 天线坐标
|
||||
float antXaxisX, float antXaxisY, float antXaxisZ, // 天线坐标系
|
||||
float antYaxisX, float antYaxisY, float antYaxisZ, //
|
||||
float antPx, float antPy, float antPZ,
|
||||
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* demx, float* demy, float* demz,long* demcls, // 地面坐标
|
||||
float* demslopex, float* demslopey, float* demslopez, float* demslopeangle,// 地面坡度
|
||||
float* Tantpattern, float Tstarttheta, float Tstartphi, float Tdtheta, float Tdphi, long Tthetapoints, long Tphipoints,// 天线方向图相关
|
||||
float* Rantpattern, float Rstarttheta, float Rstartphi, float Rdtheta, float Rdphi, long Rthetapoints, long Rphipoints,// 天线方向图相关
|
||||
float lamda, float fs, float nearrange, float Pt, long Freqnumbers, // 参数
|
||||
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen,// 地表覆盖类型-sigma插值对应函数-ulaby
|
||||
cuComplex* outecho,long* d_echoAmpFID,
|
||||
long len
|
||||
float antDirectX, float antDirectY, float antDirectZ,
|
||||
float* demx, float* demy, float* demz, long* demcls,
|
||||
float* demslopex, float* demslopey, float* demslopez, float* demslopeangle,
|
||||
float* Tantpattern, float Tstarttheta, float Tstartphi, float Tdtheta, float Tdphi, int Tthetapoints, int Tphipoints,
|
||||
float* Rantpattern, float Rstarttheta, float Rstartphi, float Rdtheta, float Rdphi, int Rthetapoints, int Rphipoints,
|
||||
float lamda, float fs, float nearrange, float Pt, int Freqnumbers,
|
||||
CUDASigmaParam* sigma0Paramslist, int sigmaparamslistlen,
|
||||
cuComplex* outecho, int* d_echoAmpFID,
|
||||
int linecount, int colcount
|
||||
);
|
||||
|
||||
extern "C" void CUDARTPCPRF(float antPx, long len);
|
||||
|
||||
|
||||
|
||||
extern "C" void CUDATestHelloWorld(float a, long len);
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue