#include <assert.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include "cuPrintf.cu"
#include "cuPrintf.cuh"
#include <stdio.h>
#include "../visualization/colormap.h"

#define DIM_X 400
#define DIM_Y 200
typedef unsigned char uchar;
//surface<void, 2> texOut; ///// maybe just do a normal array instead of a surface.
		//we may not need a surface at all.
//texture<float, cudaTextureType2D, cudaReadModeElementType> texTemplate 
texture<uchar, cudaTextureType2D, cudaReadModeElementType> texIn;
float *result;
float* v_dif;
cudaArray* srcArray;

inline void checkCUDAerrors(const char *msg)
{
	cudaError_t err = cudaGetLastError();
	if (cudaSuccess != err){
		fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err) );
		exit(1);
		}
}


float get_sum(float *diff)
{

	cublasStatus_t ret;
	cublasHandle_t handle;
	ret = cublasCreate(&handle);
	
	ret = cublasSetVector(DIM_X*DIM_Y, sizeof(*diff), diff, 1, v_dif, 1);
	//float* out = (float*) malloc(sizeof(float));
	float out;
	ret = cublasSasum(handle, DIM_X*DIM_Y, v_dif, 1, &out);
	cublasDestroy(handle);
	cudaFree(v_dif);
	return out;
}

__device__ float Template(int x)
{
	if(x < DIM_X/6 || x > DIM_X*5/6 || (x > DIM_X*2/6 && x < DIM_X*4/6)){
		return 1.0;
	}else{
		return 0.0;
	}

}

__global__
void get_diff (float *result)
{	
	//cuPrintf("Hello");
	int x 	= threadIdx.x + blockIdx.x * blockDim.x;
	int y 	= threadIdx.y + blockIdx.y * blockDim.y;
	int idx = y*DIM_X+x;	
	//idx 	= x*DIM_Y+y;
	//int idx = x*DIM_X+y;

 	//uchar4 color 		= tex2D(texIn, x, y);
	//float3 tempcolor	= make_float3(color.x, color.y, color.z);
	//float valIn		= tempcolor.x + tempcolor.y + tempcolor.z;
	float valIn		= tex2D(texIn, x, y)/255.0;
	float valTemp		= Template(x);
	result[idx] 		= abs(valIn-valTemp);
	//-valTemp;
// - valTemp;
//	#if __CUDA_ARCH__>=200
//		printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
//	#endif
	//cuPrintf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
}




void initArray(cudaGraphicsResource_t src)
{
	//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar> ();
	//cudaMallocArray(&result, &channelDesc, DIM_X, DIM_Y, 0);
	//HANDLE_ERROR(
	//	cudaGraphicsGLRegisterImage(&src,
	//		fboID,
	//		GL_TEXTURE_2D,
	HANDLE_ERROR(
		cudaGraphicsMapResources(1, &src)	
	);
	HANDLE_ERROR(
		cudaGraphicsSubResourceGetMappedArray(&srcArray, src,0,0)
		);
	HANDLE_ERROR(
		cudaBindTextureToArray(texIn, srcArray)
		);
	cudaMalloc( (void**) &result, DIM_X*DIM_Y*sizeof(float));
        checkCUDAerrors("Memory Allocation Issue 1");	
	cudaMalloc((void **) &v_dif, DIM_X*DIM_Y*sizeof(*result));
        checkCUDAerrors("Memory Allocation Issue 2");	
	//HANDLE_ERROR(
	//	cudaBindTextureToArray(texIn, ptr, &channelDesc)
	//	    );		
}

void cleanUP(cudaGraphicsResource_t src)
{
	HANDLE_ERROR(
		cudaUnbindTexture(texIn)
	);
	HANDLE_ERROR(
		cudaFree(result)
	);
	HANDLE_ERROR(
		cudaGraphicsUnmapResources(1,&src)
	);
}

extern "C"
float get_cost(cudaGraphicsResource_t src)
{
	initArray(src);
	dim3 grid(400, 200);
	dim3 block(1, 1);
	//texIn.normalized = 1;	
	get_diff <<< grid, block >>> (result);
	stim::gpu2image<float>(result, "test2.bmp", 400,200,0,255);
	float output = get_sum(result);
	cleanUP(src);
	return output;
}