 		}
 
 	}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 =


//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;

template<typename stencil_type>
struct is_stencil_true
{
	__host__ __device__
	bool operator()(const thrust::tuple<stencil_type, gdf_valid_iterator::value_type> value)
	{
		return (thrust::get<1>(value) == 1) && (thrust::get<0>(value) != 0);
	}
};


#include <thrust/functional.h>
#include <thrust/execution_policy.h>
#include <thrust/iterator/iterator_adaptor.h>
#include <thrust/iterator/transform_iterator.h>

//std lib
#include <map>
