 		}
 
 	}else if(stencil_width == 2){
-		typedef thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int16_t> >,gdf_valid_iterator> zipped_stencil_tuple;
+		typedef thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int16_t> >,gdf_valid_iterator, bit_position_iterator > zipped_stencil_tuple;
 		typedef thrust::zip_iterator<zipped_stencil_tuple> zipped_stencil_iterator;
 
-		zipped_stencil_iterator zipped_stencil_iter(thrust::make_tuple(thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int16_t * )stencil->data)),valid_iterator));
+
+		zipped_stencil_iterator zipped_stencil_iter(
+				thrust::make_tuple(
+						thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int16_t * )stencil->data)),
+						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())
+				));
 
 		if(width == 1){
 			thrust::detail::normal_iterator<thrust::device_ptr<int8_t> > input_start =

			output->size = output_end - output_start;
		}
	}else if(stencil_width == 4){
		typedef thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<int32_t> >,gdf_valid_iterator> zipped_stencil_tuple;
		typedef thrust::zip_iterator<zipped_stencil_tuple> zipped_stencil_iterator;

		zipped_stencil_iterator zipped_stencil_iter(thrust::make_tuple(thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int32_t * )stencil->data)),valid_iterator));

		if(width == 1){
			thrust::detail::normal_iterator<thrust::device_ptr<int8_t> > input_start =


//wow the freaking example from iterator_adaptpr, what a break right!
template<typename Iterator>
class repeat_iterator
		: public thrust::iterator_adaptor<
		  repeat_iterator<Iterator>, // the first template parameter is the name of the iterator we're creating
		  Iterator                   // the second template parameter is the name of the iterator we're adapting
		  // we can use the default for the additional template parameters
		  >
{
public:
	// shorthand for the name of the iterator_adaptor we're deriving from
	typedef thrust::iterator_adaptor<
			repeat_iterator<Iterator>,
			Iterator
			> super_t;
	__host__ __device__
	repeat_iterator(const Iterator &x, int n) : super_t(x), begin(x), n(n) {}
	// befriend thrust::iterator_core_access to allow it access to the private interface below
	friend class thrust::iterator_core_access;
private:
	// repeat each element of the adapted range n times
	unsigned int n;
	// used to keep track of where we began
	const Iterator begin;
	// it is private because only thrust::iterator_core_access needs access to it
	__host__ __device__
	typename super_t::reference dereference() const
	{
		return *(begin + (this->base() - begin) / n);
	}
};



typedef repeat_iterator<thrust::detail::normal_iterator<thrust::device_ptr<gdf_valid_type> > > gdf_valid_iterator;



// note: functor inherits from unary_function
struct modulus_bit_width : public thrust::unary_function<gdf_size_type,gdf_size_type>
{
	__host__ __device__
	gdf_size_type operator()(gdf_size_type x) const
	{
		return x % GDF_VALID_BITSIZE;
	}
};

typedef thrust::transform_iterator<modulus_bit_width, thrust::counting_iterator<gdf_size_type> > bit_position_iterator;


template<typename stencil_type>
struct is_stencil_true
{
	__host__ __device__
	bool operator()(const thrust::tuple<stencil_type, gdf_valid_iterator::value_type, bit_position_iterator::value_type> value)
	{
		gdf_size_type position = thrust::get<2>(value);

		return ((thrust::get<1>(value) >> position) & 1) && (thrust::get<0>(value) != 0);
	}
};

