Search code examples
c++cudaiteratorthrust

C++ CUDA Thrust how to make an iterator that contains all even indices then all odd indices?


For example, if the array has 6 elements, this iterator would like like: [0, 2, 4, 1, 3, 5]

I know about counting_iterator and constant_iterator, but how can I create the above iterator? I think it has something to do with transform_iterator, but I don't know how.


Solution

  • Here is one possible approach. You could use a permutation_iterator with a transform_iterator inside, something like this:

    $ cat t87.cu
    #include <thrust/iterator/permutation_iterator.h>
    #include <thrust/iterator/transform_iterator.h>
    #include <thrust/iterator/counting_iterator.h>
    #include <thrust/device_vector.h>
    #include <thrust/host_vector.h>
    #include <thrust/sequence.h>
    #include <thrust/copy.h>
    #include <cstdlib>
    #include <iostream>
    
    struct my_even_odd{
            int s;
            my_even_odd(size_t _s) {s = (_s+1)>>1;}
            __host__ __device__
            int operator()(int i){
                    if (i < s) return 2*i;
                    return 2*(i-s)+1;}
    };
    
    int main(int argc, char *argv[]){
      int ds = 7;
      if (argc > 1) ds = atoi(argv[1]);
      thrust::device_vector<int> d(ds);
      thrust::device_vector<int> r(d.size());
      thrust::sequence(d.begin(), d.end());
      thrust::copy_n(thrust::make_permutation_iterator(d.begin(), thrust::make_transform_iterator(thrust::counting_iterator<int>(0), my_even_odd(d.size()))), d.size(), r.begin());
      thrust::host_vector<int> h_r = r;
      thrust::copy(h_r.begin(), h_r.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
    }
    $ nvcc -o t87 t87.cu
    $ ./t87 6
    0,2,4,1,3,5,
    $ ./t87 7
    0,2,4,6,1,3,5,
    $
    

    I think that should work for even and odd vector sizes, however I have not carefully tested it and make no claims that it is defect-free.