 		filter[i] |= con;
 	}
 }
+__global__ static void genScanFilter_or_float_lth(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_geq(char *col, long tupleNum, int where, int * filter){
         int stride = blockDim.x * gridDim.x;
         int tid = blockIdx.x * blockDim.x + threadIdx.x;

					genScanFilter_or_int_geq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
				else if (rel == LEQ)
					genScanFilter_or_int_leq<<<grid,block>>>(column[index],totalTupleNum, whereValue, gpuFilter);
			}else
				genScanFilter_or<<<grid,block>>>(column[index],sn->whereAttrSize[index],sn->whereAttrType[index], totalTupleNum, gpuExp, gpuFilter);



	for(long i = tid; i<tupleNum;i+=stride){
		con = ((int*)col)[i] <= where; 
		filter[i] |= con;
	}
}
__global__ static void genScanFilter_or_float_leq(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;
	}
}
