RasterProcessTool/Toolbox/SimulationSARTool/SimulationSAR/GPUTBPImage.cu

#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___

// <20><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>
__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);// У<><D0A3><EFBFBD><EFBFBD>
			cuComplex echotemp = echoArr[prfid * freqcount + Cid];
			imgEchoArr[idx] = echotemp;
			imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echotemp, Rphi));// <20><><EFBFBD><EFBFBD>
			//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; // ÿ<><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_RFPC_SiglePRF blockSize:%d ,numBlock:%d\n",blockSize,numBlocks);
	// <20><><EFBFBD><EFBFBD> CUDA <20>˺<EFBFBD><CBBA><EFBFBD> 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