89 lines
2.4 KiB
Plaintext
89 lines
2.4 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,
|
|
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 = 2 * (R - Rnear) / LIGHTSPEED * fs;
|
|
long Rid = floorf(Ridf);
|
|
RArr[idx] = R;
|
|
if(Rid <0|| Rid >= freqcount){}
|
|
else {
|
|
float factorj = freq * 4 * PI / LIGHTSPEED;
|
|
cuComplex Rfactorj = make_cuComplex(0, factorj * R);
|
|
cuComplex Rphi =cuCexpf(Rfactorj);// 校正项
|
|
cuComplex echotime = echoArr[prfid * freqcount + Rid];
|
|
imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echotime, 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,
|
|
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,
|
|
R,
|
|
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();
|
|
}
|
|
|
|
|
|
|
|
#endif
|
|
|
|
|