diff --git a/BaseConstVariable.h b/BaseConstVariable.h index 941cf24..1aaa76b 100644 --- a/BaseConstVariable.h +++ b/BaseConstVariable.h @@ -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 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; //椭球长半轴 diff --git a/GPUTool.cu b/GPUTool.cu index 58fa1da..18fce48 100644 --- a/GPUTool.cu +++ b/GPUTool.cu @@ -8,7 +8,7 @@ #include #include #include - + #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; +} + +/** ߷ͼֵ˫Բֵ㷨ΪthetaphiϵõľͼΪݣֵͨķȡĿݡ thetaphi */ - __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 ˺Ϳijߴ int blockSize = 256; // ÿ߳ int numBlocks = (len + blockSize - 1) / blockSize; // pixelcount С // CUDA ˺ - CUDA_DistanceAB << > > ( Ax, Ay, Az, Bx, By, Bz, R, len); + CUDA_Test_HelloWorld << > > (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 ˺Ϳijߴ + int blockSize = 256; // ÿ߳ + int numBlocks = (len + blockSize - 1) / blockSize; // pixelcount С + // CUDA ˺ + CUDA_DistanceAB << > > (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 << > > (sX, sY, sZ,tX, tY, tZ, RstX,RstY, RstZ, len); + CUDA_make_VectorA_B << > > (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 << > > (Vx,Vy,Vz,R, len); + CUDA_Norm_Vector << > > (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 << > > (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 << > > ( RstX, RstY, RstZ, - antXaxisX, antXaxisY, antXaxisZ, - antYaxisX, antYaxisY, antYaxisZ, - antZaxisX, antZaxisY, antZaxisZ, - antDirectX, antDirectY, antDirectZ, - thetaAnt, phiAnt - , len); + CUDA_SatelliteAntDirectNormal << > > (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 << > > ( 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 << > > (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<<>>( - 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 << > > ( + 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 \ No newline at end of file +#endif + + diff --git a/GPUTool.cuh b/GPUTool.cuh index 64d38fa..dd19e48 100644 --- a/GPUTool.cuh +++ b/GPUTool.cuh @@ -7,6 +7,8 @@ #include #include +#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