#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 = 2 * (R - Rnear) / 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));// 校正项
			imgArr[idx] = cuCaddf(imgArr[idx], cuCmulf(echoArr[Rid] , Rphi));// 矫正
		}
	}
}


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; // 每个块的线程数
	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,
		  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