 	bit_position_iterator bit_position_iter(thrust::make_counting_iterator<gdf_size_type>(0), modulus_bit_width());
 	gdf_valid_iterator valid_iterator(thrust::detail::make_normal_iterator(thrust::device_pointer_cast(stencil->valid)),GDF_VALID_BITSIZE);
 	//TODO: can probably make this happen with some kind of iterator so it can work on any width size
- 
+
 	//zip the stencil and the valid iterator together
 	typedef thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int8_t> >,gdf_valid_iterator, bit_position_iterator > zipped_stencil_tuple;
 	typedef thrust::zip_iterator<zipped_stencil_tuple> zipped_stencil_iterator;

	}
};

/*
damn im passing out, ok how im going to do this, ill derefence the fucking entire character, offsetting with the iterator i make however the fuck that works right?
		then i will make it so that i zip a counting iterator as well and use that counting iterator to figure out which bit to and

				thrust::copy_if(thrust::cuda::par.on(stream),input_start,input_start + lhs->size,zipped_stencil_iter,output_start,is_stencil_true<thrust::detail::normal_iterator<thrust::device_ptr<int8_t> >::value_type >());
		output->size = output_end - output_start;
	}

	gdf_size_type num_values = lhs->size;
	//TODO:BRING OVER THE BITMASK!!!
	//need to store a prefix sum
	//align to size 8
	thrust::device_vector<gdf_valid_type> valid_bit_mask; //we are expanding the bit mask to an int8 because I can't envision an algorithm that operates on the bitmask that
	if(num_values % GDF_VALID_BITSIZE != 0){
		valid_bit_mask.resize(num_values + (GDF_VALID_BITSIZE - (num_values % GDF_VALID_BITSIZE))); //align this allocation on GDF_VALID_BITSIZE so we don't have to bounds check
	}else{
		valid_bit_mask.resize(num_values);
	}

	// doesn't require the use for a prefix sum which will have size 8 * num rows which is much larger than this

	typedef thrust::tuple<gdf_valid_iterator, bit_position_iterator > mask_tuple;
	typedef thrust::zip_iterator<mask_tuple> zipped_mask;


	zipped_mask  zipped_mask_iter(
			thrust::make_tuple(
					valid_iterator,
					thrust::make_transform_iterator<modulus_bit_width, thrust::counting_iterator<gdf_size_type> >(
							thrust::make_counting_iterator<gdf_size_type>(0),
							modulus_bit_width())
			)
	);

	typedef thrust::transform_iterator<is_bit_set, zipped_mask > bit_set_iterator;
	bit_set_iterator bit_set_iter = thrust::make_transform_iterator<is_bit_set,zipped_mask>(
			zipped_mask_iter,
			is_bit_set()
	);

	//copy the bitmask to device_vector of int8
	thrust::copy(thrust::cuda::par.on(stream), bit_set_iter, bit_set_iter + num_values, valid_bit_mask.begin());

	//remove the values that don't pass the stencil
	thrust::remove_if(thrust::cuda::par.on(stream),valid_bit_mask.begin(), valid_bit_mask.begin() + num_values,zipped_stencil_iter, is_stencil_true<thrust::detail::normal_iterator<thrust::device_ptr<int8_t> >::value_type >());

	//recompact the values and store them in the output bitmask
	//we can group them into pieces of 8 because we aligned this earlier on when we made the device_vector
	thrust::detail::normal_iterator<thrust::device_ptr<int64_t> > valid_bit_mask_group_8_iter =
			thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int64_t *) valid_bit_mask.data().get()));


	//you may notice that we can write out more bytes than our valid_num_bytes, this only happens when we are not aligned to  GDF_VALID_BITSIZE bytes, becasue the
	//arrow standard requires 64 byte alignment, this is a safe assumption to make
	thrust::transform(thrust::cuda::par.on(stream), valid_bit_mask_group_8_iter, valid_bit_mask_group_8_iter + ((num_values + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE),
			thrust::detail::make_normal_iterator(thrust::device_pointer_cast(output->valid)),bit_mask_pack_op());

	cudaStreamSynchronize(stream);

