 							genScanFilter_or_int_leq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
 					}
 
+				} else if (sn->whereAttrType[index] == FLOAT){
+					if(where->andOr == AND){
+						int rel = where->exp[i].relation;
+						float whereValue = *((int*) where->exp[i].content);
+						if(rel==EQ)
+							genScanFilter_and_float_eq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if(rel == GTH)
+							genScanFilter_and_float_gth<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if(rel == LTH)
+							genScanFilter_and_float_lth<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if(rel == GEQ)
+							genScanFilter_and_float_geq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if (rel == LEQ)
+							genScanFilter_and_float_leq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+					}else{
+						int rel = where->exp[i].relation;
+						float whereValue = *((int*) where->exp[i].content);
+						if(rel==EQ)
+							genScanFilter_or_float_eq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if(rel == GTH)
+							genScanFilter_or_float_gth<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if(rel == LTH)
+							genScanFilter_or_float_lth<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if(rel == GEQ)
+							genScanFilter_or_float_geq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+						else if (rel == LEQ)
+							genScanFilter_or_float_leq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
+					}
 				}else{
 					if(where->andOr == AND)
 						genScanFilter_and<<<grid,block>>>(column[index],sn->whereAttrSize[index],sn->whereAttrType[index], totalTupleNum, gpuExp, gpuFilter);

	int blockNum = sn->tn->tupleNum / block.x + 1;

	if(blockNum<1024)
		grid = blockNum; 

	int threadNum = grid.x * block.x;
	long totalTupleNum = sn->tn->tupleNum;

	}
}

__global__ static void genScanFilter_or_float_gth(char *col, long tupleNum, float where, int * filter){
        int stride = blockDim.x * gridDim.x;
        int tid = blockIdx.x * blockDim.x + threadIdx.x;
	int con;

	for(long i = tid; i<tupleNum;i+=stride){
		con = ((float*)col)[i] > where; 
		filter[i] |= con;
	}
}
__global__ static void genScanFilter_or_int_lth(char *col, long tupleNum, int where, int * filter){
        int stride = blockDim.x * gridDim.x;
        int tid = blockIdx.x * blockDim.x + threadIdx.x;
