#include #include #include #include #include #include #include #include #include "BaseConstVariable.h" #include "GPUTool.cuh" #include "GPUTBPImage.cuh" #ifdef __CUDANVCC___ // 定义参数 __device__ cuComplex cuCexpf(cuComplex d) { float factor = exp(d.x); return make_cuComplex(factor * cos(d.y), factor * sin(d.y)); } __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 << > > ( 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