KJ135
/
BaseLibraryCPP
mirror of http://172.16.0.12:5000/LAMPSARToolSoftware/BaseLibraryCPP.git
4
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

修改RTPC相关代码

main
陈增辉 2024-12-23 10:51:01 +08:00
parent 798e02352f
commit 8bc08ec6f2
10 changed files with 616 additions and 89 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
BaseConstVariable.h
View File

@ -3,7 +3,7 @@
#ifndef BASECONSTVARIABLE_H
#define BASECONSTVARIABLE_H
#define EIGEN_USE_MKL_ALL
#define EIGEN_NO_DEBUG
//#define EIGEN_NO_DEBUG
#define EIGEN_USE_BLAS
@ -14,6 +14,7 @@
#define __CUDANVCC___ // 定义CUDA函数
//#define __PRFDEBUG__
@ -29,7 +30,7 @@
#define r2d 180/3.141592653589793238462643383279
#define d2r 3.141592653589793238462643383279/180
#define LIGHTSPEED 299792458
#define PRECISIONTOLERANCE 1e-9
#define PRECISIONTOLERANCE 1e-6
#define Radians2Degrees(Radians) Radians*PI_180
#define Degrees2Radians(Degrees) Degrees*T180_PI
#define EARTHWE 0.000072292115

BIN
BaseLibraryCPP.rar
View File

Binary file not shown.

45
BaseTool.cpp
View File

@ -33,6 +33,7 @@
#include <gsl/gsl_vector.h> /* 提供了向量结构*/
#include <gsl/gsl_matrix.h>
#include <gsl/gsl_multifit.h>
#include <qcoreapplication.h>
@ -500,7 +501,7 @@ long FindValueInStdVectorLast(std::vector<double>& list, double& insertValue)
long high = list.size() - 1;
long mid = -1;
// 处理边界
if (list[high] <= insertValue)
if (list[high]+ PRECISIONTOLERANCE < insertValue)
{
return -1;
}
@ -584,3 +585,45 @@ QVector<SatellitePos> SatelliteAntPos2SatellitePos(QVector<SatelliteAntPos> pose
}
return antposes;
}
QString getDebugDataPath(QString filename)
{
QString folderName = "debugdata";
QString appDir = QCoreApplication::applicationDirPath();
QString folderpath = JoinPath(appDir, folderName);
if (!QDir(folderpath).exists()) {
QDir(appDir).mkdir(folderName);
}
QString datapath = JoinPath(folderpath, filename);
QFile datafile(datapath);
if (datafile.exists()) {
datafile.remove();
}
return datapath;
}
std::vector<std::string> split(const std::string& str, char delimiter) {
std::vector<std::string> tokens;
std::string token;
std::istringstream tokenStream(str);
while (std::getline(tokenStream, token, delimiter)) {
tokens.push_back(token);
}
return tokens;
}
Eigen::VectorXd linspace(double start, double stop, int num) {
Eigen::VectorXd result(num);
double step = (stop - start) / (num - 1); // 计算步长
for (int i = 0; i < num; ++i) {
result[i] = start + i * step; // 生成等间距数值
}
return result;
}

6
BaseTool.h
View File

@ -111,7 +111,7 @@ QVector<SatelliteAntPos> SatellitePos2SatelliteAntPos(QVector<SatellitePos> pose
QVector<SatellitePos> SatelliteAntPos2SatellitePos(QVector<SatelliteAntPos> poses);
QString getDebugDataPath(QString filename);
std::vector<std::string> split(const std::string& str, char delimiter);
Eigen::VectorXd linspace(double start, double stop, int num);
#endif

18
EchoDataFormat.cpp
View File

@ -249,6 +249,24 @@ QString EchoL0Dataset::getEchoDataFilename()
return GPSPointFilePath;
}
void EchoL0Dataset::initEchoArr(std::complex<double> init0)
{
long blockline = Memory1MB * 2000 / 8 / 2 / this->PlusePoints;
long start = 0;
for (start = 0; start < this->PluseCount; start = start + blockline) {
long templine = start + blockline < this->PluseCount ? blockline : this->PluseCount - start;
std::shared_ptr<std::complex<double>> echotemp = this->getEchoArr(start, templine);
for (long i = 0; i < templine; i++) {
for (long j = 0; j < this->PlusePoints; j++) {
echotemp.get()[i * this->PlusePoints + j] = init0;
}
}
this->saveEchoArr(echotemp, start, templine);
qDebug() << "echo init col : " << start << "\t-\t" << start + templine << "\t/\t" << this->PluseCount;
}
}
// Getter 和 Setter 方法实现
long EchoL0Dataset::getPluseCount() { return this->PluseCount; }
void EchoL0Dataset::setPluseCount(long pulseCount) { this->PluseCount = pulseCount; }

3
EchoDataFormat.h
View File

@ -108,6 +108,9 @@ public:
QString getGPSPointFilename();
QString getEchoDataFilename();
void initEchoArr(std::complex<double> init0);
private: // ²úÆ·Ãû³ÆÉèÖÃ
QString folder;
QString filename;

393
GPUTool.cu
View File

@ -29,7 +29,8 @@ __device__ cuComplex cuCexpf(cuComplex x)
// 定义仿真所需参数
__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);
float sigma= param.p1 + param.p2 * exp(-param.p3 * theta) + param.p4 * cos(param.p5 * theta + param.p6);
return sigma;
}
__device__ CUDAVector GPU_VectorAB(CUDAVector A, CUDAVector B) {
@ -96,6 +97,12 @@ __device__ float GPU_BillerInterpAntPattern(float* antpattern,
float searththeta, float searchphi) {
float stheta = searththeta;
float sphi = searchphi;
if (stheta > 90) {
return 0;
}
else {}
float pthetaid = (stheta - starttheta) / dtheta;//
float pphiid = (sphi - startphi) / dphi;
@ -104,6 +111,7 @@ __device__ float GPU_BillerInterpAntPattern(float* antpattern,
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)
{
@ -124,10 +132,10 @@ __device__ float GPU_BillerInterpAntPattern(float* antpattern,
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); // dB-> 线性
//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)
@ -170,7 +178,7 @@ __global__ void CUDA_B_DistanceA(float* Ax, float* Ay, float* Az, float Bx, floa
__global__ void CUDA_make_VectorA_B(float sX, float sY, float sZ, float* tX, float* tY, float* tZ, float* RstX, float* RstY, float* RstZ, long len) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
RstX[idx] = sX - tX[idx];
RstX[idx] = sX - tX[idx]; // 地面->天
RstY[idx] = sY - tY[idx];
RstZ[idx] = sZ - tZ[idx];
}
@ -221,17 +229,80 @@ __global__ void CUDA_SatelliteAntDirectNormal(float* RstX, float* RstY, float* R
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 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);
float Rx = Xst / RstNorm;
float Ry = Yst / RstNorm;
float Rz = Zst / RstNorm;
float Xx = AntXaxisX / AntXaxisNorm;
float Xy = AntXaxisY / AntXaxisNorm;
float Xz = AntXaxisZ / AntXaxisNorm;
float Yx = AntYaxisX / AntYaxisNorm;
float Yy = AntYaxisY / AntYaxisNorm;
float Yz = AntYaxisZ / AntYaxisNorm;
float Zx = AntZaxisX / AntZaxisNorm;
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);
// 计算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)));
thetaAnt[idx] = ThetaAnt;
phiAnt[idx] = PhiAnt;
float PhiAnt = atanf(Yant / Xant); // -pi/2 ~pi/2
if (abs(Yant) < PRECISIONTOLERANCE) { // X轴上
PhiAnt = 0;
}
else if (abs(Xant) < PRECISIONTOLERANCE) { // Y轴上，原点
if (Yant > 0) {
PhiAnt = PI / 2;
}
else {
PhiAnt = -PI / 2;
}
}
else if (Xant < 0) {
if (Yant > 0) {
PhiAnt = PI + PhiAnt;
}
else {
PhiAnt = -PI+PhiAnt ;
}
}
else { // Xant>0 X 正轴
}
if (isnan(PhiAnt)) {
printf("V=[%f,%f,%f];norm=%f;thetaAnt=%f;phiAnt=%f;\n", Xant, Yant, Zant,Norm, ThetaAnt, PhiAnt);
}
//if (abs(ThetaAnt - 0) < PRECISIONTOLERANCE) {
// PhiAnt = 0;
//}
//else {}
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
// , AntZaxisX, AntZaxisY, AntZaxisZ
// , phiAnt[idx]
// , thetaAnt[idx]
//);
}
}
@ -352,7 +423,6 @@ __global__ void CUDA_RTPC(
if (timeID < 0 || timeID >= Freqnumbers) {
timeID = 0;
amp = 0;
}
else {}
@ -367,6 +437,127 @@ __global__ void CUDA_RTPC(
}
__global__ void CUDA_TBPImage(
float* antPx, float* antPy, float* antPz,
float* imgx, float* imgy, float* imgz,
cuComplex* echoArr, cuComplex* imgArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
//printf("\nidx:\t %d %d %d\n", idx, linecount, plusepoint);
if (idx < rowcount * colcount) {
float R = sqrtf(powf(antPx[prfid] - imgx[idx], 2) + powf(antPy[prfid] - imgy[idx], 2) + powf(antPz[prfid] - imgz[idx], 2));
float Ridf = ((R - Rnear) * 2 / LIGHTSPEED) * fs;
long Rid = floorf(Ridf);
if(Rid <0|| Rid >= freqcount){}
else {
float factorj = freq * 4 * PI / LIGHTSPEED;
cuComplex Rphi =cuCexpf(make_cuComplex(0, factorj * R));// 校正项
imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echoArr[Rid] , Rphi));// 矫正
}
}
}
__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;
FreqID[idx] = timeID;
}
}
__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,
starttheta, startphi, dtheta, dphi, thetapoints, phipoints,
temptheta, tempphi) ;
gain[idx] = antPatternGain;
}
}
//__global__ void Sigma0InterpPixel(long* demcls, float* demslopeangle, CUDASigmaParam* sigma0Paramslist, float* localangle, float* sigma0list, long sigmaparamslistlen, long len)
//{
// long idx = blockIdx.x * blockDim.x + threadIdx.x;
// if (idx < len) {
// long clsid = demcls[idx];
// if(clsid<=)
// sigma0list[idx] = 0;
// }
//}
__global__ void CUDA_InterpSigma(
long* demcls, float* sigmaAmp, float* localanglearr, long len,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen) {
long idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < len) {
long clsid = demcls[idx];
float localangle = localanglearr[idx] * r2d;
CUDASigmaParam tempsigma = sigma0Paramslist[clsid];
//printf("cls:%d;localangle=%f;\n",clsid, localangle);
if (localangle < 0 || localangle >= 90) {
sigmaAmp[idx] = 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
) {
sigmaAmp[idx] = 0;
}
else {
float sigma = GPU_getSigma0dB(tempsigma, localangle);
sigma = powf(10.0, sigma / 10.0);// 后向散射系数
//printf("cls:%d;localangle=%f;sigma0=%f;\n", clsid, localangle, sigma);
sigmaAmp[idx] = sigma;
}
}
}
//错误提示
void checkCudaError(cudaError_t err, const char* msg) {
if (err != cudaSuccess) {
@ -380,6 +571,7 @@ void checkCudaError(cudaError_t err, const char* msg) {
extern "C" void* mallocCUDAHost(long memsize) {
void* ptr;
cudaMallocHost(&ptr, memsize);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
@ -387,6 +579,7 @@ extern "C" void* mallocCUDAHost(long memsize) {
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
return ptr;
}
@ -400,7 +593,7 @@ extern "C" void FreeCUDAHost(void* ptr) {
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
// GPU参数内存声明
@ -414,7 +607,7 @@ extern "C" void* mallocCUDADevice(long memsize) {
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
return ptr;
}
@ -428,7 +621,7 @@ extern "C" void FreeCUDADevice(void* ptr) {
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
// GPU 内存数据转移
@ -443,7 +636,7 @@ extern "C" void HostToDevice(void* hostptr, void* deviceptr, long memsize) {
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
extern "C" void DeviceToHost(void* hostptr, void* deviceptr, long memsize) {
@ -455,7 +648,7 @@ extern "C" void DeviceToHost(void* hostptr, void* deviceptr, long memsize) {
// Possibly: exit(-1) if program cannot continue....
}
#endif // __CUDADEBUG__
cudaDeviceSynchronize();
}
extern "C" void CUDATestHelloWorld(float a,long len) {
@ -463,7 +656,7 @@ extern "C" void CUDATestHelloWorld(float a,long len) {
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_Test_HelloWorld << <blockSize, numBlocks >> > (a, len);
CUDA_Test_HelloWorld << <numBlocks, blockSize >> > (a, len);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
@ -474,12 +667,53 @@ extern "C" void CUDATestHelloWorld(float a,long len) {
cudaDeviceSynchronize();
}
void CUDATBPImage(float* antPx, float* antPy, float* antPz,
float* imgx, float* imgy, float* imgz,
cuComplex* echoArr, cuComplex* imgArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount)
{
int blockSize = 256; // 每个块的线程数
int numBlocks = (rowcount * colcount + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
//printf("\nCUDA_RTPC_SiglePRF blockSize:%d ,numBlock:%d\n",blockSize,numBlocks);
// 调用 CUDA 核函数 CUDA_RTPC_Kernel
CUDA_TBPImage << <numBlocks, blockSize >> > (
antPx, antPy, antPz,
imgx, imgy, imgz,
echoArr, imgArr,
freq, fs, Rnear, Rfar,
rowcount, colcount,
prfid, freqcount
);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDATBPImage 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);
CUDA_DistanceAB << <numBlocks, blockSize >> > (Ax, Ay, Az, Bx, By, Bz, R, 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 BdistanceAs(float* Ax, float* Ay, float* Az, float Bx, float By, float Bz, float* R, long len) {
@ -487,7 +721,15 @@ extern "C" void BdistanceAs(float* Ax, float* Ay, float* Az, float Bx, float By,
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_B_DistanceA << <blockSize, numBlocks >> > (Ax, Ay, Az, Bx, By, Bz, R, len);
CUDA_B_DistanceA << <numBlocks, blockSize >> > (Ax, Ay, Az, Bx, By, Bz, R, 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();
}
@ -496,7 +738,15 @@ 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 << <numBlocks, blockSize >> > (sX, sY, sZ, tX, tY, tZ, RstX, RstY, RstZ, 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();
}
@ -505,7 +755,15 @@ 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 << <numBlocks, blockSize >> > (Vx, Vy, Vz, R, 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();
}
@ -513,7 +771,15 @@ extern "C" void cosAngle_VA_AB(float* Ax, float* Ay, float* Az, float* Bx, float
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_cosAngle_VA_AB << <blockSize, numBlocks >> > (Ax, Ay, Az, Bx, By, Bz, anglecos, len);
CUDA_cosAngle_VA_AB << <numBlocks, blockSize >> > (Ax, Ay, Az, Bx, By, Bz, anglecos, 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();
}
@ -528,25 +794,76 @@ 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,
CUDA_SatelliteAntDirectNormal << <numBlocks, blockSize >> > (RstX, RstY, RstZ,
antXaxisX, antXaxisY, antXaxisZ,
antYaxisX, antYaxisY, antYaxisZ,
antZaxisX, antZaxisY, antZaxisZ,
antDirectX, antDirectY, antDirectZ,
thetaAnt, phiAnt
, 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 AntPatternInterpGain(float* anttheta, float* antphi, float* gain,
float* antpattern, float starttheta, float startphi, float dtheta, float dphi, int thetapoints, int phipoints, long len) {
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
//printf("\nCUDA_RTPC_SiglePRF blockSize:%d ,numBlock:%d\n", blockSize, numBlocks);
CUDA_AntPatternInterpGain << <numBlocks, blockSize >> > ( anttheta,antphi, gain,
antpattern,
starttheta, startphi, dtheta, dphi, thetapoints, phipoints,
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 CUDARTPCPRF(float antPx, long len) {
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
printf("\nCUDA_RTPC_SiglePRF blockSize:%d ,numBlock:%d\n", blockSize, numBlocks);
CUDA_Test_HelloWorld << <blockSize, numBlocks >> > (antPx, len);
CUDA_Test_HelloWorld << <numBlocks, blockSize >> > (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 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)
{
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_calculationEcho << <numBlocks, blockSize >> > (sigma0, TransAnt, ReciveAnt,
localangle, R, slopeangle,
nearRange, Fs, pt, lamda, FreqIDmax,
echoAmp, FreqID,
len);
#ifdef __CUDADEBUG__
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
@ -610,7 +927,25 @@ CUDA_RTPC << <numBlocks, blockSize >> > (
}
extern "C" void CUDAInterpSigma(
long* demcls,float* sigmaAmp, float* localanglearr,long len,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen) {// 地表覆盖类型-sigma插值对应函数-ulaby
int blockSize = 256; // 每个块的线程数
int numBlocks = (len + blockSize - 1) / blockSize; // 根据 pixelcount 计算网格大小
// 调用 CUDA 核函数
CUDA_InterpSigma << <numBlocks, blockSize >> > (
demcls, sigmaAmp, localanglearr, len,
sigma0Paramslist, sigmaparamslistlen
);
#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

38
GPUTool.cuh
View File

@ -58,6 +58,9 @@ 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 AntPatternInterpGain(float* anttheta, float* antphi, float* gain, float* antpattern, float starttheta, float startphi, float dtheta, float dphi, int thetapoints, int phipoints,long len);
extern "C" void CUDA_RTPC_SiglePRF(
float antPx, float antPy, float antPZ,
float antXaxisX, float antXaxisY, float antXaxisZ,
@ -75,22 +78,41 @@ extern "C" void CUDA_RTPC_SiglePRF(
);
extern "C" void CUDARTPCPRF(float antPx, long len);
extern "C" void CUDATestHelloWorld(float a, long len);
extern "C" void CUDATBPImage(
float* antPx,
float* antPy,
float* antPz,
float* imgx,
float* imgy,
float* imgz,
cuComplex* echoArr,
cuComplex* imgArr,
float freq, float fs, float Rnear, float Rfar,
long rowcount, long colcount,
long prfid, long freqcount
);
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 CUDAInterpSigma(
long* demcls, float* sigmaAmp, float* localanglearr, long len,
CUDASigmaParam* sigma0Paramslist, long sigmaparamslistlen);
#endif
#endif
/**
*

30
ImageOperatorBase.cpp
View File

@ -819,8 +819,6 @@ void gdalImage::saveImage(Eigen::MatrixXd data, int start_row = 0, int start_col
}
int datarows = data.rows();
int datacols = data.cols();
@ -1174,7 +1172,7 @@ RasterExtend gdalImage::getExtend()
}
gdalImage CreategdalImage(const QString& img_path, int height, int width, int band_num,
Eigen::MatrixXd gt, QString projection, bool need_gt, bool overwrite)
Eigen::MatrixXd gt, QString projection, bool need_gt, bool overwrite, bool isEnvi)
{
if(exists_test(img_path.toUtf8().constData())) {
if(overwrite) {
@ -1187,11 +1185,21 @@ gdalImage CreategdalImage(const QString& img_path, int height, int width, int ba
}
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES"); // 注锟斤拷锟绞斤拷锟斤拷锟斤拷锟?1锟?7
GDALDriver* poDriver = GetGDALDriverManager()->GetDriverByName("GTiff");
GDALDriver* poDriver = nullptr;
if (isEnvi) {
poDriver = GetGDALDriverManager()->GetDriverByName("ENVI");
}
else {
poDriver = GetGDALDriverManager()->GetDriverByName("GTiff");
}
GDALDataset* poDstDS = poDriver->Create(img_path.toUtf8().constData(), width, height, band_num,
GDT_Float32, NULL); // 锟斤拷锟斤拷锟斤拷
if(need_gt) {
if (!projection.isEmpty()) {
poDstDS->SetProjection(projection.toUtf8().constData());
}
double gt_ptr[6] = { 0 };
for(int i = 0; i < gt.rows(); i++) {
for(int j = 0; j < gt.cols(); j++) {
@ -1932,6 +1940,18 @@ ErrorCode MergeRasterInGeoCoding(QVector<gdalImage> imgdslist, gdalImage resulti
return ErrorCode::SUCCESS;
}
bool saveEigenMatrixXd2Bin(Eigen::MatrixXd data, QString dataStrPath)
{
Eigen::MatrixXd gt = Eigen::MatrixXd::Zero(2, 3);
gdalImage img=CreategdalImage(dataStrPath, data.rows(), data.cols(), 1, gt, "", false, false,true);
img.saveImage(data, 0,0,1);
return true;
}
gdalImageComplex::gdalImageComplex(const QString& raster_path)
{
omp_lock_t lock;
@ -1942,7 +1962,7 @@ gdalImageComplex::gdalImageComplex(const QString& raster_path)
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
GDALDataset* rasterDataset = (GDALDataset*)(GDALOpen(
raster_path.toUtf8().constData(), GA_ReadOnly)); // 锟斤拷只斤拷式锟斤拷取斤拷影锟斤拷
raster_path.toUtf8().constData(), GA_ReadOnly));
this->width = rasterDataset->GetRasterXSize();
this->height = rasterDataset->GetRasterYSize();
this->band_num = rasterDataset->GetRasterCount();

99
ImageOperatorBase.h
View File

@ -226,7 +226,7 @@ public:
};
// 创建影像
gdalImage CreategdalImage(const QString& img_path, int height, int width, int band_num, Eigen::MatrixXd gt, QString projection, bool need_gt = true, bool overwrite = false);
gdalImage CreategdalImage(const QString& img_path, int height, int width, int band_num, Eigen::MatrixXd gt, QString projection, bool need_gt = true, bool overwrite = false, bool isEnvi = false);
gdalImage CreategdalImage(const QString& img_path, int height, int width, int band_num, Eigen::MatrixXd gt, long espgcode, GDALDataType eType=GDT_Float32, bool need_gt = true, bool overwrite = false,bool isENVI=false);
gdalImageComplex CreategdalImageComplex(const QString& img_path, int height, int width, int band_num, Eigen::MatrixXd gt, QString projection, bool need_gt = true, bool overwrite = false);
@ -268,12 +268,16 @@ ErrorCode MergeRasterProcess(QVector<QString> filepath, QString outfileptah, QSt
ErrorCode MergeRasterInGeoCoding(QVector<gdalImage> inimgs, gdalImage resultimg,gdalImage maskimg, ShowProessAbstract* dia = nullptr);
// 保存矩阵转换为envi文件，默认数据格式为double
bool saveEigenMatrixXd2Bin(Eigen::MatrixXd data, QString dataStrPath);
template<typename T>
std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, int rows_count, int cols_count, int band_ids, GDALREADARRCOPYMETHOD method)
{
std::shared_ptr<T> result = nullptr;
omp_lock_t lock;
omp_init_lock(&lock);
omp_set_lock(&lock);
@ -303,7 +307,6 @@ std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, i
}
}
delete[] temp;
}
else if (gdal_datatype == GDT_UInt16) {
@ -416,7 +419,81 @@ std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, i
}
delete[] temp;
}
else if (gdal_datatype == GDT_CFloat32) {
//else if (gdal_datatype == GDT_CFloat32) {
// if (std::is_same<T, std::complex<double>>::value || std::is_same<T, std::complex<float>>::value) {
// float* temp = new float[rows_count * cols_count * 2];
// demBand->RasterIO(GF_Read, start_col, start_row, cols_count, rows_count, temp, cols_count, rows_count, gdal_datatype, 0, 0);
// result = std::shared_ptr<T>(new T[rows_count * cols_count], delArrPtr);
// if (method == GDALREADARRCOPYMETHOD::MEMCPYMETHOD) {
// std::memcpy(result.get(), temp, rows_count * cols_count);
// }
// else if (method == GDALREADARRCOPYMETHOD::VARIABLEMETHOD) {
// long count = rows_count * cols_count;
// for (long i = 0; i < count; i++) {
// result.get()[i] = T(temp[i * 2],
// temp[i * 2 + 1]);
// }
// }
// delete[] temp;
// }
// else {
// result = nullptr;
// }
//}
//else if (gdal_datatype == GDT_CFloat64 ) {
// if (std::is_same<T, std::complex<double>>::value || std::is_same<T, std::complex<float>>::value) {
// double* temp = new double[rows_count * cols_count * 2];
// demBand->RasterIO(GF_Read, start_col, start_row, cols_count, rows_count, temp, cols_count, rows_count, gdal_datatype, 0, 0);
// result = std::shared_ptr<T>(new T[rows_count * cols_count], delArrPtr);
// if (method == GDALREADARRCOPYMETHOD::MEMCPYMETHOD) {
// std::memcpy(result.get(), temp, rows_count * cols_count);
// }
// else if (method == GDALREADARRCOPYMETHOD::VARIABLEMETHOD) {
// long count = rows_count * cols_count;
// for (long i = 0; i < count; i++) {
// result.get()[i] = T(temp[i * 2],
// temp[i * 2 + 1]);
// }
// }
// delete[] temp;
// }
// else {
// result = nullptr;
// }
//}
else {
}
GDALClose((GDALDatasetH)rasterDataset);
omp_unset_lock(&lock); // 锟酵放伙拷斤拷
omp_destroy_lock(&lock); // 劫伙拷斤拷
// GDALDestroy(); // or, DllMain at DLL_PROCESS_DETACH
return result;
}
template<typename T>
std::shared_ptr<T> readDataArrComplex(gdalImageComplex& imgds, int start_row, int start_col, int rows_count, int cols_count, int band_ids, GDALREADARRCOPYMETHOD method)
{
std::shared_ptr<T> result = nullptr;
omp_lock_t lock;
omp_init_lock(&lock);
omp_set_lock(&lock);
GDALAllRegister();
CPLSetConfigOption("GDAL_FILENAME_IS_UTF8", "YES");
GDALDataset* rasterDataset = (GDALDataset*)(GDALOpen(imgds.img_path.toUtf8().constData(), GA_ReadOnly)); // 锟斤拷只斤拷式锟斤拷取斤拷影锟斤拷
GDALDataType gdal_datatype = rasterDataset->GetRasterBand(1)->GetRasterDataType();
GDALRasterBand* demBand = rasterDataset->GetRasterBand(band_ids);
rows_count = start_row + rows_count <= imgds.height ? rows_count : imgds.height - start_row;
cols_count = start_col + cols_count <= imgds.width ? cols_count : imgds.width - start_col;
//Eigen::MatrixXd datamatrix(rows_count, cols_count);
if (gdal_datatype == GDT_CFloat32) {
if (std::is_same<T, std::complex<double>>::value || std::is_same<T, std::complex<float>>::value) {
float* temp = new float[rows_count * cols_count * 2];
demBand->RasterIO(GF_Read, start_col, start_row, cols_count, rows_count, temp, cols_count, rows_count, gdal_datatype, 0, 0);
@ -427,13 +504,18 @@ std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, i
else if (method == GDALREADARRCOPYMETHOD::VARIABLEMETHOD) {
long count = rows_count * cols_count;
for (long i = 0; i < count; i++) {
result.get()[i] = std::complex<float>(temp[i * 2 ],
result.get()[i] = T(temp[i * 2],
temp[i * 2 + 1]);
}
}
delete[] temp;
}
else {
result = nullptr;
}
}
else if (gdal_datatype == GDT_CFloat64) {
if (std::is_same<T, std::complex<double>>::value || std::is_same<T, std::complex<float>>::value) {
double* temp = new double[rows_count * cols_count * 2];
demBand->RasterIO(GF_Read, start_col, start_row, cols_count, rows_count, temp, cols_count, rows_count, gdal_datatype, 0, 0);
@ -444,12 +526,16 @@ std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, i
else if (method == GDALREADARRCOPYMETHOD::VARIABLEMETHOD) {
long count = rows_count * cols_count;
for (long i = 0; i < count; i++) {
result.get()[i] = std::complex<double>(temp[i * 2],
result.get()[i] = T(temp[i * 2],
temp[i * 2 + 1]);
}
}
delete[] temp;
}
else {
result = nullptr;
}
}
else {
}
GDALClose((GDALDatasetH)rasterDataset);
@ -460,7 +546,6 @@ std::shared_ptr<T> readDataArr(gdalImage& imgds, int start_row, int start_col, i
}
//--------------------- 图像分块 ------------------------------