 
 
 
-__device__ void histogramAdd(hist_type* H, const hist_type * hist_col, const int size){
+__device__ void histogramAdd(hist_type* H, const hist_type * hist_col){
 	int32_t tx = threadIdx.x;
-	for(; tx<size;tx+=blockDim.x){
+	for(; tx<256;tx+=blockDim.x){
 		H[tx]+=hist_col[tx];
 	}
 }
-__device__ void histogramSub(hist_type* H, const hist_type * hist_col, const int size){
+
+__device__ void histogramSub(hist_type* H, const hist_type * hist_col){
 	int32_t tx = threadIdx.x;
-	for(; tx<size;tx+=blockDim.x){
+	for(; tx<256;tx+=blockDim.x){
 		H[tx]-=hist_col[tx];
 	}
 }
 
-__device__ void histogramClear(hist_type* H, const int size){
+__device__ void histogramAddAndSub(hist_type* H, const hist_type * hist_colAdd,const hist_type * hist_colSub){
 	int32_t tx = threadIdx.x;
-	for(; tx<size;tx+=blockDim.x){
+	for(; tx<256;tx+=blockDim.x){
+		H[tx]+=hist_colAdd[tx]-hist_colSub[tx];
+	}
+}
+
+__device__ void histogramMultipleAdd(hist_type* H, const hist_type * hist_col,int histCount){
+	int32_t tx = threadIdx.x;
+	for(; tx<256;tx+=blockDim.x){
+		hist_type temp=H[tx];;
+		for(int i=0; i<histCount; i++)
+		    temp+=hist_col[(i<<8)+tx];
+		H[tx]=temp;
+	}
+}
+
+
+__device__ void histogramClear(hist_type* H){
+	int32_t tx = threadIdx.x;
+	for(; tx<256;tx+=blockDim.x){
 		H[tx]=0;
 	}
 }

}


__device__ void historgramMedian(hist_type* H, const int32_t size, const int32_t medPos, int32_t* retval){
	int32_t tx=threadIdx.x;
	if(tx==0){
		int32_t sum=0;
		*retval=size;
		for(int32_t i=0; i<size; i++){
			sum+=H[i];
			if(sum>=medPos){
				*retval=i;
				return;
			}
		}

	}
}

__device__ void ClearDest(im_type* dest,int32_t rows,int32_t cols){
	int32_t tx=threadIdx.x;
	int32_t totalPixels=rows*cols;
	for(; tx<totalPixels;tx+=blockDim.x)
		dest[tx]=0;
}

/////////////////////////////////////////
/////////////////////////////////////////


int print_u, print_v;

typedef unsigned short hist_type;
typedef unsigned char im_type;

#define MF_IM_SIZE 512
#define MF_HIST_SIZE 256

#include "../lookup.h"


#define CUDA(call, ...) do {                        \
