81 lines
2.1 KiB
Plaintext
81 lines
2.1 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,
|
|||
|
|
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));// У<><D0A3><EFBFBD><EFBFBD>
|
|||
|
|
imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echoArr[Rid] , Rphi));// <20><><EFBFBD><EFBFBD>
|
|||
|
|
}
|
|||
|
|
}
|
|||
|
|
}
|
|||
|
|
|
|||
|
|
|
|||
|
|
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)
|
|||
|
|
{
|
|||
|
|
int blockSize = 256; // ÿ<><C3BF><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>߳<EFBFBD><DFB3><EFBFBD>
|
|||
|
|
int numBlocks = (rowcount * colcount + blockSize - 1) / blockSize; // <20><><EFBFBD><EFBFBD> pixelcount <20><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>С
|
|||
|
|
//printf("\nCUDA_RTPC_SiglePRF blockSize:%d ,numBlock:%d\n",blockSize,numBlocks);
|
|||
|
|
// <20><><EFBFBD><EFBFBD> CUDA <20>˺<EFBFBD><CBBA><EFBFBD> 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();
|
|||
|
|
}
|
|||
|
|
|
|||
|
|
|
|||
|
|
|
|||
|
|
#endif
|
|||
|
|
|
|||
|
|
|