
#include "DFT_GPU_Kernel.h"


//Kernel Functions
/*
Discrete Fourier Transform
dir=1 forward transform
dir=-1 inverse transform
*/
__global__ void DFT_Kernel(float *Dst_real_dev,float *Dst_imag_dev,
	                       float *signal_real_dev,float *signal_imag_dev,int N_DFT,int dir){
	const int ix = blockDim.x * blockIdx.x + threadIdx.x;
	if(ix<N_DFT){
		int k;
		float theta;
		float cos_theta,sin_theta;
		Dst_real_dev[ix] = 0.0f;
		Dst_imag_dev[ix] = 0.0f;
		theta = dir * 2.0f * PI * ix / (float)N_DFT;
		for (k=0;k<N_DFT;k++){
			cos_theta = cosf(k * theta);
			sin_theta = sinf(k * theta);
			Dst_real_dev[ix] += (signal_real_dev[k] * cos_theta + signal_imag_dev[k] * sin_theta);
			Dst_imag_dev[ix] += (signal_imag_dev[k] * cos_theta - signal_real_dev[k] * sin_theta);
		}

		if (dir == 1){
			Dst_real_dev[ix]/=(float)N_DFT;
			Dst_imag_dev[ix]/=(float)N_DFT;
		}
	}
}

void DFT_GPU(float *Dst_real_dev,float *Dst_imag_dev,float *signal_real_dev,float *signal_imag_dev,int N_DFT,int dir,dim3 grid,dim3 threads){
	DFT_Kernel<<<grid,threads>>>(Dst_real_dev,Dst_imag_dev,signal_real_dev,signal_imag_dev,N_DFT,dir);
}