#include "heat_wraper.h"

float* CHeatTrasfer::reserveDeviceFloatMatrix(){
	float* t;
	cudaMalloc((void**)&t,x*y*sizeof(float));
	deviceMatrices.push_back(t);
	return t;
}
float* CHeatTrasfer::initConstTemp(){
	float* temp = reserveMatrix(false);
	  for (int i=0; i<DIM*DIM; i++) {
         temp[i] = 0;
         int x = i % DIM;
         int y = i / DIM;
         if ((x>300) && (x<600) && (y>310) && (y<601))
             temp[i] = MAX_TEMP;
    }
    temp[DIM*100+100] = (MAX_TEMP + MIN_TEMP)/2;
    temp[DIM*700+100] = MIN_TEMP;
    temp[DIM*300+300] = MIN_TEMP;
    temp[DIM*200+700] = MIN_TEMP;
    for (int y=800; y<900; y++) {
        for (int x=400; x<500; x++) {
            temp[x+y*DIM] = MIN_TEMP;
        }
    }
	return temp;
}
float* CHeatTrasfer::reserveMatrix(bool device){
	int size = x*y*sizeof(float);
	float* t;
	if(device){
		cudaMalloc((void**)&t,size);
		deviceMatrices.push_back(t);
	}else{
		t= (float*)malloc(size);
		hostMatrices.push_back(t);
	}
	return t;
}

void CHeatTrasfer::cudaCopy(float* src,float* dest,cudaMemcpyKind mem){
	cudaMemcpy(dest,src,x*y*sizeof(float),mem);
}

// texture memmory
texture<float,2> texConst;
texture<float,2> texFront;
texture<float,2> texBack;

CHeatTrasfer::~CHeatTrasfer(){
	// delete deviceMatrices
	float * ptr;
	vector<float*>::const_iterator i;
	// deviceMatrices.pop_back()
		for(i=deviceMatrices.begin(); i!=deviceMatrices.end(); ++i){
			ptr = (*i);
			cudaFree(ptr);
			ptr=NULL;
		// delete ptr
		}
		for(i=hostMatrices.begin(); i!=hostMatrices.end(); ++i){
			ptr = (*i);
			free(ptr);
			ptr=NULL;
		// delete ptr
		}
		cudaUnbindTexture(texConst);
		cudaUnbindTexture(texFront);
		cudaUnbindTexture(texBack);
}



void CHeatTrasfer::setUpTextureMemmory(float* dev_const,float* dev_front,float* dev_back){
	
	cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
	cudaBindTexture2D(NULL,texConst,dev_const,desc,x,y,sizeof(float)*x);
	cudaBindTexture2D(NULL,texFront,dev_front,desc,x,y,sizeof(float)*x);
	cudaBindTexture2D(NULL,texBack,dev_back,desc,x,y,sizeof(float)*x);
	frontTex = dev_front;
	backTex = dev_back;
}


// threads map 1:=1
__global__ void  set(float a,float*data){
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;
	int i = x + y * blockDim.x*gridDim.x;
	data[i] = a;
}



__global__ void copy_const_kernel( float *iptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
	float c = tex2D(texConst,x,y);
    if (c != 0) 
		iptr[offset] = c;
}


__global__ void blend_kernel( float *outSrc,bool swaped) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
	float t,l,c,r,b;
	if(swaped)
	{	
		t = tex2D(texFront,x,y+1);
		l = tex2D(texFront,x-1,y);
		c = tex2D(texFront,x,y);
		r = tex2D(texFront,x+1,y);
		b = tex2D(texFront,x,y-1);
	}
	else{
		t = tex2D(texBack,x,y+1);
		l = tex2D(texBack,x-1,y);
		c = tex2D(texBack,x,y);
		r = tex2D(texBack,x+1,y);
		b = tex2D(texBack,x,y-1);
	}
    outSrc[offset] =   c + SPEED * ( t+l+r+b-c*4);
}
void CHeatTrasfer::setTo(float v,float* dev){
	set<<<getNBlocks(),getNThreads()>>>(v,dev);
}

void CHeatTrasfer::cpy_const(float* dest){
	copy_const_kernel<<<getNBlocks(),getNThreads()>>>(dest);
}


void CHeatTrasfer::blend(float* in,bool swaped){
	blend_kernel<<<getNBlocks(),getNThreads()>>>(in,swaped);
}