108 lines
2.6 KiB
Plaintext
108 lines
2.6 KiB
Plaintext
|
|
|
|
#include <iostream>
|
|
#include <memory>
|
|
#include <cmath>
|
|
#include <complex>
|
|
#include <device_launch_parameters.h>
|
|
#include <cuda_runtime.h>
|
|
#include <cublas_v2.h>
|
|
#include <cuComplex.h>
|
|
|
|
#include "BaseConstVariable.h"
|
|
#include "GPUTool.cuh"
|
|
#include "GPUTBPImage.cuh"
|
|
|
|
#ifdef __CUDANVCC___
|
|
|
|
|
|
|
|
__global__ void CUDA_TBPImage(
|
|
float* antPx, float* antPy, float* antPz,
|
|
float* imgx, float* imgy, float* imgz,
|
|
float* RArr,
|
|
long* Cids,
|
|
cuComplex* echoArr,
|
|
cuComplex* imgArr,
|
|
cuComplex* imgEchoArr,
|
|
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 px = antPx[prfid];
|
|
float py = antPy[prfid];
|
|
float pz = antPz[prfid];
|
|
|
|
float tx = imgx[idx];
|
|
float ty = imgy[idx];
|
|
float tz = imgz[idx];
|
|
|
|
float R = sqrtf((px-tx) * (px - tx) + (py-ty) * (py-ty) + (pz-tz) * (pz-tz));
|
|
float Cidf= 2 * (R - Rnear) / LIGHTSPEED * fs;
|
|
long Cid = floorf(Cidf);
|
|
RArr[idx] = R;
|
|
Cids[idx] = Cid;
|
|
if(Cid <0|| Cid >= freqcount){}
|
|
else {
|
|
float factorj = freq * 4 * PI / LIGHTSPEED;
|
|
cuComplex Rfactorj = make_cuComplex(0, factorj * R);
|
|
cuComplex Rphi =cuCexpf(Rfactorj);// 校正项
|
|
cuComplex echotemp = echoArr[prfid * freqcount + Cid];
|
|
imgEchoArr[idx] = echotemp;
|
|
imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echotemp, Rphi));// 矫正
|
|
//printf("R=%f;Rid=%d;factorj=%f;Rfactorj=complex(%f,%f);Rphi=complex(%f,%f);\n", R, Rid, factorj,
|
|
// Rfactorj.x, Rfactorj.y,
|
|
// Rphi.x, Rphi.y);
|
|
}
|
|
}
|
|
}
|
|
|
|
|
|
extern "C" void CUDATBPImage(float* antPx, float* antPy, float* antPz,
|
|
float* imgx, float* imgy, float* imgz,
|
|
float* R,
|
|
long* Cids,
|
|
cuComplex* echoArr,
|
|
cuComplex* imgArr,
|
|
cuComplex* imgEchoArr,
|
|
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_RFPC_SiglePRF blockSize:%d ,numBlock:%d\n",blockSize,numBlocks);
|
|
// 调用 CUDA 核函数 CUDA_RFPC_Kernel
|
|
|
|
CUDA_TBPImage << <numBlocks, blockSize >> > (
|
|
antPx, antPy, antPz,
|
|
imgx, imgy, imgz,
|
|
R, Cids,
|
|
echoArr, imgArr, imgEchoArr,
|
|
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));
|
|
exit(2);
|
|
}
|
|
#endif // __CUDADEBUG__
|
|
cudaDeviceSynchronize();
|
|
}
|
|
|
|
|
|
|
|
#endif
|
|
|
|
|