Using strided iterator to find unique arbitrarily-size tuples in cuda : I'm trying to write an kernel that finds unique tuples in a collection, but with one several important caveat: the application should be able to handle tuples of arbitrary dimensions i.e. unknown at compile-time. Furthermore, the size of the tuple may exceed 10. These two considerations rule out the thrust tuple. As such I've tried to represent the tuple in structure of array (SoA) form using the stride iterator implementation shamelessly ripped off from [here][1], as shown below:

	#include <vector>
	#include <iostream>
	#include <thrust/unique.h>
	#include <thrust/host_vector.h>
	#include <thrust/device_vector.h>
	#include <thrust/iterator/counting_iterator.h>
	#include <thrust/iterator/transform_iterator.h>
	#include <thrust/iterator/permutation_iterator.h>

	typedef unsigned short ushort;

	template<typename Iterator>
	struct StridedIterator
	{
		typedef typename thrust::iterator_difference<Iterator>::type difference_type;

		struct StrideFunctor : public thrust::unary_function<difference_type,difference_type>
		{
			difference_type m_stride;

			StrideFunctor(difference_type stride ): m_stride( stride )
			{
			}

			__host__ __device__
			difference_type operator()(const difference_type& i) const
			{ 
				return m_stride * i;
			}
		};

		typedef typename thrust::counting_iterator<difference_type>                   CountingIterator;
		typedef typename thrust::transform_iterator<StrideFunctor, CountingIterator>  TransformIterator;
		typedef typename thrust::permutation_iterator<Iterator,TransformIterator>     PermutationIterator;

		typedef PermutationIterator Type;

		StridedIterator( Iterator startIter , Iterator endIter , difference_type stride )
			: m_startIter( startIter )
			, m_endIter( endIter )
			, m_stride( stride )
		{
		}

		Type begin() const
		{
		    return PermutationIterator( m_startIter , TransformIterator( CountingIterator(0) , StrideFunctor( m_stride ) ) );
		}

		Type end() const
		{
		    return begin() + ( ( m_endIter - m_startIter ) + ( m_stride - 1 ) ) / m_stride;
		}

		Iterator 		m_startIter;
		Iterator 		m_endIter;
		difference_type m_stride;
	};

	int main()
	{
		// AoS version
		/*
		2 , 1 ,   -- unique
		1 , 2 ,   -- unique
		1 , 1 ,   -- unique
		1 , 1 ,
		1 , 3 ,   -- unique
		2 , 2 ,   -- unique
		3 , 1 ,   -- unique
		2 , 1 ,
		1 , 2 ,
		1 , 1 ,
		2 , 2 ,
		2 , 2
		*/

		// 'SoA' -- twelve 2-tuples, 6 of which are unique
		std::vector<ushort> input{
							2 , 1 , 1 , 1 , 1 , 2 , 3 , 2 , 1 , 1 ,   2 , 2 , 
							1 , 2 , 1 , 1 , 3 , 2 , 1 , 1 , 2 , 1 ,   2 , 2
						};

		thrust::host_vector<ushort> hData = input;
		ushort dim = 2;

		thrust::device_vector<ushort> data = hData;
		typedef thrust::device_vector<ushort>::iterator Iterator;

		StridedIterator<Iterator> stridedIter( data.begin() , data.end() , dim );
		auto iter = thrust::unique( stridedIter.begin() , stridedIter.end() );

		std::cout << stridedIter.end() - stridedIter.begin() << std::endl;
		std::cout << iter - stridedIter.begin() << std::endl;
	}

Unfortunately instead of getting the value of 6 post call to `thrust::unique` I'm getting 9. Please advise.



  [1]: https://stackoverflow.com/a/42235487/181783